不允许从 __device__ 函数调用 __host__ 函数的 cuda::cub 错误

Posted 2023-02-22

【中文标题】不允许从 __device__ 函数调用 __host__ 函数的 cuda::cub 错误

【英文标题】：cuda::cub error calling a __host__ function from a __device__ functionis not allowed

【发布时间】：2018-06-05 01:28:29

【问题描述】：

我使用 cub::DeviceReduce::Sum 来计算向量的总和，但它给了我错误：

error: calling a __host__ function("cub::DeviceReduce::Sum<double *, double *> ") from a __device__ function("dotcubdev") is not allowed
error: identifier "cub::DeviceReduce::Sum<double *, double *> " is undefined in device code

代码示例如下：

__device__ void sumcubdev(double* a, double *sum, int N)

    // Declare, allocate, and initialize device-accessible pointers 
     //for input and output
    // Determine temporary device storage requirements
    void     *d_temp_storage = NULL;
    size_t   temp_storage_bytes = 0;
     cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, a, sum, N);
    // Allocate temporary storage
    cudaMalloc(&d_temp_storage, temp_storage_bytes);
    // Run sum-reduction
     cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, a, sum, N);

代码在“main”正文中可以成功运行，但在函数中无法运行。

【参考方案1】：

要从设备代码中使用 cub 设备范围的函数，有必要构建您的项目以支持 CUDA 动态并行性。在cub documentation 中，这里指出：

使用注意事项 动态并行。可以在支持 CUDA 动态并行的设备上的内核代码中调用 DeviceReduce 方法。

例如，您可以编译显示的代码：

$ cat t1364.cu
#include <cub/cub.cuh>
__device__ void sumcubdev(double* a, double *sum, int N)

    // Declare, allocate, and initialize device-accessible pointers
     //for input and output
    // Determine temporary device storage requirements
    void     *d_temp_storage = NULL;
    size_t   temp_storage_bytes = 0;
     cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, a, sum, N);
    // Allocate temporary storage
    cudaMalloc(&d_temp_storage, temp_storage_bytes);
    // Run sum-reduction
     cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, a, sum, N);

$ nvcc -arch=sm_35 -dc t1364.cu
$

（CUDA 9.2，CUB 1.8.0）

这意味着 CUB 将启动子内核来完成工作。

这不是关于如何使用 CUDA 动态并行 (CDP) 的完整教程。以上只是编译命令，省略了链接步骤。 cuda 标签有很多关于 CDP 的问题，您可以在两个 blog articles 和 programming guide 中阅读它，还有 CUDA sample projects 显示如何编译和使用它。