从跨行存储器访问有效的存储器带宽

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%。