-1

カーネル内にシーケンシャル セクションがあり、実際に速度が低下しています。ただし、内側のループを取り除く方法がわかりません。ここに何か提案はありますか？

__global__ void myKernel( int keep, int inc, int width, int* d_Xnum,
 int* d_Xco, bool* d_Xvalid,int* d_A )
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  int j = blockIdx.y * blockDim.y + threadIdx.y;

  int k1;

  if( i < keep && j <= i){
    int counter = 0;

    for(k1 = 0; k1 < inc; k1++){
      if(d_Xvalid[j*inc + k1] == 0)
         counter += (d_Xvalid[i*inc + d_Xco[j*width + k1]]);
    }

    d_A[i*keep+j] = inc - d_Xnum[i] - counter;
  }
}

削除k1することで、コードがかなり高速化されると思います。しかし、私はそれを使用してそれを行う方法がわかりませんcounter。提案、アイデア、考えは大歓迎です！このカーネルは次のように呼ばれます。

         ...
  int t = 32;
  int b = keep/(32)+1;
  int b2 = (inc/32)+1;
  dim3 thread (t, t);
  dim3 block (b, inc);

  // kernel call
  myKernel<<<block, thread>>>(k, inc, width, d_Xnum,
                  d_Xco, d_Xvalid, d_A);
  cudaThreadSynchronize();
            ...

keep約9000とinc約20000です

4

1 に答える 1

2

これはあなたの質問に対する正確な答えではありませんが、おそらくコードを最適化し、k1を取り除くため、合計の並列削減を実装するのに役立つ可能性がありif( i < keep && j <= i)ます。テクスチャを使用してこれらの読み取り専用ベクトルにアクセスするなど、GPU モデルに応じて実装できる他の最適化があります。

インデックスを生成する方法が原因で、他のスレッドが終了するのを待って多くのスレッドが停止します。スレッドを起動していますが、実際に何かを行っているkeep*incのは最大数のみです (条件のため)。keep*(keep+1)/2j <= i

以下の変更で改善できると思います。

  1. 起動keep*(keep+1)/2スレッド

  2. コードに次のことを行います

    __global__ void myKernel( int keep, int inc, int width, int* d_Xnum,
int* d_Xco, bool* d_Xvalid,int* d_A )
{
  int k = blockIdx.x * blockDim.x + threadIdx.x;
  int i = (int)(sqrt(0.25+2.0*k)-0.5); 
  int j = k - i*(i+1)/2;

  int k1;
  if( i < keep && j < inc){
    int counter = 0;

    for(k1 = 0; k1 < inc; k1++){
      if(d_Xvalid[j*inc + k1] == 0)
         counter += (d_Xvalid[i*inc + d_Xco[j*width + k1]]);
    }

    d_A[i*keep+j] = inc - d_Xnum[i] - counter;
  }
}

あなたがしていること (スレッドをkeep = 4起動4*4 = 16する場合、最良のシナリオでは。そのinc > keepように思われる場合、さらに多くのスレッドを起動している場合) は、(各「ボックス」はスレッドです) と見なすことができます。

_________________________________
| i = 0 | i = 0 | i = 0 | i = 0 |
| j = 0 |   -   |   -   |   -   |
_________________________________
| i = 1 | i = 1 | i = 1 | i = 1 |
| j = 0 | j = 1 |   -   |   -   |
_________________________________
| i = 2 | i = 2 | i = 2 | i = 2 |
| j = 0 | j = 1 | j = 2 |   -   |
_________________________________
| i = 3 | i = 3 | i = 3 | i = 3 |
| j = 0 | j = 1 | j = 2 | j = 3 |
_________________________________

インデックスを追加しkて生成iしj、必要に応じて生成することをお勧めします（keep = 4起動(4*(4+1)/2 = 10スレッド用）

_________________________________________________________________________________
| k = 0 | k = 0 | k = 1 | k = 0 | k = 1 | k = 2 | k = 0 | k = 1 | k = 2 | k = 3 |
| i = 0 | i = 1 | i = 1 | i = 2 | i = 2 | i = 2 | i = 3 | i = 3 | i = 3 | i = 3 |
| j = 0 | j = 0 | j = 1 | j = 0 | j = 1 | j = 2 | j = 0 | j = 1 | j = 2 | j = 3 |
_________________________________________________________________________________

これはで行うことができます

  • i = (int)(sqrt(0.25+2*k)-0.5)

  • j = k - i*(i+1)/2

これをレシピとして受け入れるか、その背後にある数学を少し調べることができます.

ここにたどり着くには、 ( k = 0+1+2+...+i = i*(i+1)/2j = 0 ) があることを知っています。ここで、この方程式を解くと、 の方程式が得られます(int のキャストは切り捨てられ、正しい結果が得られるようになります)。取得するには、 0の場合の値を減算する必要があります。i*(i+1)/2 = kij!=0jkji*(i+1)/2

于 2012-10-03T02:12:42.503 に答える