CUDAによる減算：Nとは何ですか？

Question

はNVIDIAによると、thisは最速の合計削減カーネルです：

template <unsigned int blockSize> 
__device__ void warpReduce(volatile int *sdata, unsigned int tid) { 
if (blockSize >= 64) sdata[tid] += sdata[tid + 32]; 
if (blockSize >= 32) sdata[tid] += sdata[tid + 16]; 
if (blockSize >= 16) sdata[tid] += sdata[tid + 8]; 
if (blockSize >= 8) sdata[tid] += sdata[tid + 4]; 
if (blockSize >= 4) sdata[tid] += sdata[tid + 2]; 
if (blockSize >= 2) sdata[tid] += sdata[tid + 1]; 
} 
template <unsigned int blockSize> 
__global__ void reduce6(int *g_idata, int *g_odata, unsigned int n) { 
extern __shared__ int sdata[]; 
unsigned int tid = threadIdx.x; 
unsigned int i = blockIdx.x*(blockSize*2) + tid; 
unsigned int gridSize = blockSize*2*gridDim.x; 
sdata[tid] = 0; 
while (i < n) { sdata[tid] += g_idata[i] + g_idata[i+blockSize]; i += gridSize; } 
__syncthreads(); 
if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); } 
if (blockSize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); } 
if (blockSize >= 128) { if (tid < 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads(); } 
if (tid < 32) warpReduce(sdata, tid); 
if (tid == 0) g_odata[blockIdx.x] = sdata[0]; 
} 

しかし、私は "n" は、パラメータを理解していません。すべての手がかりは？ whileループでバッファオーバーフローが発生するため、減らす配列のサイズとは思えません。

Answer 1

あなたはスライドの入力ミスを発見したと思います（恐らくwhile(i + blockDim.x < n)のようなものでしょう）。

あなたはCUDA SDKサンプル"reduction"でソースコードを見てみる場合は、最新のreduce6の体は次のようになります。

template <class T, unsigned int blockSize, bool nIsPow2> 
__global__ void 
reduce6(T *g_idata, T *g_odata, unsigned int n) 
{ 
    T *sdata = SharedMemory<T>(); 

    // perform first level of reduction, 
    // reading from global memory, writing to shared memory 
    ... 

    T mySum = 0; 

    // we reduce multiple elements per thread. The number is determined by the 
    // number of active thread blocks (via gridDim). More blocks will result 
    // in a larger gridSize and therefore fewer elements per thread 
    while (i < n) 
    {   
     mySum += g_idata[i]; 
     // ensure we don't read out of bounds -- this is optimized away for powerOf2 sized arrays 
     if (nIsPow2 || i + blockSize < n) 
      mySum += g_idata[i+blockSize]; 
     i += gridSize; 
    } 

が範囲外に防ぎwhile内の明示的なチェックに注意してください。 g_idataにアクセスしてください。あなたの最初の疑惑は正しいです。 nは、単純にg_idata配列のサイズです。