CUDA中的非方形矩阵乘法

对于我的GPU编程类，我们的任务是完成非方形矩阵乘法程序的某些部分。具体来说，内核函数和初始化线程块和内核网格维度。CUDA中的非方形矩阵乘法

我已经将我的代码基于CUDA C编程指南的矩阵乘法代码，但我没有像使用结构那样使用结构，而是修改了我的参数（因为我们不允许更改参数）。我们提供了3个矩阵A，B和C，以及它们的尺寸 - m x k，k x n和m x n。凡结构，用于A.height，我用维数m，它曾经B.width，我用n维等

我碰到的几个问题，其中第一个是我的程序没有通过所包含的测试，这验证了产品矩阵C的正确性。我假设在我的矩阵乘法代码中存在错误，那么这个问题可能是由我调整结构代码引起的。

#include <stdio.h> 
__global__ void mysgemm(int m, int n, int k, const float *A, const float *B, 
     float* C) { 

    /******************************************************************** 
    * 
    * Compute C = A x B 
    * where A is a (m x k) matrix 
    * where B is a (k x n) matrix 
    * where C is a (m x n) matrix 
    * 
    ********************************************************************/ 

    // INSERT KERNEL CODE HERE 
    // Each thread computes one element of C 
    // by accumulating results into Cvalue 
    float Cvalue = 0; 
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
    int col = blockIdx.x * blockDim.x + threadIdx.x; 

    for (int e = 0; e < k; ++e){ 
     Cvalue += (A[row * k + e]) * (B[e * n + col]); 
    } 
    C[row * n + col] = Cvalue; 
}

我的其他问题，我甚至更确定的，涉及到代码初始化线程块和核心网的尺寸。

// Initialize thread block and kernel grid dimensions --------------------- 
    const unsigned int BLOCK_SIZE = 16; // Use 16x16 thread blocks 
//INSERT CODE HERE 
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); 
    dim3 dimGrid(n/dimBlock.x, m/dimBlock.y); 
// Invoke CUDA kernel ----------------------------------------------------- 
//INSERT CODE HERE 
    mysgemm<<<dimGrid, dimBlock>>>(m, n, k, A, B, C);

我明白dimBlock，但我不明白dimGrid，并没有一个适当的想法来使用它作为参数。当我按原样运行代码时，如果我通过的矩阵的维数不是2的幂，那么内核甚至不会启动。如果我使用2的幂，则测试仍然失败。

我道歉，如果我太罗嗦了。这是我的第一篇文章，我想给尽可能多的细节。希望有人能帮助我解决这些问题。

来源

2013-09-25 JoBo

关于cuda矩阵乘法有很多问题，几乎考虑了所有可能的变体。像[这一个]例如（http://stackoverflow.com/questions/18815489/cuda-tiled-matrix-matrix-multiplication-with-shared-memory-and-matrix-size-whic）。也许你应该回顾一些已经被要求提出想法/提示/线索的问题。 –

你的代码目前只能当m和n是16的倍数，这是你的块大小。

你现在可以做的两件事情，使它在任意大小的工作。

充分利用网格尺寸足够大，以代替支付使用地板的n/blockdim.x的整个矩阵C.为你做了，你可以通过
使用该值的小区
```
(n+blockdim.x-1)/blockdim.x 
```
完成第1步后，由于天花板操作，您所乘的矩阵会稍大一些。您可以通过在内核中添加if子句将乘法限制为结果矩阵C的确切大小。

有关更多详细信息，请参阅CUDA文档，尤其是编程指南。

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

来源

2013-09-25 07:22:37 kangshiyin

下面的内核，我下面张贴是一个变种我在

CUDA: Tiled matrix-matrix multiplication with shared memory and matrix size which is non-multiple of the block size

张贴在它不使用共享内存。

