CUDA矩阵加法计时,按行与 按列

Question

我目前正在学习CUDA,正在进行一些练习.其中之一是实现以3种不同方式添加矩阵的内核:每个元素1个线程,每行1个线程,每列1个线程.矩阵是方形的,并且实现为1D向量,我只需将其索引

A[N*row + col]

直觉上,我预计第一个选项由于线程开销而变得最慢,第二个选项是单个线程处理相邻数据以来最快的选项.

在CPU上,使用8000 x 8000的密集矩阵,我得到:

Adding on CPU - Adding down columns
Compute Time Taken: 2.21e+00 s
Adding on CPU - Adding across rows
Compute Time Taken: 2.52e-01 s

因此,由于更多的缓存命中,所以大约加速了一个数量级.在具有相同矩阵的GPU上,我得到:

Adding one element per thread 
Compute Time Taken: 7.42e-05 s
Adding one row per thread 
Compute Time Taken: 2.52e-05 s
Adding one column per thread 
Compute Time Taken: 1.57e-05 s

这对我来说不直观.对于最后一种情况,30-40%的加速度在大约1000×1000矩阵之上是一致的.请注意,这些时序仅是内核执行,不包括主机和设备之间的数据传输.下面是我的两个内核进行比较.

__global__
void matAddKernel2(float* A, float* B, float* C, int N)
{
        int row = threadIdx.x + blockDim.x * blockIdx.x;
        if (row < N)
        {
                int j;
                for (j = 0; j < N; j++)
                {
                        C[N*row + j] = A[N*row + j] + B[N*row + j];
                }
        }
}



__global__
void matAddKernel3(float* A, float* B, float* C, int N)
{
        int col = threadIdx.x + blockDim.x * blockIdx.x;
        int j;

        if (col < N)
        {
                for (j = 0; j < N; j++)
                {
                        C[col + N*j] = A[col + N*j] + B[col + N*j];
                }
        }
}

我的问题是,为什么GPU线程似乎不会从相邻数据的工作中受益,这有助于它获得更多的缓存命中？

Answer 1

GPU线程确实受益于处理相邻数据,你缺少的是GPU线程不是像CPU线程那样的独立线程,它们在一个名为warp的组中工作.warp将32个线程组合在一起,并以类似于单个CPU线程的方式工作,执行宽度为32的SIMD指令.

所以实际上每列使用一个线程的代码是最有效的,因为warp中的相邻线程正在从内存访问相邻的数据位置,这是访问全局内存的最有效方式.

您将在CUDA文档中找到详细信息.