c++ - 矩阵 vector 乘积 CUDA 性能

Question

我在上一个主题中找到了一些关于 cuda 矩阵 vector 积的代码: Matrix-vector multiplication in CUDA: benchmarking & performance我首先想知道为什么作者没有为 dA(矩阵)使用共享内存？

然后，为什么列主排序比行主排序快？

代码如下:

    template<typename T>
__global__ void matvec_kernel(const T * __restrict__ dA, const T * __restrict__ dx, T * __restrict__ dy, const unsigned int nRows, const unsigned int nCols)
{
    const unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;

    __shared__ T x_shared[BLOCK_SIZE];

    T y_val = 0.0;

    #pragma unroll
    for (unsigned int m = 0; m < ((nCols + BLOCK_SIZE - 1)/ BLOCK_SIZE); ++m)
    {
        if ((m * BLOCK_SIZE + threadIdx.x) <  nCols) x_shared[threadIdx.x] = dx[threadIdx.x + m * BLOCK_SIZE];
        else                                         x_shared[threadIdx.x] = 0.f;
        __syncthreads();

        #pragma unroll
        for (unsigned int e = 0; e < BLOCK_SIZE; ++e) {
            // --- Column-major ordering - faster
            y_val += dA[tid + (e + BLOCK_SIZE * m) * nRows] * x_shared[e];
            // --- Row-major ordering - slower
            //y_val += dA[tid * nCols + (e + BLOCK_SIZE * m)] * x_shared[e];
        }

        __syncthreads();
    }

    if (tid < nRows) dy[tid] = y_val;

我在这两个问题上思考了 1 天，这就是我来这里的原因。

非常感谢!

Answer 1

这里的共享内存起到缓存的作用。 vector 的分量会被多次读取，而矩阵的分量在计算过程中只会被读取一次。这就是代码只缓存 vector 而不缓存矩阵的原因。

列优先矩阵更快，因为在读取矩阵时，线程是沿着矩阵列组织的。 Col-major 因此确保了 coalesced global memory access .如果矩阵是行优先的，则 CUDA 内核应该以不同的方式实现以实现最佳性能。