3

我目前正在学习 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 线程似乎不能从处理相邻数据中受益，这会帮助它获得更多的缓存命中？

4

1 回答 1

5

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

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

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

于 2013-05-31T20:39:58.490 回答