CUDA 和线程块开销

Question

我实现了一个计算矩阵乘法的简单内核。这是代码：

__global__ void MatMultKernel(Mat A, Mat B, Mat C){

int i;
double val=0.0;
int ix=blockDim.x*blockIdx.x+threadIdx.x;  
int iy=blockDim.y*blockIdx.y+threadIdx.y;
if(ix<A.nrows && iy<B.nrows){
    for(i=0;i<A.nrows;i++)
        val+=A.value[iy*A.nrows+i]*B.value[i*B.nrows+ix];
    C.value[iy*C.nrows+ix]=val;
 }
}

我通过改变线程和块配置测量了这个内核的运行时间。

我发现如果线程分配在列块中（例如 dim3 (1,256,1)），执行时间总是更差，这是什么原因？

Answer 1

首先，请允许我指出您的代码仅适用于方阵，通常您应该在代码的所有三个位置使用 A.ncols 而不是 A.nrows。

性能差异是由于内存访问。您将所有三个矩阵都以行优先格式存储，代码执行以下操作：每个线程访问 A 的第 iy 行和 [=15= 的第 ix 列] 并计算它们的点积。请注意，warp 中的所有线程始终同时执行相同的指令，因此在代码的串行循环中，i 对于 warp 中的所有线程始终相同。在您的代码中，块形状很重要，因为：

如果块大小为(256, 1, 1)，则一个块中的每个线程具有相同的iy，但不同ix。让我们看一下 B 的访问模式：同一个 warp 中的所有线程总是访问 B 的同一行，因为在 B.value[i*B.nrows+ix] 中 i 是相同的并且ix 不同，所以负载可以是 coalesced.
如果块大小为 (1, 256, 1)，则情况发生转换，因此您可能希望合并来自 A 的负载。但事实并非如此，因为 iy 确定了行，并且两个相邻线程访问的值偏移了 A.ncols.

C 的访问模式与 B 相同，但重要性要低得多。对于 2D 块，情况介于两种 1D 情况之间。

如果您想进一步优化代码，可以使用 CUDA Programming Guide 中所示的共享内存。

CUDA 和线程块开销

CUDA and thread blocks overhead

cuda

gpgpu