CUDA内核中全局内存负载事务的计算

Question

对于CUDA内核中全局内存负载事务的计算，我有一个疑问，因为配置值与我的期望值不匹配。以以下简单的矩阵乘法代码为例：

__global__ void matmul_kernel(float *A, float* B, float *C, int n)
{
    int i, j, k;
    float c;

    i = blockIdx.x;
    for(i=i; i < n; i += gridDim.x){
        j = threadIdx.x;
        c = 0.0;
        for(k = 0; k < n; k++)
            c += A[i*n + k]*B[k*n + j];
        C[i*n + j] = c;
    }
}

​

dim3 grid(1,1,1);
dim3 block(128,1,1);
n = 128;
matmul_kernel<<<grid, block>>>(A, B, C, n);

我以最简单的矩阵乘法为例。在上面的CUDA实现中，我将i循环迭代映射为块索引，j循环映射到每个线程块中的线程索引。线程块和网格都是一维的。

请不要将重点放在此实现的性能上。我知道这是没有效率的，因为我只是把它用于实验目的。

在这个实现中，因为我在每个块中分配了128个线程，所以j循环可以完全并行化。但是我只为i循环分配了一个块，所以它会循环n次数。下图显示了k=0时执行的状态。在这种状态下，128个线程访问A的第一个元素和B的128个第一行元素。我在Quadro K6000上执行这个CUDA代码，它使用开普勒40架构，我打开了L1缓存。由于对B的128个访问是合并的，所以加载的数量是128*4/128 = 4 (第一个128个是128个元素，第二个128个是L1缓存行大小(以字节为单位)，4个是浮点类型的字节)。对于A的128个访问，因为它们访问相同的元素，一个缓存行加载就足够了。所以全局负载的数量是4+1=5。但这只是k=0时加载的数量。k将被循环128次，i也将被循环128次，因此加载的总数为5*128*128=81920。然而，可分析的全局负载是131072。此值等于(4+4)*128*128。似乎A在k=0上的负载数是4而不是1。有人能解释为什么配置值与我的期望值不匹配吗？

​

​

Answer 1

只有一点是你错过的。全局内存访问仅用于翘曲中的线程(参见程序编制指南)。在你的例子中，有4个翘曲。每个元素都需要一个内存事务。