在CUDA中,每个线程都知道它在网格中的块索引和块内的线程索引.但似乎没有明确可用的两个重要值:
假设网格是一维(又名线性的,即blockDim.y和blockDim.z是1),可以明显地获得这些如下:
enum : unsigned { warp_size = 32 };
auto lane_id = threadIdx.x % warp_size;
auto warp_id = threadIdx.x / warp_size;
Run Code Online (Sandbox Code Playgroud)
如果您不信任编译器来优化它,您可以将其重写为:
enum : unsigned { warp_size = 32, log_warp_size = 5 };
auto lane_id = threadIdx.x & (warp_size - 1);
auto warp_id = threadIdx.x >> log_warp_size;
Run Code Online (Sandbox Code Playgroud)
这是最有效的事情吗?对于每个线程来说,计算它仍然需要很多浪费.
(受这个问题的启发.)