使用共享内存时不执行 CUDA 内核代码

Posted 2023-03-23

技术标签:

【中文标题】使用共享内存时不执行 CUDA 内核代码【英文标题】：CUDA kernel code does not execute when using shared memory 【发布时间】：2022-01-06 17:41:18 【问题描述】：

我正在学习使用共享内存来优化 cuda 代码。我遵循了 Nvidia 材料中的大部分实现。但我发现我的设备代码从未执行过。任何人都可以帮我弄清楚为什么？我错过了什么？谢谢。

#include <stdio.h>
#include <cuda_runtime.h>
#include <chrono>
#define BLOCKSIZE 16

typedef struct 
    int height;
    int width;
    int stride;
    float *element;
 Matrix;

void initData(float *p, int size)
    for (int t=0; t<size; t++)
        p[t] = (float)(rand()&0xffff)/1000.0f;
    


__device__ float getElement(Matrix a, int row, int col)

    return a.element[row*a.stride+col];


__device__ Matrix getSubM(Matrix a, int row, int col)

    Matrix res;
    res.height = BLOCKSIZE;
    res.width = BLOCKSIZE;
    res.stride = a.width;
    res.element = &a.element[row*BLOCKSIZE*a.stride+col*BLOCKSIZE];

    return res;


__device__ void setElement(Matrix a, int row, int col, float val)

    a.element[row*a.stride+col] = val;


__global__ void shmMM(Matrix a, Matrix b, Matrix c)


    int blockRow = blockDim.y;
    int blockCol = blockDim.x;

    Matrix Csub = getSubM(c, blockRow, blockCol);

    int row = threadIdx.y;
    int col = threadIdx.x;

    float tmp = 0;

    for (int i=0; i < a.width/BLOCKSIZE; i++)
       
        Matrix a_sub = getSubM(a, blockRow, i);
        Matrix b_sub = getSubM(b, i, blockCol);
        __shared__ float A[BLOCKSIZE][BLOCKSIZE];
        __shared__ float B[BLOCKSIZE][BLOCKSIZE];

        A[row][col] = getElement(a, row, col);
        B[row][col] = getElement(b, row, col); 
        __syncthreads();
        for (int e = 0; e < BLOCKSIZE; e++)
        
            tmp += A[row][e]*B[e][col];
        
        __syncthreads();
    
    //printf("debug: %f.\n", tmp);
    setElement(Csub, row, col, tmp);


int main()

    Matrix a, b, c;
    int size = 1<<12;
    a.height = a.width = size;
    b.height = b.width = size;
    c.height = c.width = size;
    a.stride = a.width;
    b.stride = b.width;
    c.stride = c.width;
    float *a_h, *b_h, *c_h;
    cudaMallocHost((float**)&a_h, a.height*a.width*sizeof(float));
    cudaMallocHost((float**)&b_h, b.height*b.width*sizeof(float));
    initData(a_h, a.height*a.width);
    initData(b_h, b.height*b.width);
    c_h = (float*)malloc(c.height*c.width*sizeof(float));
    float *a_d, *b_d, *c_d;
    cudaMalloc((float**)&a.element, a.height*a.width*sizeof(float));
    cudaMalloc((float**)&b.element, b.height*b.width*sizeof(float));
    cudaMalloc((float**)&c.element, c.height*c.width*sizeof(float));
    cudaMemcpy(a.element, a_h, a.height*a.width*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(b.element, b_h, b.height*b.width*sizeof(float), cudaMemcpyHostToDevice);
    dim3 block(BLOCKSIZE, BLOCKSIZE);
    dim3 grid((b.width-1)/block.x+1, (a.height-1)/block.y+1);
    //naiveMM<<<block, grid>>>(a, b, c);
    shmMM<<<block, grid>>>(a, b, c);
    cudaMemcpy(c_h, c.element, c.height*c.width*sizeof(float), cudaMemcpyDeviceToHost);

    cudaDeviceSynchronize();

    cudaFree(a_h);
    cudaFree(b_h);
    free(c_h);
    cudaFree(a.element);
    cudaFree(b.element);
    cudaFree(c.element);
    return 0;

由于没有报告编译错误和运行时错误，我无法弄清楚。

【问题讨论】：

【参考方案1】：

因为没有报告编译错误和运行时错误。

如果您未能使用proper CUDA error checking，您将不会收到任何报告的运行时错误。 我建议您在 CUDA 代码遇到问题时随时使用。根据您的 GPU，使用诸如 cuda-memcheck 或 compute-sanitizer 之类的清理程序运行代码也是一种很好的做法。

如果你做了任何这些，你会在你的内核启动时得到一个无效的配置参数错误。那本来应该或应该将您的注意力集中在此代码上：

dim3 block(BLOCKSIZE, BLOCKSIZE);
dim3 grid((b.width-1)/block.x+1, (a.height-1)/block.y+1);
//naiveMM<<<block, grid>>>(a, b, c);
shmMM<<<block, grid>>>(a, b, c);

问题是你的块和网格参数颠倒了，应该是：