使用共享内存时不执行 CUDA 内核代码
Posted
技术标签:
【中文标题】使用共享内存时不执行 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);
问题是你的块和网格参数颠倒了,应该是:
shmMM<<<grid, block>>>(a, b, c);
我并不是说我已经完全调试了您的应用程序。但这就是原因的来源:
CUDA 内核代码不执行
这些代码行也不正确:
cudaFree(a_h);
cudaFree(b_h);
但这不是您所问问题的根源。 cudaMallocHost
对应的空闲操作是cudaFreeHost
,上面提到了here
【讨论】:
啊,我刚试过compute-sanitizer
,它列出了你提到的所有错误。而且cuda错误检查也是我需要学习的一个很好的编码实践。非常感谢,罗伯特。以上是关于使用共享内存时不执行 CUDA 内核代码的主要内容,如果未能解决你的问题,请参考以下文章