c++ - Cuda 矩阵示例 block 大小

Question

我刚刚开始学习 CUDA，并且一直在查看 NVIDIA 网站上的示例。具体来说，我已经实现了矩阵乘法的非共享版本(第一个示例是非共享版本，即使它在共享内存部分):

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

我在更改 block 大小时遇到​​输出问题。 NVIDIA 代码的默认 block 大小为 16，当我将两个矩阵相乘时，这为我提供了正确的输出。但是，如果我将 block 大小更改为大于 16 的任何值(同时仍然是 16 的倍数)，我将得到矩阵中所有元素的零输出。我也在我的笔记本电脑上对此进行了测试，并注意到超过 32 而不是 16 的任何结果都是相同的。有人可以解释发生了什么吗？我在 SLI 中有两个 9800GTX+ 视频卡，所以我的最大块大小应该为 (512,512,1)。为什么我只能做16个？

此外，我注意到矩阵乘法的共享版本(也在 NVIDIA 页面上)中存在相同的行为。

我没有发布我的代码，因为如果我直接从 NVIDIA 网站复制代码，我会遇到同样的问题。

我非常感谢任何对此或资源的帮助，以了解有关此类 CUDA 详细信息的更多信息。

谢谢!

我已按要求附上代码:

    #include "stdio.h"
    #include <cuda.h>
    #include <assert.h>
    #include <time.h>
    #include <math.h>

    // This is an example CUDA program that compares the timings of a matrix multiplication.
    // The comparisons are between the CPU, GPU, and the GPU with shared memory.

    #define BLOCK_SIZE 32

    typedef struct {

    int width;
    int height;
    int stride;
    float* elements;

    } Matrix;

    typedef void (*FuncPtr)(Matrix& A, Matrix& B, Matrix& C);

    void multiplyMatrix(Matrix& A, Matrix& B, Matrix& C);

    // Helper declarations
    void initializeMatrix(Matrix& A, int rows, int cols, float val);
    void copyMatrix(Matrix& dest, Matrix& src);
    void freeMatrix(Matrix& A);
    void printError(cudaError_t err);
    void printMat(Matrix& A);
    void setVal(Matrix& A, float val);
    double applyMultFunc(FuncPtr func, Matrix& A, Matrix& B, Matrix& C, int numOfIters);

    // CUDA declarations
    __global__ void cudaMultMat(Matrix A, Matrix B, Matrix C);


   int main() {

       printf("Beginning Matrix Multiplication Comparison\n");

       // Initialize matrix
       Matrix A, B, C;
       int rowsA = 32;
       int colsA = 32;
       int colsB = 32;
       initializeMatrix(A, rowsA, colsA, 5.0f);
       initializeMatrix(B, colsA, colsB, 2.0f);
       initializeMatrix(C, rowsA, colsB, 0.0f);

       // C = A * B using CPU, GPU, and GPU with shared memory
       FuncPtr gpuMatMult = &multiplyMatrix;
       int numOfIterations = 100;
       double multTime = applyMultFunc(gpuMatMult, A, B, C, numOfIterations);

       printMat(C); 

       // Update user
       printf("Normal Mat Mult Time: %f\n", multTime);


       // Cleanup
       freeMatrix(A);
       freeMatrix(B);
       freeMatrix(C);

       printf("\nPress Enter to continue...\n");
       getchar();

       return 0;

  }

  void multiplyMatrix(Matrix& A, Matrix& B, Matrix& C) {

    // Initialize device matrices
    Matrix deviceA, deviceB, deviceC;
    copyMatrix(deviceA, A);
    copyMatrix(deviceB, B);
    copyMatrix(deviceC, C);

    // Initialize number of blocks and threads
    dim3 numOfThreadsPerBlock(BLOCK_SIZE, BLOCK_SIZE);
    int xSize = (C.width + numOfThreadsPerBlock.x - 1) / numOfThreadsPerBlock.x;
    int ySize = (C.height + numOfThreadsPerBlock.y - 1) / numOfThreadsPerBlock.y;
    dim3 numOfBlocks(xSize, ySize);

    // Call CUDA kernel
    cudaMultMat<<<numOfBlocks, numOfThreadsPerBlock>>>(deviceA, deviceB, deviceC);
    printError(cudaThreadSynchronize());
    printError(cudaMemcpy(C.elements, deviceC.elements, C.height * C.width * sizeof(float), cudaMemcpyDeviceToHost));

    // Free cuda memory
    printError(cudaFree(deviceA.elements));
    printError(cudaFree(deviceB.elements));
    printError(cudaFree(deviceC.elements));

  }



 // CUDA definitions

 // GPU matrix multiplication (non-shared memory)
 __global__ void cudaMultMat(Matrix A, Matrix B, Matrix C) {

    // If the matrices are of the wrong size then return
    if(A.width != B.height) {
        return;
    }

    // Initialize the indexes into the grid
    int col = (blockDim.x * blockIdx.x) + threadIdx.x;
    int row = (blockDim.y * blockIdx.y) + threadIdx.y;

    // Initialize the result
    float cVal = 0.0f;

    // Find the result for the dot product of a row of A and a column of B
    for(int i = 0; i < A.width; i++) {

        cVal += A.elements[row * A.width + i] * B.elements[i * B.width + col];

     }

     // If we are in bounds then save the result
     if(row < C.height && col < C.width) {
        C.elements[row * C.width + col] = cVal;
     }

  } 

  // Helper functions
  void initializeMatrix(Matrix& A, int rows, int cols, float val) {

    A.width = cols;
    A.height = rows;
    A.stride = A.width;
    int numOfElements = A.width * A.height;
    A.elements = (float*) malloc(numOfElements * sizeof(float));
    for(int i = 0; i < numOfElements; i++) {
        A.elements[i] = val;
    }

   }

   void copyMatrix(Matrix& dest, Matrix& src) {

    dest.width = src.width;
    dest.height = src.height;
    dest.stride = src.stride;
    int size = src.width * src.height * sizeof(float);
    printError(cudaMalloc(&dest.elements, size)); 
    printError(cudaMemcpy(dest.elements, src.elements, size, cudaMemcpyHostToDevice));

   }

   void freeMatrix(Matrix& A) {
    free(A.elements);
   }

   void printError(cudaError_t err) {
    if(err != 0) {
        printf("CUDA ERROR: %s\n", cudaGetErrorString(err));
        getchar();
    }

    }

    void printMat(Matrix& A) {

    printf("*********************************\n");
    for(int i = 0; i < A.height; i++) {
         for(int j = 0; j < A.width; j++) {
             int index = i * A.width + j;
             printf("%2.1f, ", A.elements[index]); 
         }
         printf("\n");
     }

  }

  void setVal(Matrix& A, float val) {

     for(int i = 0; i < A.width * A.height; i++) {
          A.elements[i] = val;
     }

  }

  double applyMultFunc(FuncPtr func, Matrix& A, Matrix& B, Matrix& C, int numOfIters) {

    clock_t startTime = clock();
    for(int i = 0; i < numOfIters; i++) {
        func(A, B, C);
     } 
     clock_t endTime = clock();
     return (double) (endTime - startTime) / CLOCKS_PER_SEC;

   }

Answer 1

当您增加 block 大小时，您超出了 GPU 的每 block 线程数规范。

无论您如何创建 block ，9800GTX 的每个 block 都有 512 个线程的限制。 16*16 = 256 没问题。 32 x 32 = 1024 这是不对的。在这种情况下，内核无法运行，因此输出不正确。

你的笔记本电脑可能有一个更新的 GPU，它支持每 block 1024 个线程，所以 32 x 32 是可以的，但更大的就不行了。

如果添加proper cuda error checking您可以通过代码确认这一点。请注意，此代码似乎具有 cuda 错误检查，但在内核调用上实现的检查不完整。研究我提供的链接，您会看到不同之处。如果您修改带有完整错误检查的代码，您将看到该错误。