跨步内存访问的有效内存带宽

Question

假设我有一个执行跨步内存访问的内核，如下所示：

__global__ void strideExample (float *outputData, float *inputData, int stride=2) 
{
        int index = (blockIdx.x * blockDim.x + threadIdx.x) * stride;
        outputData[index] = inputData[index]; 
}

据我所知，步长为 2 的访问将导致 50% load/store 效率，因为交易中涉及的一半元素未被使用（成为带宽浪费）。我们如何继续计算更大步幅的 load/store 效率？提前致谢！

Answer 1

总的来说：

load efficiency = requested loads / effective loads

其中requested loads是软件请求读取的字节数，effective loads是硬件实际[=28]读取的字节数=]必须阅读。相同的公式适用于商店。

完美合并的访问效率为 1。

您的代码正好请求 (blockIdx.x * blockDim.x + threadIdx.x) * sizeof(float) 字节。假设 outputData 正确对齐（cudaMalloc 返回的指针也是如此），硬件将必须读取 (blockIdx.x * blockDim.x + threadIdx.x) * sizeof(float) * stride 字节，四舍五入到事务大小（SM/L1 为 128 字节, L1/L2).

32 字节

假设您的区块大小足够大，交易大小的四舍五入可以忽略不计，您可以将等式简化为 1 / stride，在这种情况下，负载效率约为 16.7%。

跨步内存访问的有效内存带宽

effective memory bandwidth from strided memory access

cuda

gpu

stride

memory-bandwidth