在1-D网格中计算warp id/lane id的最有效方法是什么？

Question

在CUDA中,每个线程都知道它在网格中的块索引和块内的线程索引.但似乎没有明确可用的两个重要值:

• 它的指数作为其经线内的一条车道(其"车道ID")
• 经线的索引,它是块内的一个通道(它的"warp id")

假设网格是一维(又名线性的,即blockDim.y和blockDim.z是1),可以明显地获得这些如下:

enum : unsigned { warp_size = 32 };
auto lane_id = threadIdx.x % warp_size;
auto warp_id = threadIdx.x / warp_size;

如果您不信任编译器来优化它,您可以将其重写为:

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;

这是最有效的事情吗？对于每个线程来说,计算它仍然需要很多浪费.

_{(受这个问题的启发.)}

Answer 1

天真的计算目前是最有效的.

注意:此答案已经过大量编辑.

尝试完全避免计算是非常诱人的 - 因为如果你深入了解这两个值似乎已经可用了.

您可以看到,nVIDIA GPU具有特殊寄存器,您的(已编译)代码可以读取这些寄存器以访问各种有用信息.有一个这样的登记册threadIdx.x; 另一个持有blockDim.x; 另一个 - 时钟滴答计数; 等等.C++作为一种语言显然没有这些暴露; 事实上,CUDA也没有.但是,编译CUDA代码的中间表示(名为PTX)确实暴露了这些特殊寄存器(因为PTX 1.3,即CUDA版本> = 2.1).

其中两个特殊寄存器是%warpid和%laneid.现在,CUDA支持使用asm关键字在CUDA代码中内联PTX代码- 就像它可以用于主机端代码直接发出CPU汇编指令一样.使用这种机制,可以使用这些特殊寄存器:

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

__forceinline__ __device__ unsigned warp_id()
{
    // this is not equal to threadIdx.x / 32
    unsigned ret; 
    asm volatile ("mov.u32 %0, %warpid;" : "=r"(ret));
    return ret;
}

......但这里有两个问题.

第一个问题 - 正如@Patwie建议的那样 - 是%warp_id不会给你你真正想要的东西 - 它不是网格环境中的扭曲索引,而是在物理SM(它可以容纳如此多的扭曲)的上下文中一次居住),那两个不一样.所以不要使用%warp_id.

至于%lane_id它确实给你正确的值,但是它具有误导性的非性能:即使它是一个"寄存器",它也不像寄存器文件中的常规寄存器,具有1个周期的访问延迟.它是一个特殊的寄存器,在实际硬件中使用S2R指令检索,该指令可能表现出长延迟.

底线:只需从线程ID计算warp ID和线程ID.我们无法解决这个问题 - 现在.

Answer 2

另一个答案是非常 危险的！自己计算lane-id 和warp-id。

#include <cuda.h>
#include <iostream>

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

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

__global__ void kernel() {
  const int actual_warpid = get_warp_id();
  const int actual_laneid = get_lane_id();
  const int expected_warpid = threadIdx.x / 32;
  const int expected_laneid = threadIdx.x % 32;
  if (expected_laneid == 0) {
    printf("[warp:] actual: %i  expected: %i\n", actual_warpid,
           expected_warpid);
    printf("[lane:] actual: %i  expected: %i\n", actual_laneid,
           expected_laneid);
  }
}

int main(int argc, char const *argv[]) {
  dim3 grid(8, 7, 1);
  dim3 block(4 * 32, 1);

  kernel<<<grid, block>>>();
  cudaDeviceSynchronize();
  return 0;
}

这给出了类似的东西

[warp:] actual: 4  expected: 3
[warp:] actual: 10  expected: 0
[warp:] actual: 1  expected: 1
[warp:] actual: 12  expected: 1
[warp:] actual: 4  expected: 3
[warp:] actual: 0  expected: 0
[warp:] actual: 13  expected: 2
[warp:] actual: 12  expected: 1
[warp:] actual: 6  expected: 1
[warp:] actual: 6  expected: 1
[warp:] actual: 13  expected: 2
[warp:] actual: 10  expected: 0
[warp:] actual: 1  expected: 1
...
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0
[lane:] actual: 0  expected: 0

另请参阅 PTX 文档

一个预定义的只读特殊寄存器，用于返回线程的扭曲标识符。扭曲标识符在 CTA 内提供唯一的扭曲编号，但不在网格内的跨 CTA 之间提供唯一的扭曲编号。对于单个 warp 中的所有线程，warp 标识符将相同。

请注意，%warpid 是易失性的，并返回读取时线程的位置，但其值可能在执行期间发生变化，例如，由于抢占后线程的重新调度。

因此，它是调度程序的 warp-id，但不保证它与虚拟 warp-id 匹配（从 0 开始计数）。

该文档清楚地表明了这一点：

因此，如果内核代码中需要这样的值，则应使用 %ctaid 和 %tid 来计算虚拟扭曲索引；%warpid 主要用于启用分析和诊断代码来采样和记录信息，例如工作场所映射和负载分布。

如果你认为，好吧，让我们使用 CUB 来实现：这甚至会影响cub::WarpId()

返回调用线程的 warp ID。Warp ID 保证在 warp 之间是唯一的，但可能不对应于线程块内从零开始的排名。

编辑：使用%laneid似乎是安全的。