CUDA lane ID vs threadadix.基于X的计算

Question

最容易通过cub::LaneId()或下面的函数来解释:

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

与计算车道ID为threadIdx.x & 31相比

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

__ballot_sync()文档在其mask参数中谈到车道ID，据我所知，它返回每个车道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一个预定义的，只读的特殊寄存器，返回线程在经线内的通道。通道标识符的取值范围是0到WARP_SZ-1。

这个寄存器将始终包含正确的值，而threadIdx.x & 31假设warp大小为32。然而，对于迄今为止的所有GPU代，warpsize都是32，因此对于旧的和当前的架构，计算通道将是相同的。但是，不能保证总是如此。

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

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

注意，即使在单个代码路径中，warp中的线程也可以发散。因此，__activemask()和__ballot(1)可能只返回当前代码路径上线程的一个子集。

Answer 2

这两种方法在1D网格中产生相同的值吗?

是的(而CUDA的翘曲大小是32)。参见这个问题:

在一维网格中计算翘曲id/车道id的最有效方法是什么?

但是我想这样写:

enum { warp_size = 32 };
// ...
inline unsigned lane_id() {
constexpr const auto lane_id_mask = warp_size - 1;
return threadIdx.x & lane_id_mask;
}

，如果您想要更加专业，您总是可以静态断言，以确保翘曲大小是2的幂:-P

那么下列断言永远不会失败吗?

代码看起来很奇怪。为什么要根据线程ID或经线数左移?我不明白为什么这不会失败。