Redukcja sumy za pomocą CUDA: Co to jest N?

Question

Według NVIDIA this jest najszybszym jądro redukcja 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]; 
} 

Jednak nie rozumiem parametru „n”. Jakieś wskazówki? Nie sądzę, że jest to rozmiar tablicy do zredukowania, ponieważ w pętli while wystąpiłoby przepełnienie bufora.

Answer 1

Wierzę, że odkryłeś literówkę na slajdach (prawdopodobnie powinno to być coś w rodzaju while(i + blockDim.x < n)).

Jeśli przyjrzeć kodu źródłowego w próbce CUDA SDK "reduction", ciało ostatniego reduce6 wygląda następująco:

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

Uwaga wyraźne check w while który zapobiega poza granicami dostęp do g_idata. Twoje początkowe podejrzenie jest poprawne; n to po prostu rozmiar tablicy g_idata.