行列ベクトル積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

ここで共有メモリはキャッシュとして機能します。ベクトルの成分は複数回読み取られますが、行列の成分は計算中に1回だけ読み取られます。そのため、コードはベクトルのみをキャッシュしますが、マトリックスはキャッシュしません。

マトリックスを読み取ると、スレッドがマトリックスの列に沿って編成されるため、カラムの主要マトリックスが高速になります。コルメジャーはcoalesced global memory accessを保証します。行列が行優先である場合、最大限のパフォーマンスを達成するために、CUDAカーネルを別の方法で実装する必要があります。