问带计数器的循环的CUDA索引

EN

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

看起来您正在使用计数器无间隙地包装您的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;
        }
    }
}

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

例如，如果你有二维的块，你可以使用Thriadix.x和Thriadix.y作为索引，并添加工作组的偏移量(blokidx.x*block Dim.x)作为偏移量。

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

如果可能，您可以使用cudpp或thrust (库，它们实现诸如remove_if或compact之类的并行函数，就像您在示例中所拥有的那样)。

Cudpp

Thrust

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