什么是在一维网格中计算warp id / lane id的最有效方式？

Question

什么是在一维网格中计算warp id / lane id的最有效方式？

14

在CUDA中，每个线程都知道它在网格中的块索引和块内线程索引。但两个重要的值似乎对它并不明确可用：

- 它作为warp中的lane的索引（其“lane id”） - 它作为warp所在块中的lane的索引（其“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;

那是最有效的做法吗？每个线程都必须计算这个似乎仍然会产生很多浪费。

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

- einpoklum

CUDA编程指南在其示例中多次使用threadIdx.x＆0x1f来获取lane id，这相当于threadIdx.x％32。 - Azmisov

2个回答

6

另一个答案非常危险！自己计算车道标识和线程块标识。

#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文档。

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

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

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

文档已经表明:

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

如果您认为可以使用CUB解决此问题：这甚至会影响 cub::WarpId()

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

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

- Patwie

根据之前的评论和您的回答，我想删除我的答案。所以，您介意编辑您的答案，警告不要使用%warp_id寄存器吗？此外，您对使用%lane_id有任何疑虑吗？ - einpoklum

随意修改我的回答。对于%lane_id，我编辑了代码。我只是在我的代码中进行错误调试，将您的回答视为终极智慧，并尝试在自己的代码行中找到错误。 - Patwie

你的代码现在有问题：1. 你正在将%lane_id分配给lane_id。2. 你只检查了车道0的车道ID。 - einpoklum

我修改了代码。似乎 lane_id 没有受到影响。 - Patwie

Patwie：所以车道ID没问题；但你只从0号车道打印它。 - einpoklum

让我们在聊天中继续这个讨论。 - Patwie

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- einpoklum · Accepted Answer

目前，朴素计算是最有效的。

注意：此答案已经进行了大幅编辑。

在计算这两个值时，很容易尝试避免计算——因为如果你深入了解，这些值似乎已经可以利用。你知道，nVIDIA GPU 有一些特殊寄存器，你的（编译后）代码可以读取这些寄存器以访问各种有用的信息。其中一个寄存器保存 threadIdx.x，另一个保存 blockDim.x，另一个则是时钟节拍计数。C++ 并没有暴露这些特殊寄存器，显然，实际上，CUDA 也没有。然而，CUDA 编译成的中间表示形式，称为 PTX，确实暴露了这些特殊寄存器（自 CUDA 2.1 起，即 PTX 1.3 版本以上）。

其中两个特殊寄存器是 %warpid 和 %laneid。现在，CUDA 支持在 CUDA 代码内联 PTX 代码使用 asm关键字——就像它可以用于主机代码以直接发出 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指令检索的，可以展示较长的延迟时间。由于你几乎肯定已经在一个寄存器中拥有threadIdx.x的值，所以对这个值应用位掩码比检索%lane_id更快。

底线：只需从线程ID计算warp ID和lane ID。我们无法绕过这一点 - 至少目前是这样。