0

この次の CUDA-C の部分並列和アルゴリズムの実装を理解するのを手伝ってもらえますか? partialSum共有配列の初期フィルアップを理解するのに問題があります[3 行目から 8 行目]。私は今それを何時間も追跡しましたが、次のコードで開始する必要がある理由がわかりませ2*blockIdx.x*blockDim.x;んblockIdx.x*blockDim.x;。

ホスト コード:

numOutputElements = numInputElements / (BLOCK_SIZE<<1);
 if (numInputElements % (BLOCK_SIZE<<1)) {
     numOutputElements++;
 }
#define BLOCK_SIZE 512
dim3 dimGrid(numOutputElements, 1, 1);
dim3 dimBlock(BLOCK_SIZE, 1, 1);
total<<<dimGrid, dimBlock>>>(deviceInput, deviceOutput, numInputElements);

カーネルコード:

1    __global__ void total(float * input, float * output, int len) {
2    
3    __shared__ float partialSum[2*BLOCK_SIZE];
4      
5      unsigned int t = threadIdx.x;
6      unsigned int start = 2*blockIdx.x*blockDim.x;
7      partialSum[t] = input[start + t];
8      partialSum[blockDim.x + t] = input[start + blockDim.x + t];
9     
10    for (unsigned int stride = blockDim.x; stride >=1; stride >>=1)
11      {
12       __syncthreads();
13        
14       if (t < stride)
15         partialSum[t] += partialSum[t + stride];
16      }
17      output[blockIdx.x] = partialSum[0];   
18  }

合計する要素が 10 個あり、ブロックサイズを 4、ブロックあたり 4 スレッドにすることを選択したとします。したがって、3 つのブロックが使用されますよね? 【ワープのサイズなどは一旦忘れましょう】

blockIdx.x が 2 (2 つの要素を持つ最後のブロック) の場合、開始は (2*2*4=)16 になり、10 より大きく、input長さを超えます (したがって、partialSum[t]との両方partialSum[blockDim.x + t]が変更されず、block2の共有メモリは空のままです。) もしそうなら、配列の最後の 2 つの要素が失われます!!

blockIdx.x、blockDim.x を間違った方法で取得していると思われます。誰かが私を修正してもらえますか? お願いします！

4

2 に答える 2

3

ブロックの半分だけを起動し、ブロックごとに 2 倍の作業を行います。これを行う利点は、部分的な合計を保存するために必要なスクラッチ スペースが半分になることです (半分のブロックしか起動していないため)。

リダクション (この場合は合計) を行う通常の方法は、次のようにすることです。

1    __global__ void total(float * input, float * output, int len) {
2    
3    __shared__ float partialSum[BLOCK_SIZE];
4      
5     unsigned int t = threadIdx.x;
6     unsigned int start = blockIdx.x*blockDim.x;
7     partialSum[t] = 0;
8     for (int T = start; T < len; T += blockDim.x * gridDim.x) 
9        partialSum[t] += input[T];
10    for (unsigned int stride = blockDim.x/2; stride >=1; stride >>=1)
11      {
12       __syncthreads();
13        
14       if (t < stride)
15         partialSum[t] += partialSum[t + stride];
16      }
17      output[blockIdx.x] = partialSum[0];   
18  }

len = 1024したがって、とがあれば、 BLOCK_SIZE = 256<= 4 ブロックなら何でも起動できます。

さまざまな数のブロックを起動したときに、8 行目と 9 行目に含まれる for ループで何が起こるかを見てみましょう。また、出力には要素数 == ブロック数が必要であることに注意してください。

  • Blocks == 4つまり、blockDim.x * gridDim.x= 256 x 4 = 1024 なので、1 回だけ繰り返されます。出力への結合されていない書き込みの数 = 4。
  • Blocks == 2つまり、blockDim.x * gridDim.x= 256 x 2 = 512 なので、2 回繰り返します。出力への結合されていない書き込みの数 = 2。
  • Blocks == 1つまり、blockDim.x * gridDim.x= 256 x 1 = 256 なので、4 回繰り返します。出力への結合されていない書き込みの数 = 1。

したがって、より少ないブロックを起動すると、メモリ フット プリントが削減され、グローバルな書き込みも削減されるという利点があります。ただし、並列処理が減少します。

理想的には、アルゴリズムに最適な組み合わせをヒューリスティックに見つける必要があります。または、それを行う既存のライブラリを使用することもできます。

問題のカーネルは、パフォーマンスを向上させるために半分のブロックを起動することを選択しています。ただし、2 倍の共有メモリを使用する必要はない場合があります。

于 2013-01-12T07:38:25.840 に答える
1

ブロックのカウントに問題がある場合があります。

合計する要素が 10 個あり、ブロックサイズを 4、ブロックあたり 4 スレッドにすることを選択した場合、使用中のブロックは 2 つだけになります。

カーネルコードによると、各スレッドはグローバルデバイス mem の 2 つの要素を担当するためです。

各スレッドが読み込む入力要素を以下に示します。コードに範囲チェックが見られません。したがって、10要素に十分なゼロパディングがあると思います。

blockIdx.x           : 0 0 0 0  1 1 1 1  2 2 2 2  3 3 3 3
threadIdx.x          : 0 1 2 3  0 1 2 3  0 1 2 3  0 1 2 3
linear thread id     : 0 1 2 3  4 5 6 7  8 9 a b  c d e f

Idx of the element     0 1 2 3  8 9
read by the thread   : 4 5 6 7

そのoutput[0]ため、要素 0 ～ 7output[1]の合計が格納され、要素 8 ～ 9 の合計が格納されます。何も失われていないと思います。

CUDA での並列削減の最適化のカーネル 4 を参照して、なぜ2*. 低速のカーネル 3 と @Pavan の回答で指定されたカーネルは、各スレッドが 1 つの要素のみを担当する同様の実装です。

于 2013-01-12T12:26:51.803 に答える