Матричный векторный продукт CUDA performance

Question

Я нашел код о векторном продукте cuda matrix в предыдущей теме: 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 должно быть реализовано по-другому, чтобы достичь максимальной производительности.