c - 如何实现 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

结果的行和列限制 (m x n)，以及因此所需的线程数组的大小仅由输出矩阵大小定义，输出矩阵大小由 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 .