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

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% 的加载/存储效率，因为事务中涉及的一半元素没有被使用（成为浪费的带宽）。我们如何继续计算更大步幅的加载/存储效率？提前致谢！

Answer 1

一般来说：

load efficiency = requested loads / effective loads

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

完全合并的访问效率为 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%。