2010-12-08 16 views
0

を実行するために100倍の時間をとります。(私はすべてのビットが間違っているものを見るために変えるんだ、とコンパイラは何をやっているので、それはめちゃくちゃです)パフォーマンスの問題は、私はこれを持って

__device__ inline int f(int i, int j, int value) 
{ 
    int x; 
    int y; 
    int delta; 

    int* p = p_new_solution; 
    int pitch = p_new_solution_pitch; 

    int* p_row_i = (int*)((char*)p + i * pitch); 
    int p_i = p_row_i[threadIdx.x + blockIdx.x * blockDim.x]; 

    int* p_row_j = (int*)((char*)p + j * pitch); 
    int p_j = p_row_j[threadIdx.x + blockIdx.x * blockDim.x]; 

    delta = (tex2D(A_matrix, i, i) - tex2D(A_matrix, j, j)) * (tex2D(B_matrix, p_j, p_j) - tex2D(B_matrix, p_i, p_i)); 
    delta += (tex2D(A_matrix, i, j) - tex2D(A_matrix, j, i)) * (tex2D(B_matrix, p_j, p_i) - tex2D(B_matrix, p_i, p_j)); 


    for(int k = 0 ; k < n ; k++) 
    { 
     int* p_row = (int*)((char*)p + k * pitch); 
     int p_k = p_row[threadIdx.x + blockIdx.x * blockDim.x]; 

     int A_ki = tex2D(A_matrix, k, i); 
     int A_kj = tex2D(A_matrix, k, j); 
     int A_ik = tex2D(A_matrix, i, k); 
     int A_jk = tex2D(A_matrix, j, k); 
     int B_pkpj = tex2D(B_matrix, p_k, p_j); 
     int B_pkpi = tex2D(B_matrix, p_k, p_i); 
     int B_pjpk = tex2D(B_matrix, p_j, p_k); 
     int B_pipk = tex2D(B_matrix, p_i, p_k); 

     x = (A_ki - A_kj); 
     x *= (B_pkpj - B_pkpi); 

     y = (A_ik - A_jk); 
     y *= (B_pjpk - B_pipk); 

     x += y; 
    } 

    x -= ((tex2D(A_matrix, i, i) - tex2D(A_matrix, i, j)) * (tex2D(B_matrix, p_i, p_j) - tex2D(B_matrix, p_i, p_i))) + 
      ((tex2D(A_matrix, i, i) - tex2D(A_matrix, j, i)) * (tex2D(B_matrix, p_j, p_i) - tex2D(B_matrix, p_j, p_i))); 

    x -= ((tex2D(A_matrix, j, i) - tex2D(A_matrix, j, j)) * (tex2D(B_matrix, p_j, p_j) - tex2D(B_matrix, p_j, p_i))) + 
      ((tex2D(A_matrix, i, j) - tex2D(A_matrix, j, j)) * (tex2D(B_matrix, p_j, p_j) - tex2D(B_matrix, p_j, p_j))); 


    x += delta; 
    x *= 2; 

    return value; 
    //return x; 
} 

問題はこれらの2つのリターンステートメントにあります。もしi return valueなら、カーネル全体は300msのようになります。もしiになると約33000 msとなります。これの問題は何ですか?私はいくつか__syncthreads()を試しましたが、それでも同じ悪い時間がありました。

これらの戻り関数は最終的なコードではありません。戻り値を選択するにはelse文が必要ですが、valueまたはvalue + xになります。このif else文も長すぎます。

ありがとうございます。

+0

cuda-gdbで(戻り値バージョンで)実行してみてください。コンパイル後にx、y、pなどが存在しないことがわかります。実際には、この関数内のどの行をもステップすることはできません。迷惑なのは、異なるデバッグ最適化レベルでコンパイルしても、コンパイラを停止する方法を理解できなかったことです。 – jmilloy

答えて

2

測定している時間は変数を返す時間ではなく、xを計算する時間です。 NVCCは、xを返さない場合、その結果は決して使用されないので、何もしていないコードがたくさんあることを検出します。無駄なコードを取り除き、関数を高速化します。

+0

うん、私はそれをいくつかの時間前に分かった...ダング..それは速くなることはできません:/ – hfingler

1

return valueは、関数の引数のうちの1つを返すだけで、関数全体をノーオペレーションにします。私はそれが完全に最適化されていると思います。 return xの場合、実際の作業は33秒かかります。

関連する問題