了解内存重播和进行中的请求

Question

我正在尝试了解矩阵转置如何能够更快地从列与行中天真地读取。（示例来自 Professional CUDA C Programming）矩阵按行存储在内存中，即 (0,1),(0,2),(0,3)...(1,1),(1,2)

__global__ void transposeNaiveCol(float *out, float *in, const int nx, const int ny) {
    unsigned int ix = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int iy = blockDim.y * blockIdx.y + threadIdx.y;

    if (ix < nx && iy < ny) {
           out[iy*nx + ix] = in[ix*ny + iy]; // 
           // out[ix*ny + iy] = in[iy*nx + ix]; // for by row
    }
}

这是我不明白的地方：transposeNaiveCol() 的负载吞吐量是 642.33 GB/s，tranposeNaiveRow() 的负载吞吐量是 129.05 GB/s。作者说：

The results show that the highest load throughput is obtained with cached, strided reads. In the case of cached reads, each memory request is serviced with a 128-byte cache line. Reading data by columns causes each memory request in a warp to replay 32 times (because the stride is 2048 data elements), resulting in good latency hiding from many in-flight global memory reads and then excellent L1 cache hit ratios once bytes are pre-fetched into L1 cache.

我的问题：我认为 aligned/coalesced 读取是理想的，但在这里似乎跨步读取可以提高性能。

为什么读取缓存行会导致性能下降这种情况？
总的来说重放不是一件坏事吗？它在这里提到它导致 "good latency hiding"。

Answer 1

有效负载吞吐量并不是决定内核性能的唯一指标！具有完美合并负载的内核将始终具有比等效的非合并内核更低的有效负载吞吐量，但是单独对其执行没有任何说明时间：最后，真正重要的一个指标是内核完成所需的挂钟时间，作者没有提及。

也就是说，内核通常分为两类：

计算绑定内核，其性能可以通过尝试隐藏指令延迟来提高：保持流水线满（最大化 ILP）。
I/O 绑定内核，可以通过尝试隐藏内存延迟来提高其性能：保持数据传输（最大化带宽）。

矩阵转置的计算强度非常低，因此受到 I/O 限制，因此要获得更好的性能，您应该尝试增加带宽使用量。

为什么列转置在最大化带宽使用方面更好？

在行转置的情况下，读取被合并：每个 warp 提供一个 128 字节的事务，即每个线程 4 个字节。这 128 个字节被放入缓存但从未被重复使用，因此缓存在这种情况下实际上没有用。

在列转置的情况下，读取不会合并：每个 warp 得到 32 个 128 字节的事务，所有这些都将进入 L1 并将在接下来的 31 次重放中重复使用（假设它们没有被踢出缓存）。对于非常高的有效负载吞吐量和最大缓存使用率，这是非常低的负载效率。

您当然可以通过简单地为每个线程请求更多数据（例如通过每个线程加载 32 float 或 8 float4）或使用 CUDA 的预取在行转置中获得相同的效果能力。

了解内存重播和进行中的请求

Understanding Memory Replays and In-Flight Requests

caching

cuda