2013-04-03 32 views
2

我正在嘗試使用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

奇怪的是

  1. 這兩種方法的結果是不同的。這可能是由於函數調用開銷造成的,但無論如何奇怪的是,這種開銷在(N1,N2)=(2048,2048)時發生變化。
  2. 這兩種方法的結果幾乎是獨立的矩陣大小。
  3. 與表達式的「手動編碼」版本相比,結果有很大不同(我假設庫應該具有生產效率 - 性能權衡)。

需要注意的是,任何操作之前,我在熱身使用的代碼

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` 

回答

1

ArrayFire使用即時編譯器進行元素明智的操作(這包括算術,邏輯,三角和其他數學運算)。

這意味着直到D_D的值是由用戶或另一非JIT功能請求類似

D_D = A_D + B_D + C_D + 3.; 

被存儲作爲表達。

如果您使用af::eval()函數或eval()方法,您可以強制對這些表達式進行評估。

因此,對於您的特定問題,請使用D_D.eval()這兩種方法。對於第一種方法,您還需要執行af::sync()timeit()不需要明確同步。

+0

非常感謝您的回答。我已經應用了您的建議,但似乎仍然存在問題。我進一步編輯了我的帖子,以呈現最新的結果。看來,如果我「增加GPU的預熱量」,那麼我會得到更合理的結果。這些是確定的時機嗎?我錯過了什麼? – JackOLantern

+0

謝謝。我已經用新的結果更新了我的帖子,包括你推薦的附加'af :: sync()'。他們看起來很合理,因爲時間隨着矢量的大小而增加。但是,我現在應該對'timeit'事件採取任何行動(第二種方法)以獲得類似的結果嗎?我仍然會持續2ms的時間。提前感謝您的幫助。 – JackOLantern