Matrix Vektor Produkt CUDA Leistungen

Question

Ich habe etwas Code über eine Cuda Matrix Vektor Produkt in einem früheren Thema gefunden: Matrix-vector multiplication in CUDA: benchmarking & performance Ich fragte mich zuerst, warum der Autor nicht gemeinsam genutzten Speicher für dA (die Matrix) verwendet?

Und dann, warum die Spalte Hauptbestellung ist schneller als Reihe Hauptbestellung? Hier

ist der Code:

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; 

}

Ich bin jetzt für 1 Tag auf diese beiden Fragen zu denken, und das ist, warum ich hier bin.

Vielen Dank!

Answer 1

Gemeinsamer Speicher fungiert hier als Cache. Die Komponenten des Vektors werden mehrfach gelesen, aber die Komponenten der Matrix werden nur einmal während der Berechnung gelesen. Deshalb speichert der Code nur den Vektor, aber nicht die Matrix.

Die Spaltenhauptmatrix ist schneller, weil beim Lesen der Matrix die Threads entlang der Matrixspalten organisiert sind. Col-Dur sorgt somit für die coalesced global memory access. Wenn die Matrix eine Zeilenhauptreihe hat, sollte der CUDA-Kernel auf andere Weise implementiert werden, um eine maximale Leistung zu erzielen.