矩阵向量积 CUDA 性能

Question

我在上一个主题中找到了一些关于 cuda 矩阵向量积的代码： 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

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

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