CUDA：如何在 GPU 中将数组的所有元素相加为一个数？
CUDA: how to sum all elements of an array into one number within the GPU?
首先，让我声明，我完全知道我的问题已经被问到：Block reduction in 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 内生成的所有值求和并将结果保存在一个名为 result 的 int 数组中，位于此类内的某个位置对应于当前迭代的数组。它必须在内核内部或至少仍在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-Beta/x86_website/projects/reduction/doc/reduction.pdf
然而，虽然我完全理解此类幻灯片中描述的代码步骤以及一般优化，但我不明白如果代码实际上是一个完整的数组（和一个不明确的维度）。有人能解释一下吗，并给我举个例子说明它是如何工作的（即如何从输出数组中获取一个数字）？
现在，回到我开头提到的那个问题(Block reduction in CUDA)。请注意，其接受的答案仅建议阅读我在上面 link 编辑的 pdf - 而不是 讨论如何处理代码生成的输出数组。在评论中，那里的 OP 提到 he/she 能够通过对 CPU 处的输出数组求和来完成这项工作——这是我做不到的，因为这意味着每次迭代都要下载输出数组我的 while 循环。最后，link 中的第三个答案建议使用库来完成此操作 - 但我有兴趣学习这样做的本机方法。
或者，我也对关于如何实现我上面描述的任何其他命题非常感兴趣。
如果你添加 2 个连续的数字，并保存结果，在你保存这些数字的任何插槽中，你只需要 运行，多次相同的内核，继续减少 2对数组的总和求幂，如本例所示：
对值求和的数组：
[·1,·2,·3,·4,·5,·6,·7,·8,·9,·10]
首先 运行 n/2 线程，对连续的数组元素求和，并将其存储在每个线程的 "left" 上，数组现在看起来像：
[·3,2,·7,4,·11,6,·15,8,·19,10]
运行 相同的内核，运行 n/4 线程，现在添加每 2 个元素，并将其存储在最左边的元素上，数组现在看起来像：
[·10,2,7,4,·26,6,15,8,·19,10]
运行同一个内核，运行n/8个线程，现在每4个元素相加，存入数组最左边的元素，得到：
[·36,2,7,4,26,6,15,8,·19,10]
运行最后一次，单线程将每8个元素相加，存入数组最左边的元素，得到：
[55,2,7,4,26,6,15,8,19,10]
这样，你只需要运行你的内核以一些线程作为参数，就可以得到最后的redux，在第一个元素（55）中查看"dots"（· ) 查看数组中的哪些元素是 "active" 对它们求和，每个 运行.
关于块并行缩减的规范信息你已经找到了，我就不再赘述了。如果您不想自己编写很多新代码来执行此操作，我建议您查看 CUB 库 block_reduce implementation，它提供了一个最佳的块式缩减操作，并添加了大约 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);
}
这里的关键是 atomic function 用于使用给定块的结果安全地更新内核 运行 结果，而不会出现内存竞争。您绝对 必须 在 运行 内核之前初始化 iter_result ，否则代码将无法运行，但这是基本的内核设计模式。
首先，让我声明，我完全知道我的问题已经被问到：Block reduction in 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 内生成的所有值求和并将结果保存在一个名为 result 的 int 数组中，位于此类内的某个位置对应于当前迭代的数组。它必须在内核内部或至少仍在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-Beta/x86_website/projects/reduction/doc/reduction.pdf
然而，虽然我完全理解此类幻灯片中描述的代码步骤以及一般优化，但我不明白如果代码实际上是一个完整的数组（和一个不明确的维度）。有人能解释一下吗，并给我举个例子说明它是如何工作的（即如何从输出数组中获取一个数字）？
现在，回到我开头提到的那个问题(Block reduction in CUDA)。请注意，其接受的答案仅建议阅读我在上面 link 编辑的 pdf - 而不是 讨论如何处理代码生成的输出数组。在评论中，那里的 OP 提到 he/she 能够通过对 CPU 处的输出数组求和来完成这项工作——这是我做不到的，因为这意味着每次迭代都要下载输出数组我的 while 循环。最后，link 中的第三个答案建议使用库来完成此操作 - 但我有兴趣学习这样做的本机方法。
或者，我也对关于如何实现我上面描述的任何其他命题非常感兴趣。
如果你添加 2 个连续的数字，并保存结果，在你保存这些数字的任何插槽中，你只需要 运行，多次相同的内核，继续减少 2对数组的总和求幂，如本例所示：
对值求和的数组：
[·1,·2,·3,·4,·5,·6,·7,·8,·9,·10]
首先 运行 n/2 线程，对连续的数组元素求和，并将其存储在每个线程的 "left" 上，数组现在看起来像：
[·3,2,·7,4,·11,6,·15,8,·19,10]
运行 相同的内核，运行 n/4 线程，现在添加每 2 个元素，并将其存储在最左边的元素上，数组现在看起来像：
[·10,2,7,4,·26,6,15,8,·19,10]
运行同一个内核，运行n/8个线程，现在每4个元素相加，存入数组最左边的元素，得到：
[·36,2,7,4,26,6,15,8,·19,10]
运行最后一次，单线程将每8个元素相加，存入数组最左边的元素，得到：
[55,2,7,4,26,6,15,8,19,10]
这样，你只需要运行你的内核以一些线程作为参数，就可以得到最后的redux，在第一个元素（55）中查看"dots"（· ) 查看数组中的哪些元素是 "active" 对它们求和，每个 运行.
关于块并行缩减的规范信息你已经找到了，我就不再赘述了。如果您不想自己编写很多新代码来执行此操作，我建议您查看 CUB 库 block_reduce implementation，它提供了一个最佳的块式缩减操作，并添加了大约 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);
}
这里的关键是 atomic function 用于使用给定块的结果安全地更新内核 运行 结果，而不会出现内存竞争。您绝对 必须 在 运行 内核之前初始化 iter_result ，否则代码将无法运行，但这是基本的内核设计模式。