cuda - CUDA 中的协作组

Question

自从 CUDA 9 发布以来，显然可以将不同的线程和 block 分组到同一组中，以便您可以一起管理它们。这对我来说非常有用，因为我需要启动一个包含多个 block 的内核并等待所有 block 都同步(cudaThreadSynchronize() 对我来说不值得，因为线程同步后我必须继续在我的内核中工作)。

我的想法是将这些线程 block 包含在同一个组中，并等待所有线程都同步，如 Nvdia 主页示例所示。

他们做这样的事情:

__device__ int reduce_sum(thread_group g, int *temp, int val)
{
    int lane = g.thread_rank();

// Each iteration halves the number of active threads
// Each thread adds its partial sum[i] to sum[lane+i]
for (int i = g.size() / 2; i > 0; i /= 2)
{
    temp[lane] = val;
    g.sync(); // wait for all threads to store
    if(lane<i) val += temp[lane + i];
    g.sync(); // wait for all threads to load
}

我的问题是如何将这些 block 分组到 g 组中。这就是我最初启动内核的方式:

asap << <5, 1000 >> > (cuda_E2, cuda_A2, cuda_temp, Nb, *binM, Nspb);

每当我尝试使用 thread_group 时，编译器都会说它是未定义的。我正在使用 cooperative_groups.h header 。

有人知道怎么处理吗？提前致谢。

Answer 1

引自documentation :

Cooperative Groups requires CUDA 9.0 or later. To use Cooperative Groups, include the header file:

#include <cooperative_groups.h>

and use the Cooperative Groups namespace:

using namespace cooperative_groups; 

Then code containing any intra-block Cooperative Groups functionality can be compiled in the normal way using nvcc.

命名空间是您所缺少的。