循环如何以块、扭曲和线程的形式映射到 GPU 中？答案

【问题标题】：How a loop is mapped into the GPU, in blocks, warps and threads?循环如何以块、扭曲和线程的形式映射到 GPU 中？
【发布时间】：2021-04-04 16:53:21
【问题描述】：

我需要澄清如何使用 OpenACC 在设备上映射循环。我也不确定块、扭曲和线程的作用。

如果我有这样的循环：

#pragma acc parallel loop
for(i=0; i<1024; i++){
  vector[i] += 1;
}

而且我的 GPU 支持“每个块的最大线程数 = 1024”。 如何将循环并行化为块？ 我的第一个想法是单个块足以处理操作，因为向量有 1024 个元素。在这种情况下，我认为该块由 1024 个线程组成，每个线程对应于具有不同索引 i 的操作vector[i] += 1;。

我对线程的理解是否正确？

我会有 32 条 32 条线的经线。他们是如何被处决的？可以同时运行吗？

【问题讨论】：

标签： multithreading gpu openacc

【解决方案1】：

OpenACC 本身并没有规定到目标设备的 gang 和 vector 映射，这是由编译器实现完成的。虽然您没有具体说明，但我的回答假设您使用的是 NVIDA HPC 编译器（又名 PGI），而不是 GNU。

以 NVIDIA GPU 为目标时，“帮派”映射到 CUDA 块，“向量”映射到线程块的 x 维度。由于您没有指定向量长度，因此编译器很可能使用 8 个组（块），每个组有 128 个向量（线程）。您可以通过在编译期间添加标志“-Minfo=accel”以查看编译器反馈消息并设置环境变量“NV_ACC_TIME=1”以在运行后获取简单配置文件来验证这一点。

如果要强制编译器在每个块中使用 1024 个线程，请将子句“vector_length(1024)”添加到并行循环指令中。在这种情况下，性能在使用 1x1024 或 8x128 计划之间并不重要，但请尝试一下。

我会有 32 条 32 条线的经线。他们是如何被处决的？都可以它们中的哪些同时运行？

是的，您总共有 32 个经纱，每个经纱有 32 个线程。 Warp 在 SIMT 模式下执行，单指令多线程，这意味着所有线程都在同一时间执行相同的指令，只是在不同的数据上。

在 NVIDIA 设备上，您拥有流式多处理器 (SM)，其中最多可以同时执行 2048 个线程或 64 个 warp。由于您只使用 1024 个线程，是的，这些线程将同时运行。

请注意，SM 数量因设备而异，但例如，V100 有 80 个 SM，因此仅使用一个 SM 的一半会严重影响设备的利用率。

【讨论】：