0

我有一个嵌套循环，中间有一个计数器。我已经设法将 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

4

3 回答 3

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;
        }
    }
}
于 2012-10-01T11:03:39.110 回答
1

请注意，您可以使用 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 的原子增量功能。

于 2012-10-01T11:18:34.340 回答
1

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

库普

推力

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

于 2012-10-01T13:48:13.427 回答