c++ - CUDA channel ID 与基于 threadIdx.x 的计算

Question

通过 cub::LaneId() 或类似以下的函数来解释是最简单的:

inline __device__ unsigned get_lane_id() {
  unsigned ret;
  asm volatile("mov.u32 %0, %laneid;" : "=r"(ret));
  return ret;
}

与将 channel ID 计算为 threadIdx.x & 31 相比。

这两种方法在一维网格中产生相同的值吗？

__ballot_sync() 文档在其 mask 参数中谈到了 channel ID，据我了解，它返回每个 channel ID 设置的位。那么下面的断言永远不会失败吗？

int nWarps = /*...*/;
bool condition = /*...*/;
if(threadIdx.x < nWarps) {
  assert(__activemask() == ((1u<<nWarps)-1));
  uint32_t res = __ballot_sync(__activemask(), condition);
  assert(bool(res & (1<<threadIdx.x)) == condition);
}

Answer 1

来自 PTX ISA 文档:https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-laneid

%laneid A predefined, read-only special register that returns the thread's lane within the warp. The lane identifier ranges from zero to WARP_SZ-1.

该寄存器将始终包含正确的值，而 threadIdx.x & 31假设扭曲大小为 32。但是，对于迄今为止的所有 GPU 代，扭曲大小均为 32，因此对于旧架构和当前架构，计算 channel 将相同。但是，不能保证情况总是如此。

关于你关于断言的问题。使用独立线程调度，不能保证 warp 中的所有线程都会执行 __activemask()同时。我认为断言可能会失败。

引用编程指南:https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#independent-thread-scheduling-7-x

Note that threads within a warp can diverge even within a single code path. As a result, __activemask() and __ballot(1) may return only a subset of the threads on the current code path.