如何计算正在启动的 CUDA 线程数？

Posted 2023-02-16

【中文标题】如何计算正在启动的 CUDA 线程数？

【英文标题】：How do i calculate the number of CUDA threads being launched?

【发布时间】：2015-07-25 14:20:49

【问题描述】：

我有一张 CUDA 卡： Cuda Compute 能力 (3.5) 如果我有一个诸如 >> 之类的调用，内核中发生的迭代次数是多少？我以为是（2000 * 512），但测试并不能证明这一点？我还想确认我计算变量的方式是正确的。

情况是，在内核中，我根据线程号递增传递的全局内存号：

  int thr = blockDim.x * blockIdx.x + threadIdx.x;
  worknumber = globalnumber + thr;

所以，当我返回 CPU 时，我想知道有多少增量，以便我可以跟踪，这样当我调用内核 GPU 来处理我的下一组数字时，我不会重复或跳过数字.

编辑：

__global__ void allin(uint64_t *lkey, const unsigned char *d_patfile)


    uint64_t kkey;
    int tmp;
    int thr = blockDim.x * blockIdx.x + threadIdx.x;
    kkey = *lkey + thr;

if (thr > tmp) 
    tmp = thr;
    printf("%u \n", thr);
    

【问题讨论】：

您将不得不解释迭代的含义。 CUDA 在设计上不是以任何方式迭代的

当我说迭代时，我的意思是内核中的代码运行了多少次。线程总数是多少？当我说线程时，我不完全确定我的命名约定是否正确。迭代的定义：“将计算机指令序列重复指定次数或直到满足条件”。这不正是发生的事情吗？

你能把代码 sn-p 扩展到你的内核中吗？我猜你看到了内存竞争的影响，但是没有看到代码就不可能说

从 c/c++ 的角度来看，您的内核没有意义。在if 语句中测试之前，您没有将tmp 初始化为任何值。我认为编译器会对此发出警告。 <<<2000,512>>> 创建的线程或“迭代”数确实是 2000*512。来自 cuda 内核的printf 有各种限制，因此使用它来验证是否启动了大量线程可能不起作用。

如果你想自己证明2000*512这个数字，那么创建一个单独的__device__全局变量，初始化为零，然后让每个线程做atomicAdd(&var, 1);之后，复制变量返回主机代码并打印出来。

【参考方案1】：

如果您使用配置<<<X,Y>>> 启动内核，并且您没有违反任何 CUDA 使用规则，那么启动的线程数实际上将是 X*Y（或者如果我们对它进行适当的修改，正在谈论 2 或 3 维线程块和/或网格，即X.x*X.y*X.z*Y.x*Y.y*Y.z)。

来自 CUDA 内核的printf 有各种limitations。因此，从 CUDA 内核生成大量 printf 输出通常是不明智的，并且对于验证在大型网格中启动的线程数可能没有用。

如果您想跟踪实际启动的线程数，您可以使用全局变量并让每个线程自动更新它。像这样的：

$ cat t848.cu
#include <stdio.h>

__device__ unsigned long long totThr = 0;

__global__ void mykernel()

  atomicAdd(&totThr, 1);


int main()

  mykernel<<<2000,512>>>();
  unsigned long long total;
  cudaMemcpyFromSymbol(&total, totThr, sizeof(unsigned long long));
  printf("Total threads counted: %lu\n", total);

$ nvcc -o t848 t848.cu
$ cuda-memcheck ./t848
========= CUDA-MEMCHECK
Total threads counted: 1024000
========= ERROR SUMMARY: 0 errors
$

请注意，原子操作可能相对较慢。出于性能原因，我不建议经常使用这样的代码。但是如果你想说服自己启动的线程数，它应该给出正确的答案。

【讨论】：

是的，我正在学习 GPU 编码。这仅用于调试和考虑。感谢您提供深思熟虑的答案。非常感谢！