CUDA 矩阵加法时序，按行与按行比较按栏目

我目前正在学习 CUDA，并正在做一些练习。其中之一是实现以 3 种不同方式添加矩阵的内核：每个元素 1 个线程、每行 1 个线程和每列 1 个线程。矩阵是方阵，并被实现为一维向量，我只需用以下命令对其进行索引

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 x 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 线程似乎没有从处理相邻数据中受益，这将有助于它获得更多缓存命中？

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

因此，实际上，每列使用一个线程的代码是最有效的，因为扭曲内的相邻线程正在访问内存中的相邻数据位置，这是访问全局内存的最有效方法。

您将在以下位置找到详细信息CUDA文档 http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#simt-architecture.