Как реализовать умножение матрицы NxM CUDA?

Question

Я уже искал об этом на нескольких страницах, но я только нашел реализации для NxN-матриц.

Как я могу определить dimBlock и dimGrid, чтобы рассчитать это? Кроме того, как играть с:

int row = blockIdx.y*blockDim.y+threadIdx.y; 
int col = blockIdx.x*blockDim.x+threadIdx.x; 

с их пределами?

Answer 1

Для наивного матричного умножения это довольно тривиально. Неквадратное умножение матриц должно иметь вида:

C(rowsA x colsB) = A(rowsA x colsA) x B(colsA x colsB) 
    (m)  (n)  (m) ^  ^ (n) 
           |   | 
          must be the same 

Строка и пределы столбцов результата (MXN), а также размер массива нитей, необходимый поэтому просто определяется размером выходной матрицы, который определяется строками A и столбцами B. Следовательно, m=rowsA=rowsC и n=colsB=colsC. Что-то вроде этого:

template <typename T> 
__global__ void mm_kernel(const T *A, const T *B, T *C, int m, int n, int colsA){ 

    int row = blockIdx.y*blockDim.y+threadIdx.y; 
    int col = blockIdx.x*blockDim.x+threadIdx.x; 
    T sum = 0; 
    if ((row < m) && (col < n)){ 
    for (int i = 0; i < colsA; i++) sum += A[colsA*row + i] * B[i*n+col]; 
    C[row*n+col] = sum;} 
} 

в коде хоста вам нужно будет создать сетку, как это:

const int m = 1000; // determines size of output matrix 
const int n = 2000; 
const int blkdim = 16; 
dim3 dimBlock(blkdim,blkdim); 
dim3 dimGrid((n+dimBlock.x-1)/dimBlock.x, (m+dimBlock.y-1)/dimBlock.y); 

(все выше закодирована в браузере, не тестируется)

Если вы попробуйте сделать shared-memory optimized version, это становится заметно более сложным для неквадратных размеров. Но если производительность - это то, что вам нужно, вы должны использовать CUBLAS.