我正在嘗試使用ArrayFire評估簡單的GPU元素矩陣操作的性能。ArrayFire中的時序
特別地,考慮
int N1 = something;
int N2 = something;
array A_D = constant(1.,N1*N2,1,f64);
array B_D = constant(1.,N1*N2,1,f64);
array C_D = constant(1.,N1*N2,1,f64);
array D_D = constant(1.,N1*N2,1,f64);
我想執行以下指令
D_D = A_D + B_D + C_D + 3.;
我使用兩種方法的定時。第一個是
timer time_last;
time_last = timer::start();
D_D = A_D + B_D + C_D + 3.;
double elapsed = timer::stop(time_last);
printf("elapsed time using start and stop = %g ms \n",1000.*elapsed);
的第二個被定義下列函數
void timing_test()
{
int N1 = something;
int N2 = something;
array A_D = constant(1.,N1*N2,1,f64);
array B_D = constant(1.,N1*N2,1,f64);
array C_D = constant(1.,N1*N2,1,f64);
array D_D = constant(1.,N1*N2,1,f64);
D_D = A_D + B_D + C_D + 3.;
}
,然後調用
printf("elapsed time using timeit %g ms \n", 1000.*timeit(timing_test));
我已獲得下列結果:
(N1,N2)=(256,256)
第一種方法= 0.0456ms
第二種方法= 0.264ms
(N1,N2)=(512,512)
第一種方法=0.0451ms
秒方法= 0.264ms
(N1,N2)=(1024,1024)
第一方法= 0.0457ms
第二種方法= 0.263ms
(N1,N2)=(2048,2048)
第一方法=0.127ms
秒方法= 0.265ms
我還使用了遵循根據
的「手編碼」版本cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
eval_matrix_wrap_handcoded(A_D,B_D,C_D,D_D,N1*N2);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
template <class T1, class T2, class T3, class T4>
__global__ inline void evaluation_matrix_handcoded(T1 *A_D, T2 *B_D, T3 *C_D, T4 *D_D, int NumElements)
{
const int i = blockDim.x * blockIdx.x + threadIdx.x;
if(i < NumElements) D_D[i]=A_D[i]+B_D[i]+C_D[i]+3.;
}
__host__ void eval_matrix_wrap_handcoded(double *A_D, double *B_D, double *C_D, double *D_D, int NumElements)
{
dim3 dimGrid(iDivUp(NumElements,dimBlock.x));
evaluation_matrix_handcoded<<<dimGrid,dimBlock>>>(A_D,B_D,C_D,D_D,NumElements);
}
獲得以下
(N1,N2)=(256,256)
0.0897ms
(N1,N2)=(512,512)
0.339ms
(N1,N2)=(1024,1024)
1.3ms
(N1,N2)=(2048,2048)
5.37ms
奇怪的是
- 這兩種方法的結果是不同的。這可能是由於函數調用開銷造成的,但無論如何奇怪的是,這種開銷在
(N1,N2)=(2048,2048)
時發生變化。 - 這兩種方法的結果幾乎是獨立的矩陣大小。
- 與表達式的「手動編碼」版本相比,結果有很大不同(我假設庫應該具有生產效率 - 性能權衡)。
需要注意的是,任何操作之前,我在熱身使用的代碼
array test1(1,5);
test1(0,0)=1;
test1(0,1)=2;
test1(0,2)=3;
test1(0,3)=4;
test1(0,4)=5;
有人能幫助我解釋上述結果的GPU?謝謝。
EDIT FOLLOWING PAVAN的回答
首先方法修飾以修飾以
void timing_test()
{
int N1 = something;
int N2 = something;
array A_D = constant(1.,N1*N2,1,f64);
array B_D = constant(1.,N1*N2,1,f64);
array C_D = constant(1.,N1*N2,1,f64);
array D_D = constant(1.,N1*N2,1,f64);
D_D = A_D + B_D + C_D + 3.;
D_D.eval();
}
然而
timer time_last;
time_last = timer::start();
D_D = A_D + B_D + C_D + 3.;
D_D.eval();
af::sync();
double elapsed = timer::stop(time_last);
printf("elapsed time using start and stop = %g ms \n",1000.*elapsed);
方法二,定時現在是
`(N1,N2)=(256,256)` first approach = `14.7ms` second approach = `2.04ms`
`(N1,N2)=(512,512)` first approach = `14.3ms` second approach = `2.04ms`
`(N1,N2)=(1024,1024)` first approach = `14.09ms` second approach = `2.04ms`
`(N1,N2)=(2048,2048)` first approach = `16.47ms` second approach = `2.04ms`
我仍然有不同的時間和獨立的矢量大小。
如果我修改第一種方法來
D_D = A_D + B_D + C_D + 3.;
D_D.eval();
timer time_last;
time_last = timer::start();
D_D = A_D + B_D + C_D + 3.;
D_D.eval();
af::sync();
double elapsed = timer::stop(time_last);
printf("elapsed time using start and stop = %g ms \n",1000.*elapsed);
即,I 「增加的」 GPU熱身階段,我獲得,對於第一方法,
`(N1,N2)=(256,256)` `0.19ms`
`(N1,N2)=(512,512)` `0.42ms`
`(N1,N2)=(1024,1024)` `1.18ms`
`(N1,N2)=(2048,2048)` `4.2ms`
出現更因爲時間取決於數據大小並且更接近於手工編碼。
第二個編輯 總結:我已經佔了答案和註釋,並且對於第一種方法,我使用
D_D = A_D + B_D + C_D + 3.;
D_D.eval();
timer time_last;
af::sync();
time_last = timer::start();
D_D = A_D + B_D + C_D + 3.;
D_D.eval();
af::sync();
double elapsed = timer::stop(time_last);
printf("elapsed time using start and stop = %g ms \n",1000.*elapsed);
我獲得以下(新的)結果:
`(N1,N2)=(256,256)` `0.18ms`
`(N1,N2)=(512,512)` `0.30ms`
`(N1,N2)=(1024,1024)` `0.66ms`
`(N1,N2)=(2048,2048)` `2.18ms`
非常感謝您的回答。我已經應用了您的建議,但似乎仍然存在問題。我進一步編輯了我的帖子,以呈現最新的結果。看來,如果我「增加GPU的預熱量」,那麼我會得到更合理的結果。這些是確定的時機嗎?我錯過了什麼? – JackOLantern
謝謝。我已經用新的結果更新了我的帖子,包括你推薦的附加'af :: sync()'。他們看起來很合理,因爲時間隨着矢量的大小而增加。但是,我現在應該對'timeit'事件採取任何行動(第二種方法)以獲得類似的結果嗎?我仍然會持續2ms的時間。提前感謝您的幫助。 – JackOLantern