Как вызвать функцию Thrust в потоке из ядра? - PullRequest
Новая
       29

Как вызвать функцию Thrust в потоке из ядра?

0 голосов
/ 24 сентября 2019

Я хочу сделать thrust::scatter асинхронным, вызвав его в ядре устройства (я также мог бы сделать это, вызвав его в другом потоке хоста).thrust::cuda::par.on(stream) - это функция хоста, которую нельзя вызвать из ядра устройства.Следующий код был опробован в CUDA 10.1 на архитектуре Turing. 


__global__ void async_scatter_kernel(float* first,
    float* last,
    int* map,
    float* output)
{
    cudaStream_t stream;
    cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
    thrust::scatter(thrust::cuda::par.on(stream), first, last, map, output);
    cudaDeviceSynchronize();
    cudaStreamDestroy(stream);
}

Я знаю, что thrust использует динамический параллелизм для запуска своих ядер при вызове с устройства, однако я не смог найти способ указать поток.

1 Ответ

0 голосов
/ 26 сентября 2019

Следующий код корректно компилируется для меня на CUDA 10.1.243: 

$ cat t1518.cu
#include <thrust/scatter.h>
#include <thrust/execution_policy.h>

__global__ void async_scatter_kernel(float* first,
    float* last,
    int* map,
    float* output)
{
    cudaStream_t stream;
    cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);
    thrust::scatter(thrust::cuda::par.on(stream), first, last, map, output);
    cudaDeviceSynchronize();
    cudaStreamDestroy(stream);
}

int main(){

  float *first = NULL;
  float *last = NULL;
  float *output = NULL;
  int *map = NULL;
  async_scatter_kernel<<<1,1>>>(first, last, map, output);
  cudaDeviceSynchronize();
}
$ nvcc -arch=sm_35 -rdc=true t1518.cu -o t1518
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2019 NVIDIA Corporation
Built on Sun_Jul_28_19:07:16_PDT_2019
Cuda compilation tools, release 10.1, V10.1.243
$

Необходимы (но не во всех случаях достаточные) ключи компиляции для -arch=sm_35 (или аналогичных) и -rdc=true длялюбой код, который использует динамический параллелизм CUDA.Если вы опустите, например, ключ -rdc=true, вы получите ошибку, аналогичную описанной вами: 

$ nvcc -arch=sm_35 t1518.cu -o t1518
t1518.cu(11): error: calling a __host__ function("thrust::cuda_cub::par_t::on const") from a __global__ function("async_scatter_kernel") is not allowed

t1518.cu(11): error: identifier "thrust::cuda_cub::par_t::on const" is undefined in device code

2 errors detected in the compilation of "/tmp/tmpxft_00003a80_00000000-8_t1518.cpp1.ii".
$

Так, для примера, который вы показали здесь, ваша ошибка компиляции может быть устранена либообновление до последней версии CUDA или путем указания правильной командной строки, или обоих.

...