矩陣向量乘積CUDA的性能

我發現了一些關於前一個主題中的cuda矩陣向量乘積的代碼： Matrix-vector multiplication in CUDA: benchmarking & performance 我首先想知道爲什麼作者沒有爲dA（矩陣）使用共享內存？矩陣向量乘積CUDA的性能

然後，爲什麼列主要排序比行主要排序快？

下面是代碼：

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一天，這些兩個問題，這就是爲什麼我在這裏。

非常感謝！

來源

2016-07-28 Titouan Parcollet

此處的共享內存用作緩存。矢量的組成部分將被多次讀取，但矩陣的組成部分在計算過程中只能讀取一次。這就是代碼只緩存向量而不是矩陣的原因。

列主矩陣更快，因爲在讀取矩陣時，線程沿矩陣列組織。 Col-Major因此確保coalesced global memory access。如果矩陣是主要行，CUDA內核應以不同的方式實現以實現最高性能。

來源

2016-07-28 17:59:13 kangshiyin

完美答案！非常感謝！ –

因此，爲了達到行主要的最大性能，我需要使用threadIdx.y和nRows來代替threadIdx.x/nCols（在矩陣讀取階段）？ –

@TitouanParcollet編號它與上面的內核有很大不同。上面一個使用每個矩陣行一個*線程*，除非矩陣非常大，否則在性能方面實際上並不是最優的。對於行主矩陣，可以使用每個矩陣行中的一個*線程塊*，並使用並行約簡來計算行總和。 – kangshiyin

矩陣向量乘積CUDA的性能

回答

相關問題