Ошибка cuda :: cub при вызове функции __host__ из функции __device__ не допускается

Question

Я использую 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 {}", но не может работать в функции.

Answer 1

Чтобы использовать функцию куба для всего устройства из кода устройства, необходимо создать свой проект для поддержки динамического параллелизма CUDA.В документации cub это указано здесь:

Замечания по использованию Динамический параллелизм.Методы DeviceReduce можно вызывать из кода ядра на устройствах, в которых поддерживается динамический параллелизм CUDA.

Например, вы можете скомпилировать код, который вы показали с помощью:

$ 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 Dynamic Parallelism (CDP),Выше приведена только команда компиляции и пропускается шаг ссылки.Здесь есть много вопросов о теге cuda, которые обсуждают CDP, вы можете прочитать об этом в двух статьях блога и руководстве по программированию , а также есть примеры проектов CUDA показывает, как его скомпилировать и использовать.