CUDA 9中未找到grid_group。

Question

我尝试在CUDA 9中使用协作组（Cooperative Groups），但是编译时出现了错误。
有人知道解决方案吗？

开发环境如下：

• CUDA 9
• Kepler K80
• 计算能力：3.7

#include <cstdint>
#include <iostream>
#include <vector>

#include <cooperative_groups.h>

__global__
void kernel(uint32_t values[])
{
    using namespace cooperative_groups;

    grid_group g = this_grid();
}

int main(void)
{
    constexpr uint32_t kNum = 1 << 24; 
    std::vector<uint32_t> h_values(kNum);
    uint32_t *d_values;

    cudaMalloc(&d_values, sizeof(uint32_t) * kNum);
    cudaMemcpy(d_values, h_values.data(), sizeof(uint32_t) * kNum, cudaMemcpyHostToDevice);

    const uint32_t thread_num = 256;
    const dim3 block(thread_num);
    const dim3 grid((kNum + block.x - 1) / block.x);
    void *params[] = {&d_values};

    cudaLaunchCooperativeKernel((void *)kernel, grid, block, params);

    cudaMemcpy(h_values.data(), d_values, sizeof(uint32_t) * kNum, cudaMemcpyDeviceToHost);

    cudaFree(d_values);

    return 0;
}

$ nvcc -arch=sm_37 test.cu --std=c++11 -o test
test.cu(12): error: identifier "grid_group" is undefined
test.cu(12): error: identifier "this_grid" is undefined

- kotatsu

Answer 1

grid_group 功能仅支持 Pascal 架构及更高版本。您可以尝试编译 sm_60 等架构（当然，可执行文件无法在您的 GPU 上运行）。此外，您需要启用可重定位设备代码（-rdc=true）。不幸的是，编程指南对此并不十分清楚。我在其中找不到这些信息。但是，在 devblog.nvidia.com 的一些帖子中提到了这一点。从 https://devblogs.nvidia.com/cuda-9-features-revealed/ 或者 https://devblogs.nvidia.com/cooperative-groups/ 的最后可以看到相关信息。

- havogt