3

万が一、配列の要素を効率的に分割する方法があるかどうか知りたいです。私はマトリックス値10000x10000で実行しており、他のカーネルと比較してかなりの時間がかかります。分割は費用のかかる作業であり、それを改善する方法がわかりません。

__global__ void division(int N, float* A, int* B){

  int row = blockIdx.x * blockDim.x + threadIdx.x;
  int col = blockIdx.y * blockDim.y + threadIdx.y;

  if((row < N) && (col <= row) ){
    if( B[row*N+col] >0 )
      A[row*N+col] /= (float)B[row*N+col];
  }

}

で起動されたカーネル

  int N = 10000;
  int threads = 32
  int blocks = (N+threads-1)/threads
  dim3 t(threads,threads);
  dim3 b(blocks, blocks);
  division<<< b, t >>>(N, A, B);
  cudaThreadSynchronize();

オプションB：

__global__ void division(int N, float* A, int* B){
  int k =  blockIdx.x * blockDim.x + threadIdx.x;
  int kmax = N*(N+1)/2 
  int i,j;
  if(k< kmax){
    row = (int)(sqrt(0.25+2.0*k)-0.5); 
    col = k - (row*(row+1))>>1;
    if( B[row*N+col] >0 )
      A[row*N+col] /= (float)B[row*N+col];
  }
}

で起動

  int threads =192;
  int totalThreadsNeeded = (N*(N+1)/2;
  int blocks = ( threads + (totalThreadsNeeded)-1 )/threads;
  division<<<blocks, threads >>>(N, A, B);

threadIdsが正しい場合でも、オプションBが間違った結果を返すのはなぜですか？ここに何が欠けていますか？

4

3 に答える 3

4

基本的な問題は、信じられないほど巨大なグリッド（10000x10000配列の例では1億を超えるスレッド）を起動していることです。カーネルのアクセスパターンは三角形であるため、これらのスレッドの半分は生産性を発揮しません。そのため、特に正当な理由もなく、膨大な量のGPUサイクルが無駄になっています。さらに、使用しているアクセスパターンでは、合体したメモリアクセスが許可されていないため、実際に有用な作業を行っているスレッドのパフォーマンスがさらに低下します。

私があなたの問題を正しく理解していれば、カーネルは正方形配列の下三角に対してのみ要素ごとの除算を実行しています。この場合、次のようなものを使用して同様に行うことができます。

__global__ 
void division(int N, float* A, int* B)
{
    for(int row=blockIdx.x; row<N; row+=gridDim.x) {
        for(int col=threadIdx.x; col<=row; col+=blockDim.x) {
            int val = max(1,B[row*N+col]);
            A[row*N+col] /= (float)val;
        }
    }
}

[免責事項：ブラウザで記述され、コンパイルされておらず、テストされておらず、自己責任で使用してください]

ここでは、1次元グリッドが使用され、各ブロックが一度に1行を計算します。ブロック内のスレッドは行に沿って移動するため、メモリアクセスは合体します。コメントで、GPUはTeslaC2050であるとおっしゃっています。そのデバイスは、14個のSMのそれぞれを完全に「満たす」ためにそれぞれ192スレッドの112ブロックを必要とするだけで、それぞれ8ブロックの完全な補完と、SMごとの同時スレッドの最大数を示します。したがって、起動パラメータは次のようになります。

int N = 10000;
int threads = 192;
int blocks = min(8*14, N);
division<<<blocks, threads>>>(N, A, B);

これは、現在のアプローチよりもかなり高速に実行されると思います。数値の精度がそれほど重要でない場合は、除算を近似の逆数の固有値と浮動小数点の乗算に置き換えることで、さらに高速化できる可能性があります。

于 2012-10-04T17:27:53.197 に答える
3

ifスレッドはワープと呼ばれる32のグループで実行されるため、両方の条件が1つのスレッドのみの場合、ワープ内の32のスレッドすべての分割に対して料金を支払うことになりますtrue。条件がfalse多くのスレッドに当てはまる場合は、別のカーネルで除算が不要な値を除外できるかどうかを確認してください。

intからfloatへの変換自体が遅い場合があります。その場合、前の手順で直接floatを生成し、floatの配列としてBを渡すことができる場合があります。

B配列を生成する前のステップで、反転数を生成できる場合があります。その場合、このカーネルでは除算の代わりに乗算を使用できます。(a / b == a * 1 / b)。

アルゴリズムによっては、精度の低い除算で解決できる場合があります。__fdividef(x, y)あなたが試すことができる本質的な、があります。コンパイラフラグもあり-prec-div=falseます。

于 2012-10-04T14:55:05.377 に答える
2

最初に確認するのは、統合されたメモリアクセスです。ここで非合体パターンの理由はありません。多くのメモリ帯域幅を浪費しないように、行と列を交換するだけです。

int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
...
A[row*N+col] ...

これがコンピューティング機能2.0以降で実行されている場合でも、キャッシュはこの次善のパターンを修正するのに十分な大きさではありません。

于 2012-10-04T17:30:34.110 に答える