理解翘曲平行度（费米）答案

【问题标题】：Understanding Warp Parallelism (Fermi)理解翘曲平行度（费米）
【发布时间】：2016-01-25 11:32:10
【问题描述】：

我有以下内核，其中每个线程（一维网格、一维块）只处理输入数组的一个元素。

__global__ void normalize_fft_result(double *u_device, int n0)
{
    //Use 1d data mapping;
    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    if (tid < n0)
        {
            //Normalize Result
            u_device[tid] = u_device[tid] / float(n0);
        }
}

我在 Fermi GPU 上运行它，我发现处理器将数据加载到 L1 缓存中的缓存线长度为 128 字节。我正在使用 8 个字节的双精度数，这意味着在一个事务中，一个 warp 中只有一半的线程具有可用的指令操作数（128/8=16）。这意味着为了获取另一半线程的数据，warp 需要另一个 128 B 事务。

warp 中的线程应该是并发执行的，那么在等待第二个事务期间究竟发生了什么？前 16 个线程是在等待后 16 个线程，还是在其他线程等待操作数时执行指令？

无论如何，这种数据等待不会产生不可避免的延迟吗？

【问题讨论】：

warp 调度程序将重播指令，直到所有线程完成内存加载或存储。在 CC2.x 设备上，通过发出前 16 个线程然后发出第二个线程来完成 64 位加载。如果存在额外的地址分歧（例如，每个线程读取单独的缓存行）和每个缓存未命中，则必须执行额外的重播。在 CC2.x 上，可以在 LD 指令的所有线程完成后发出来自 warp 的额外独立指令。
@GregSmith：如果您想将其添加为一个，那将是一个很好的答案。
@GregSmith 那么，这意味着warp调度程序将播放指令两次，因为需要2次加载才能获得必要的数据？确实还有一些等待经纱开始执行？
是的，warp 调度程序将至少重播指令两次。 Fermi 架构是一种延迟隐藏架构。为了隐藏延迟，您必须在每个 SM 上启动足够的 warp 以隐藏内存和执行依赖延迟。

标签： cuda gpgpu

【解决方案1】：

warp 调度程序将重播指令，直到所有线程完成内存加载或存储。在 CC2.x 设备上，通过发出前 16 个线程然后发出第二个线程来完成 64 位加载。如果存在额外的地址分歧（例如，每个线程读取单独的缓存行）和每个缓存未命中，将执行额外的重播。在 CC2.x 设备上，可以在加载或存储指令的所有线程完成后发出来自 warp 的额外独立指令。

有关全局、本地和共享内存重放的更多信息，请参阅Compute Capability 2.x 上的 CUDA 编程部分

【讨论】：