了解经线平行性（费米）

Question

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

__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 个，还是在其他线程等待操作数时执行指令？

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

Answer 1

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

有关全局、本地和共享内存回放的其他信息，请参阅 Compute Capability 2.x

上的 CUDA 编程部分

了解经线平行性（费米）

Understanding Warp Parallelism (Fermi)

cuda

gpgpu