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

Calling thrust function inside a CUDA Kernel __global___

我了解到较新版本的 CUDA 支持动态并行性，我可以在带有 thrust::device 参数的内核函数中调用像 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，这样我就可以执行 if 到 threadIdx；如果没有，块中的线程如何相互通信以表明对 thrust 的调用已经完成并且它们应该忽略它（这似乎有点虚构，因为没有系统的方法来确保用户代码）。总而言之，当我在内核中使用 thrust::device 参数调用推力函数时到底发生了什么？

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

  2. 如果您满足了 CUDA 动态并行 (CDP) 调用的所有要求（HW/SW 和编译设置），那么每个遇到推力算法调用的线程都会启动一个 CDP子内核执行推力算法（在这种情况下，CDP 子内核中的线程 do 协作）。如果不是，每个遇到推力算法调用的线程都会执行它，就好像您指定了 thrust::seq 而不是 thrust::device.

  3. 如果您希望在其他支持 CDP 的环境中避免使用 CDP activity，您可以改为指定 thrust::seq。

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

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

    或类似。

  5. 关于同步的问题before/after调用和普通的CUDA代码没什么区别。如果您需要块中的所有线程等待推力算法完成，请使用例如__syncthreads()，（以及 CDP 案例中的 cudaDeviceSynchronize()）。

信息 可能也很有趣。