在 CUDA 内核 __global___ 内调用推力函数

Question

我读到过有关 CUDA 的新版本支持动态并行性的信息，并且我可以像thrush::exclusive_scan在带有参数的内核函数内部一样调用推力函数thrust::device。

__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，这样我就可以做一个ifto threadIdx；如果没有，块中的线程如何相互通信，以确保对推力的调用已完成，并且它们应该忽略它（这似乎有点想象，因为没有系统的方法来确保用户的代码）。总结一下，当我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 代码没有什么不同。如果您需要块中的所有线程等待推力算法完成，请使用例如__syncthreads()（cudaDeviceSynchronize()在 CDP 情况下）。

这里的信息可能也令人感兴趣。