And*_*eer 6 arrays cuda sum nvidia reduction
首先,请允许我声明我已经充分意识到我的问题已经被提出:CUDA中的减少但是,正如我希望明确的那样,我的问题是对此的后续行动,我有特殊的需求使该OP发现的解决方案不合适.
所以,让我解释一下.在我当前的代码中,我在while循环的每次迭代中运行一个Cuda内核来对数组的值进行一些计算.举个例子,想象如下:
int max_iterations = 1000;
int iteration = 0;
while(iteration < max_iterations)
{
__global__ void calcKernel(int* idata, int* odata)
{
int i = blockIdx.x*blockDim.x + threadIdx.x;
if (i < n)
{
odata[i] = (idata[i] + 2) * 5;
}
}
iteration++;
}
但是,接下来我必须为GPU完成看似艰巨的任务.在调用内核的while循环的每次迭代中,我必须对在其中生成的所有值求和odata并将结果保存在一个int数组中result,该数组在此类数组中与当前迭代对应的位置.它必须在内核中完成,或者至少仍然在GPU中完成,因为由于性能限制,我只能result在完成所有迭代后最终检索数组.
一个错误的天真尝试看起来像下面这样:
int max_iterations = 1000;
int iteration = 0;
while(iteration < max_iterations)
{
__global__ void calcKernel(int* idata, int* odata, int* result)
{
int i = blockIdx.x*blockDim.x + threadIdx.x;
if (i < n)
{
odata[i] = (idata[i] + 2) * 5;
}
}
result[iteration] = 0;
for(int j=0; j < max_iterations; j++)
{
result[iteration] += odata[j];
}
iteration++;
}
当然,由于GPU跨线程分发代码,上面的代码不起作用.为了学习如何正确地做到这一点,我一直在网站上阅读有关使用CUDA进行阵列缩减的其他问题.特别是,我发现了一个非常好的NVIDIA关于这个主题的pdf,我在开头提到的前SO问题中也讨论了这个问题:http://developer.download.nvidia.com/compute/cuda/1.1-贝塔/ x86_website /项目/还原/ DOC/reduction.pdf
然而,虽然我完全理解了这些幻灯片中描述的代码的步骤,以及一般的优化,但我不知道如果代码实际上是一个完整的数组,那么这种方法可以总结 - 将数组减少到一个数字(和一个不明确的维度).有人可以请一些关于它的说明,并向我展示一个如何工作的例子(即如何从输出数组中取出一个数字)？
现在,回到我在开头提到的那个问题(CUDA中的块减少).请注意,它接受的答案只是建议人们阅读上面链接的pdf - 这里没有讨论如何处理代码生成的输出数组.在评论中,OP提到他/她能够通过在CPU上对输出数组求和来完成工作 - 这是我不能做的事情,因为这意味着每次迭代我的while循环时都要下载输出数组.最后,该链接中的第三个答案表明使用库来实现这一目标 - 但我有兴趣学习本地的方法.
或者,我也会对如何实现我上面描述的任何其他命题非常感兴趣.
你已经找到了关于块并行缩减的规范信息，所以我不会重复。如果您不想自己编写大量新代码来执行此操作，我建议您查看 CUB 库block_reduce实现，它通过向现有内核添加大约 4 行代码来提供最佳的逐块缩减操作。
在真正的问题上，如果你做这样的事情，你可以做你想做的事：
__global__ void kernel(....., int* iter_result, int iter_num) {
// Your calculations first so that each thread holds its result
// Block wise reduction so that one thread in each block holds sum of thread results
// The one thread holding the adds the block result to the global iteration result
if (threadIdx.x == 0)
atomicAdd(iter_result + iter_num, block_ressult);
}
The key here is that an atomic function is used to safely update the kernel run result with the results from a given block without a memory race. You absolutely must initialise iter_result before running the kernel, otherwise the code won't work, but that is the basic kernel design pattern.
如果您添加 2 个连续的数字，并将结果保存在保存这些数字的任何插槽中，您只需运行多次相同的内核，即可继续减少数组总和的 2 次方，例如在这个例子中：
\n\n数组求和值：
\n\n[\xc2\xb71,\xc2\xb72,\xc2\xb73,\xc2\xb74,\xc2\xb75,\xc2\xb76,\xc2\xb77,\xc2\xb78,\xc2\xb79,\xc2\xb710]\n首先运行 n/2 个线程，对连续的数组元素求和，并将其存储在每个线程的“左侧”，数组现在如下所示：
\n\n[\xc2\xb73,2,\xc2\xb77,4,\xc2\xb711,6,\xc2\xb715,8,\xc2\xb719,10]\n运行相同的内核，运行 n/4 个线程，现在添加每 2 个元素，并将其存储在最左边的元素上，现在数组将如下所示：
\n\n[\xc2\xb710,2,7,4,\xc2\xb726,6,15,8,\xc2\xb719,10]\n运行相同的内核，运行n/8个线程，现在将每4个元素相加，并存储在数组中最左边的元素中，得到：
\n\n[\xc2\xb736,2,7,4,26,6,15,8,\xc2\xb719,10]\n最后运行一次，单线程将每8个元素相加，并存储在数组最左边的元素中，得到：
\n\n[55,2,7,4,26,6,15,8,19,10]\n这样，您只需使用一些线程作为参数运行内核，即可在最后获得 redux，在第一个元素 (55) 中查看“点”(\xc2\xb7) 以查看数组中的哪些元素每次运行时都“主动”对它们求和。
\n