CUDA 缩减 - 基础知识答案

【问题标题】：CUDA reduction - basicsCUDA 缩减 - 基础知识
【发布时间】：2012-06-19 13:02:26
【问题描述】：

我正在尝试用这段代码对一个数组求和，但我被卡住了。我可能需要一些“CUDA for dummies 教程”，因为我花了很多时间在这样的基本操作上，但我无法让它工作。

这里列出了一些我不明白或不确定的事情：

我应该使用多少块 (dimGrid)？我认为应该是N/dimBlock.x/2（N=输入数组的长度），因为在内核开始时，数据从全局内存的两个“块”加载并添加到共享内存中
在原始代码中有blockSize。我将其替换为blockDim.x，因为我不知道这些变量有何不同。但是当blockSize = blockDim.x 时，gridSize = blockDim.x*2*gridDim.x 对我来说没有意义——gridSize 将大于 N。在一维数组的上下文中 *Dim.x 和 *Size 有什么区别？
主逻辑 - 在内核中，每个块总和 2*dimBlock（块中的线程）个数。当 N = 262144 和 dimBlock = 128 时，内核返回 1024 个部分和数组。然后我再次运行内核，结果 = 4 部分和。最后，在最后一次运行中，返回单个和，因为数组是由单个块处理的。
我对二进制数组求和。在第一次运行中，我可以使用uchar4 输入数据。在第二次和第三次运行中，我将使用int。

请告诉我我错过了什么

谢谢

__global__ void sum_reduction(uchar4* g_idata, int* g_odata, int N) { 

extern __shared__ int s_data[]; 

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockDim.x*2) + tid;
unsigned int gridSize = blockDim.x*2*gridDim.x;

while (i < N) {
    s_data[tid] += g_idata[i].x + g_idata[i+blockDim.x].x +
            g_idata[i].y + g_idata[i+blockDim.x].y +
            g_idata[i].z + g_idata[i+blockDim.x].z +
            g_idata[i].w + g_idata[i+blockDim.x].w;
    i += gridSize;
}
__syncthreads();

if (tid < 64) {
    s_data[tid] += s_data[tid + 64];
}
__syncthreads(); 

if (tid < 32) { 
    volatile int *s_ptr = s_data; 
    s_ptr[tid] += s_ptr[tid + 32];
    s_ptr[tid] += s_ptr[tid + 16];
    s_ptr[tid] += s_ptr[tid + 8]; 
    s_ptr[tid] += s_ptr[tid + 4];
    s_ptr[tid] += s_ptr[tid + 2]; 
    s_ptr[tid] += s_ptr[tid + 1]; 
} 
if (tid == 0) {
    g_odata[blockIdx.x] = s_data[0];
} 
}


main{
...
dim3 dimBlock(128);
dim3 dimGrid(N/dimBlock.x);
sum_reduction<<<dimGrid, dimBlock>>>(in, out, N);
...
}

【问题讨论】：

所写的内核有一个相当大的错误 - 主求和循环正在使用未初始化的内存，因此结果应该是完全不可预测的。此外，内核启动缺少共享内存大小，这会导致内核运行时错误（或者这只是一个错字？）。
@talonmies 你又救了我！我添加了int smemSize = 128 * sizeof(int); 并使用 smemSize 参数调用内核。结果是正确的。谢谢 PS我之前没有运行时错误。内存初始化不影响结果（现在确实如此）。
请添加问题的答案，而不是作为评论。

标签： cuda reduction

【解决方案1】：

像这样调用内核可以解决问题。

dim3 dimBlock(128);
dim3 dimGrid(N/dimBlock.x);
int smemSize = dimBlock.x * sizeof(int);
sum_reduction<<<dimGrid, dimBlock, smemSize>>>(in, out, N);

【讨论】：

【解决方案2】：

好的，我认为你需要重新开始。在reduction上查看来自 NVIDiA 的此分步流程指南

【讨论】：