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

Question

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

如果我有这样的循环：

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

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

的操作 vector[i] += 1;

我对线程的理解正确吗？

我会有 32 个 warp，每个 32 个线程。他们是如何执行的？他们都可以同时运行吗？

Answer 1

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

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

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

I would have so 32 warps of 32 threads. How are they executed? Can all of them run simultaneously?

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

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

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

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

How a loop is mapped into the GPU, in blocks, warps and threads?

multithreading

gpu

openacc