Warp中非活跃线程的内存操作机制探究

__device__ int arr[1]{5};

__global__ void kernel()
{
  uint32_t threadId = blockDim.x * blockIdx.x + threadIdx.x;

  uint32_t mask = __activemask();

  printf("id %u, mask %u, value %i\n", threadId, mask, arr[threadId]);
}

int main(int argc, char *[])
{
  kernel<<<1, 1>>>();
  gpuErrchk(cudaPeekAtLastError());

  gpuErrchk(cudaDeviceSynchronize());

  return 0;
}

上述代码仅启动一个线程，active mask显示仅该线程活跃，但受NVIDIA GPU特性影响，完整的一个Warp（32个线程）会被调度执行该内核。由此产生以下疑问：

• Warp中的非活跃线程如何执行内存操作？
• 其读写操作是否会实际执行？
• 是否会引发示例中数组大小为1时的内存越界问题？

解答 #

• 非活跃线程的内存操作处理：NVIDIA GPU的Warp执行单元会依据active mask掩码过滤非活跃线程的指令，非活跃线程的内存访问指令会被硬件直接忽略，不会进入实际的内存请求流程。
• 读写操作是否实际执行：不会。非活跃线程的所有计算、内存读写等操作都会被硬件抑制，这些线程仅在指令流水线中“空转”，不会产生任何实际的操作结果或内存交互行为。
• 是否会引发内存越界：不会。示例中只有线程ID为0的活跃线程会执行arr[threadId]的内存访问，其余31个非活跃线程的该访问指令被完全忽略，不会去读取超出数组arr大小的内存区域，因此不会触发内存越界错误。

本质上，Warp调度的是物理线程，但active mask决定了哪些线程的操作是有效的，非活跃线程不会参与任何有副作用的操作。

内容的提问来源于stack exchange，提问作者nikitablack