CUDA 中的合作组

Question

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

我的想法是将这些线程块包含在同一个组中，然后等待它们全部同步，正如 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
}

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

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

每当我尝试使用 thread_group 时，编译器都会说它是 undefied。我正在使用 cooperative_groups.h 标头。

有谁知道如何处理这个问题？提前致谢。

Answer 1

引用自documentation：

合作组需要 CUDA 9.0 或更高版本。使用合作社 组，包括头文件：

#include <cooperative_groups.h>

并使用合作组命名空间：

using namespace cooperative_groups; 

然后代码包含任何 块内合作组功能可以在 使用 nvcc 的正常方式。

命名空间是您所缺少的。