在CUDA 9中找不到grid_group

Question

我尝试在CUDA 9中使用协作组,但是在编译时遇到错误.
有谁知道解决方案？

开发环境如下:

• CUDA 9
• 开普勒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

Answer 1

这些grid_group功能仅在Pascal架构及更高版本中受支持.您可以尝试编译,例如,sm_60(当然,可执行文件不会在您的GPU上运行).此外,您需要启用可重定位设备代码(-rdc=true).

不幸的是,编程指南对此并不十分清楚.我在那里找不到这些信息.但是在devblog.nvidia.com上的一些帖子中提到了它:

来自https://devblogs.nvidia.com/cuda-9-features-revealed/

虽然协作组适用于所有GPU架构,但随着GPU功能的发展,某些功能不可避免地取决于架构.所有体系结构都支持基本功能,例如将小于线程块的组同步到warp粒度,而Pascal和Volta GPU支持新的网格范围和多GPU同步组.

或者在https://devblogs.nvidia.com/cooperative-groups/的最后

Pascal和Volta GPU的新功能通过启用跨越在一个甚至多个GPU上运行的整个内核启动的线程组的创建和同步,帮助协作组进一步发展.