Giảm tổng số bằng CUDA: N là gì?

Question

Theo NVIDIA, this là hạt nhân giảm tổng nhanh nhất:

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

Tuy nhiên, tôi không hiểu các tham số "n". Bất kì manh mối nào? Tôi không nghĩ rằng đó là kích thước của mảng để giảm, vì trong vòng lặp while sẽ có tràn bộ đệm.

Answer 1

Tôi tin rằng bạn đã phát hiện ra lỗi đánh máy trong các trang trình bày (có thể là một cái gì đó như while(i + blockDim.x < n)).

Nếu bạn có một cái nhìn vào mã nguồn trong mẫu CUDA SDK "reduction", cơ thể của gần đây nhất reduce6 trông như thế này:

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

Lưu ý kiểm tra rõ ràng trong while, giúp ngăn chặn ngoài giới hạn truy cập vào g_idata. Nghi ngờ ban đầu của bạn là chính xác; n chỉ đơn giản là kích thước của mảng g_idata.