Reducción de suma con CUDA: ¿Qué es N?

Question

Según NVIDIA, this es el núcleo más rápida reducción de la suma:

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]; 
} 

Sin embargo, no entiendo el parámetro "n". ¿Alguna pista? No creo que sea del tamaño de la matriz para reducir, ya que en el ciclo while habría un desbordamiento de búfer.

Answer 1

Creo que has descubierto un error tipográfico en las diapositivas (probablemente debería ser algo así como while(i + blockDim.x < n)).

Si se echa un vistazo al código fuente en la muestra CUDA SDK "reduction", el cuerpo de la más reciente reduce6 se parece a esto:

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; 
    } 

Nota la comprobación explícita dentro del while que impide fuera de los límites acceso a g_idata. Su sospecha inicial es correcta; n es simplemente el tamaño de la matriz g_idata.