2

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

我的代码基于 CUDA C 编程指南的矩阵乘法代码，但我没有像他们那样使用结构，而是修改了我的代码以仅使用给定的参数（因为我们不允许更改参数）。我们得到了 3 个矩阵 A、B 和 C，以及它们的维度——mxk、kxn 和 mxn。在结构使用 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 的幂，测试仍然会失败。

如果我太啰嗦了，我很抱歉。这是我的第一篇文章，我想提供尽可能多的细节。希望有人可以帮助我解决这些问题。

4

2 回答 2

5

我在下面发布的以下内核是我发布的内核的变体

CUDA：具有共享内存和非块大小倍数的矩阵大小的平铺矩阵-矩阵乘法

因为它不使用共享内存。

__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内核中的两个语句if是 Eric 回答中提到的语句。

为了您的方便，我在下面发布了完整的代码：

#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;

确保矩阵的完整平铺覆盖，如 Eric 回答的第 1 点所述。

于 2013-09-25T14:03:19.663 回答
1

您的代码目前仅在 m 和 n 是 16 的倍数时才有效，这是您的块大小。

您现在可以做两件事来使其适用于任意尺寸。

  1. 使网格尺寸足够大以覆盖整个矩阵C。您可以使用的ceil

     (n+blockdim.x-1)/blockdim.x

  2. 完成第 1 步后，由于天花板操作，您要相乘的矩阵会稍大一些。然后，您可以通过在内核中添加 if 子句将乘法限制为结果矩阵 C 的确切大小。

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

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

于 2013-09-25T07:22:37.337 回答