パラレルリダクション＃5最後のワープを展開

Question

私はマークハリスによる有名なスライドからの削減を研究しています。特に、最適化ステップ＃5を実装しましたが、私は間違った結果を得ました：41の代わりに17を使用しました。スライドに表示されているのと同じ番号シーケンスを使用しました。カーネルとホストコードは同じ.cuファイル内にあるので、共有配列の "extern"というコードでは省略しました。このコードTHREAD_PER_BLOCKで

#include <stdio.h> 
#include <cuda_runtime.h> 

#define THREAD_PER_BLOCK 16 

__global__ void reduce5(int *g_idata, int *g_odata) { 
    __shared__ int sdata[THREAD_PER_BLOCK]; 
    // perform first level of reduction, 
    // reading from global memory, writing to shared memory 
    unsigned int tid = threadIdx.x; 
    unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x; 
    sdata[tid] = g_idata[i] + g_idata[i+blockDim.x]; 
    __syncthreads(); 
    // do reduction in shared mem 
    for (unsigned int s=blockDim.x/2; s>32; s>>=1) { 
     if (tid < s) sdata[tid] += sdata[tid + s]; 
     __syncthreads(); 
    } 

    if (tid < 32) 
    { 
     sdata[tid] += sdata[tid + 32]; 
     sdata[tid] += sdata[tid + 16]; 
     sdata[tid] += sdata[tid + 8]; 
     sdata[tid] += sdata[tid + 4]; 
     sdata[tid] += sdata[tid + 2]; 
     sdata[tid] += sdata[tid + 1]; 
    } 
    // write result for this block to global mem 
    if (tid == 0) g_odata[blockIdx.x] = sdata[0]; 
} 

int main() 
{ 
    int inputLength=16; 
    int hostInput[16]={10,1,8,-1,0,-2,3,5,-2,-3,2,7,0,11,0,2}; 
    int hostOutput=0; 
    int *deviceInput; 
    int *deviceOutput; 

    cudaMalloc((void **)&deviceInput, inputLength * sizeof(int)); 
    cudaMalloc((void **)&deviceOutput, sizeof(int)); 

    cudaMemcpy(deviceInput, hostInput, inputLength * sizeof(int),cudaMemcpyHostToDevice); 

    reduce5<<<1,16>>>(deviceInput, deviceOutput); 

    cudaDeviceSynchronize(); 

    cudaMemcpy(&hostOutput, deviceOutput,sizeof(int), cudaMemcpyDeviceToHost); 

    printf("%d\n",hostOutput); 

    cudaFree(deviceInput); 
    cudaFree(deviceOutput); 

    return 0; 
} 

Answer 1

32の倍数でなければならず、入力データの少なくとも64の長さは、ブロックとグリッドサイズの二倍の製品でなければなりません。

これは表示されません（何らかのエラーチェックを実行していないため）、境界外の共有メモリーとグローバルなmemeoryアクセスのためスレッドとワープの削減に失敗します。

また、extern __shared__は、カーネルと他のコードが同じファイルに含まれているかどうかは関係ありません。コンパイル時に静的ではなく、実行時にその変数の共有メモリが動的に割り当てられることを意味します。割り当てのサイズは、カーネル起動構文の第3引数として渡されます。

Answer 2

同じ問題が発生し、変数がvolatileと宣言されていない場合、スレッドが実際には同期していないことが判明しました。

sdataを宣言すると、volatileを単に追加すると問題が解決します。 私の記事を参照してください：cuda Threads in A Warp appear to be not in synchronization