0

我想thrust::scatter通过在设备内核中调用它来实现异步（我也可以通过在另一个主机线程中调用它来实现）。thrust::cuda::par.on(stream)是不能从设备内核调用的主机函数。以下代码在图灵架构上使用 CUDA 10.1 进行了尝试。


__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);
}

我知道推力在从设备调用时使用动态并行性来启动其内核，但是我找不到指定流的方法。

4

1 回答 1

1

以下代码在 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
$

对于任何使用 CUDA 动态并行的代码，-arch=sm_35（或类似的）和是必要的（但并非在所有情况下都足够）编译开关。-rdc=true例如，如果您省略-rdc=trueswitch，则会收到类似于您描述的错误：

$ 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 版本或通过指定正确的命令行或两者兼而有之来消除。

于 2019-09-26T18:27:25.847 回答