凭经验确定一个warp中有多少线程

Question

是否可以编写一个 CUDA 内核来显示一个 warp 中有多少线程，而不使用任何与 warp 相关的 CUDA 设备函数并且不使用基准测试？如果有，怎么做？

Answer 1

由于您表示使用原子的解决方案会很有趣，因此我将此作为我认为可以给出答案的东西提出来，但我不确定它是否一定是您正在寻找的答案。我承认它在本质上有点统计。我提供这个仅仅是因为我发现这个问题很有趣。我不认为这是“正确”的答案，我怀疑聪明的人会想出一个“更好”的答案。不过，这可能会提供一些想法。

为了避免使用任何显式引用扭曲的东西，我认为有必要关注“隐式”扭曲同步行为。我最初走的是一条思考如何使用 if-then-else 构造的路径（它具有一些扭曲同步的含义），但为此苦苦挣扎并提出了这种方法：

#include <stdio.h>
#define LOOPS 100000

__device__ volatile int test2 = 0;
__device__ int test3 = 32767;

__global__ void kernel(){

  for (int i = 0; i < LOOPS; i++){
    unsigned long time = clock64();
//    while (clock64() < (time + (threadIdx.x * 1000)));
    int start = test2;
    atomicAdd((int *)&test2, 1);
    int end = test2;
    int diff = end - start;
    atomicMin(&test3, diff);
    }
}

int main() {

   kernel<<<1, 1024>>>();
   int result;
   cudaMemcpyFromSymbol(&result, test3, sizeof(int));
   printf("result = %d threads\n", result);
   return 0;
}

我编译：

nvcc -O3 -arch=sm_20 -o t331 t331.cu

我将其称为“统计”，因为它需要大量迭代 (LOOPS) 才能产生正确的估计 (32)。随着迭代次数的减少，“估计”增加。

我们可以通过取消注释内核中注释掉的行来应用额外的 warp 同步杠杆。对于我的测试用例*，该行未注释，即使LOOPS = 1，估计也是正确的

*我的测试用例是 CUDA 5、Quadro5000、RHEL 5.5

Answer 2

这里有几个简单的解决方案。还有其他使用 warp 同步编程的解决方案；但是，许多解决方案并不适用于所有设备。

解决方案 1：以每个块的最大线程数启动一个或多个块，读取特殊寄存器 %smid 和 %warpid，以及 blockIdx 并将值写入内存。按三个变量对数据进行分组以查找经纱大小。如果您将启动限制为单个块，这会更容易，那么您只需要 %warpid。

解决方案 2：以每个块的最大线程数启动一个块并读取特殊寄存器 %clock。这需要以下假设，这些假设可以在 CC 1.0-3.5 设备上证明是正确的：

• %clock 被定义为一个无符号的 32 位只读周期计数器，它以静默方式包装并在每个发布周期更新
• warp 中的所有线程都读取相同的 %clock 值
• 由于同一 SM 上的 warp 启动延迟和指令获取 warp，但不同的 warp 调度程序无法在同一周期发出 warp 的第一条指令

在 CC1.0 - 3.5 设备上具有相同时钟时间的块中所有线程（将来可能会更改）将具有相同的时钟时间。

解决方案 3：使用 Nsight VSE 或 cuda-gdb 调试器。翘曲状态视图显示了足够的信息来确定翘曲大小。也可以单步查看每个线程PC地址的变化。

解决方案 4：使用 Nsight VSE、Visual Profiler、nvprof 等。启动 1 个块的内核，每次启动时增加线程数。确定导致 warps_launched 的线程数从 1 变为 2 的时间。