带有计数器的循环的 CUDA 索引答案

【问题标题】：CUDA indices for loops with counters带有计数器的循环的 CUDA 索引
【发布时间】：2012-09-22 04:36:29
【问题描述】：

我有一个嵌套循环，中间有一个计数器。我已经设法将 CUDA 索引用于外部循环，但我想不出任何方法可以在这种循环中利用更多的并行性。你有过类似的工作经验吗？

int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i < Nx) {
    counter = 0;
    for (k = 0; k < Ny; k++) {

        d_V[i*Ny + k] = 0;

        if ( d_X[i*Ny + k] >= 2e2 ) {

             /* do stuff with i and k and counter i.e.*/
                d_example[i*length + counter] = k;
                    ...
             /* increment counter */
             counter++;
        }
    }
}

我看到的问题是如何处理计数器，因为k 也可以在 CUDA 中用threadIdx.y + blockIdx.y * blockDim.y 进行索引

【问题讨论】：

您也可以使用 threadIdx.y 作为“k”变量。即 2D 块在 CUDA 中是可能的。您可以将 counter 定义为 shared 变量，并将其存储回全局内存。

标签： c++ cuda gpu gpgpu

【解决方案1】：

在循环迭代之间使用计数器/循环变量是并行化的自然对立面。理想的并行循环具有可以以任何顺序运行的迭代，彼此不知道。不幸的是，一个公共变量使它既依赖于顺序又相互感知。

您似乎正在使用计数器来打包您的 d_example 数组而没有间隙。通过浪费一些内存，这种事情很可能在计算时间上更有效；如果你让 d_example 中不会被设置的元素保持为零，通过低效打包 d_example，你可以在任何昂贵的计算步骤之后对 d_example 执行过滤。

实际上，您甚至可以在读取数组时将过滤留给修改过的迭代器，它只会跳过任何零值。如果零是数组中的有效值，只需使用特定的 NaN 值或单独的掩码数组。

int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i < Nx) {
    for (k = 0; k < Ny; k++) {

        d_V[i*Ny + k] = 0;

        if ( d_X[i*Ny + k] >= 2e2 ) {

             /* do stuff with i and k and counter i.e.*/
                d_example[i*length + i*k] = k;
                d_examask[i*length + i*k] = 1;
                    ...
             /* increment counter */
        } else {
             d_examask[i*length+i*k] = 0;
        }
    }
}

【讨论】：

@Phil H：我明白你的意思。虽然它会增加设备上的内存，但值得一试。

【解决方案2】：

请注意，您可以使用 threadIDx.y 作为数组中的第二个索引。更多信息请看这里：http://www.cs.sunysb.edu/~mueller/teaching/cse591_GPU/threads.pdf

例如，如果您有二维块，您可以使用 threadix.x 和 threadix.y 作为您的索引，并添加工作组的偏移量 (blockidx.x * blockDim.x) 作为您的偏移量。

由于 GPU 上的分支非常昂贵，并且给定工作组中的所有线程将始终等待组中的所有任务继续执行，因此最好简单地计算所有元素并丢弃不需要的元素，如果这是可能，这可能会完全避免使用计数器。如果没有，最好的解决方案是在 phoad 在他的评论中指定的全局计数器上使用 CUDA api 的原子增量功能。

【讨论】：

【解决方案3】：

如果可能，您可以使用 cudpp 或推力（库，实现并行功能，如 remove_if 或 compact - 什么，你有什么例子）。

Cudpp

Thrust

您可以在这些页面上找到简单的示例以及如何使用它们。我更喜欢cudpp，因为恕我直言比推力更快。

【讨论】：