c++ - 在 CUDA 内核 __global___ 内调用推力函数

Question

我已经了解到新版本的 CUDA 支持动态并行性，并且我可以使用 thrust::在内核函数中调用诸如 thrush::exclusive_scan 之类的推力函数。设备参数。

__global__ void kernel(int* inarray, int n, int *result) {
  extern __shared__ int s[];
  int t = threadIdx.x;

  s[t] = inarray[t];
  __syncthreads();

  thrust::exclusive_scan(thrust::device, s, n, result);
  __syncthreads();
}

int main() {
  // prep work

  kernel<<<1, n, n * sizeof(int)>>>(inarray, n, result);
}

我感到困惑的是:

1. 当在内核中调用推力函数时，每个线程是否调用该函数一次，并且都对数据进行动态并行处理？
2. 如果这样做，我只需要一个线程来调用 thrust，这样我就可以对 threadIdx 执行一个 if 操作；如果没有， block 中的线程如何相互通信，以确保对推力的调用已完成，并且它们应该忽略它(这似乎有点想象，因为没有系统的方法来确保用户的代码)。总结一下，当我在内核中使用 thrust::device 参数调用推力函数时到底发生了什么？

Answer 1

1. 内核中执行推力算法的每个线程都将执行算法的单独拷贝。内核中的线程不会在单个算法调用上进行合作。

2. 如果您已满足 CUDA 动态并行 (CDP) 调用的所有要求(硬件/软件和编译设置)，则遇到推力算法调用的每个线程将启动 CDP 子内核来执行推力算法(在这种情况下，CDP 子内核中的线程执行合作)。如果没有，遇到推力算法调用的每个线程都会执行它，就像您指定了 thrust::seq 而不是 thrust::device 一样。

   <

3. 如果您希望避免在支持 CDP 的环境中进行 CDP 事件，则可以指定 thrust::seq。

4. 例如，如果您打算只执行推力算法的一个拷贝，则需要在内核代码中确保只有一个线程调用它，例如:

   if (!threadIdx.x) thrust::exclusive_scan(...  
   

   或类似内容。

5. 调用前后的同步问题与普通 CUDA 代码没有什么不同。如果您需要 block 中的所有线程等待推力算法完成，请使用例如__syncthreads()(在 CDP 情况下还有 cudaDeviceSynchronize())。

信息here可能也会引起兴趣。