__global__ void MatMulNoShared(float* A, float* B, float* C, int ARows, int ACols, int BRows, int BCols, int CRows, int CCols) { 

    float CValue = 0; 

    int Row = blockIdx.y*TILE_DIM + threadIdx.y; 
    int Col = blockIdx.x*TILE_DIM + threadIdx.x; 

    for (int k = 0; k < (TILE_DIM + ACols - 1)/TILE_DIM; k++) { 

     for (int n = 0; n < TILE_DIM; ++n) 
      if ((k*TILE_DIM + n < ACols && Row < ARows) && (k*TILE_DIM + n < BRows && Col < BCols)) 
       CValue += A[Row*ACols + k*TILE_DIM + n] * B[(k*TILE_DIM + n)*BCols + Col]; 

    } 

    if (Row < CRows && Col < CCols) C[((blockIdx.y * blockDim.y + threadIdx.y)*CCols)+(blockIdx.x*blockDim.x)+threadIdx.x]=CValue; 
}

内核两个if语句是由Eric的答复中提到的if语句。

为了您的方便起见，我在下面张贴的完整代码：

#include <stdio.h> 
#include <math.h> 
#include <conio.h> 

#define TILE_DIM 16      // Tile dimension 
#define DIMX 373        
#define DIMY 242 
#define DIMZ 533 

__global__ void MatMulNoShared(float* A, float* B, float* C, int ARows, int ACols, int BRows, int BCols, int CRows, int CCols) { 

    float CValue = 0; 

    int Row = blockIdx.y*TILE_DIM + threadIdx.y; 
    int Col = blockIdx.x*TILE_DIM + threadIdx.x; 

    for (int k = 0; k < (TILE_DIM + ACols - 1)/TILE_DIM; k++) { 

     for (int n = 0; n < TILE_DIM; ++n) 
      if ((k*TILE_DIM + n < ACols && Row < ARows) && (k*TILE_DIM + n < BRows && Col < BCols)) 
       CValue += A[Row*ACols + k*TILE_DIM + n] * B[(k*TILE_DIM + n)*BCols + Col]; 

    } 

    if (Row < CRows && Col < CCols) C[((blockIdx.y * blockDim.y + threadIdx.y)*CCols)+(blockIdx.x*blockDim.x)+threadIdx.x]=CValue; 
} 

int main() { 

    int CCols = DIMZ, CRows=DIMX, ACols=DIMY, ARows=DIMX, BCols=DIMZ, BRows=DIMY; 

    dim3 dimBlock(TILE_DIM, TILE_DIM, 1); 
    dim3 dimGrid; 

    dimGrid.x = (CCols + dimBlock.x - 1)/dimBlock.x; 
    dimGrid.y = (CRows + dimBlock.y - 1)/dimBlock.y; 

    float *deviceA, *deviceB, *deviceC; 

    float* hostA = (float*)malloc(DIMX*DIMY*sizeof(float)); 
    float* hostB = (float*)malloc(DIMY*DIMZ*sizeof(float)); 
    float* hostC = (float*)malloc(DIMX*DIMZ*sizeof(float)); 
    float* hostCp = (float*)malloc(DIMX*DIMZ*sizeof(float)); 

    for (int x = 0; x<DIMX; x++) 
     for (int y = 0; y<DIMY; y++) { 
      hostA[x*DIMY+y] = rand()/(float)RAND_MAX; 
      hostB[x*DIMY+y] = rand()/(float)RAND_MAX; 
     } 

    cudaMalloc((void **)&deviceA, DIMX*DIMY*sizeof(float)); 
    cudaMalloc((void **)&deviceB, DIMY*DIMZ*sizeof(float)); 
    cudaMalloc((void **)&deviceC, DIMX*DIMZ*sizeof(float)); 

    cudaMemcpy(deviceA, hostA, DIMX*DIMY*sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(deviceB, hostB, DIMY*DIMZ*sizeof(float), cudaMemcpyHostToDevice); 

    MatMulNoShared<<<dimGrid , dimBlock>>>(deviceA , deviceB , deviceC , ARows , ACols, BRows ,BCols , CRows , CCols); 

    cudaMemcpy(hostC, deviceC, DIMX*DIMZ*sizeof(float), cudaMemcpyDeviceToHost); 

    return 0; 
}

注意，两个指令

dimGrid.x = (CCols + dimBlock.x - 1)/dimBlock.x; 
    dimGrid.y = (CRows + dimBlock.y - 1)/dimBlock.y;

确保矩阵的全覆盖平铺，如前所述在埃里克答案的第一点。

来源

2013-09-25 14:03:19 JackOLantern

非常感谢^ _ ^ –

CUDA中的非方形矩阵乘法

回答

相关问题