CUDA8.0矩阵乘法例子解释（matrixMul.cpp）

最新推荐文章于 2026-06-16 17:54:40 发布

原创最新推荐文章于 2026-06-16 17:54:40 发布 · 5.2k 阅读

4 ·

本内容遵循CC 4.0 BY-SA版权协议

机器学习和GPU 专栏收录该内容

22 篇文章

订阅专栏

本文详细介绍了使用CUDA实现矩阵乘法的过程，包括矩阵大小定义、内存分配与拷贝、核心函数调用及参数传递等内容。

通过学习英伟达自带的例子matrixMul学CUDA库的使用。

简略部分垃圾。只说核心代码。

这个例子是实现 C=A*B的矩阵相乘

// Use a larger block size for Fermi and above
    int block_size = 32;


    //original:
    dim3 dimsA(5*2*block_size, 5*2*block_size, 1);
    dim3 dimsB(5*4*block_size, 5*2*block_size, 1);


    // reduce sizes to avoid running out of memory
    //dim3 dimsA(32,32, 1);
    //dim3 dimsB(32,32,1);

定义了矩阵的大小。 dim3是一个三维结构。xyz。分别代表长宽高。是cuda的内置结构

struct __device_builtin__ dim3
{
    unsigned int x, y, z;
};

矩阵的z高度是1。表示是一个面，可以忽略不看。

矩阵A列5*2*32 （320）行 5*2*32 （320）

矩阵B列5*4*32 （640）行 5*2*32（320）

下面说为了避免内存用完。可以减小一些宽和高。

随后调用

int matrix_result = matrixMultiply(argc, argv, block_size, dimsA, dimsB);

实现计算过程

进入matrixMultiply以后的代码如下

unsigned int size_A = dimsA.x * dimsA.y;
unsigned int mem_size_A = sizeof(float) * size_A;
float *h_A = (float *)malloc(mem_size_A);
unsigned int size_B = dimsB.x * dimsB.y;
unsigned int mem_size_B = sizeof(float) * size_B;
float *h_B = (float *)malloc(mem_size_B);

// Initialize host memory
const float valB = 0.01f;
constantInit(h_A, size_A, 1.0f);
constantInit(h_B, size_B, valB);

size_A和size_B分别是AB矩阵元素的个数。

mem_size_A和mem_size_B分别是矩阵所需要的内存的大小，这里每个元素都是浮点数。

h_A和h_B则是分配好的内存的起始地址。

随后对A B的数据进行初始化。A每个元素赋值为1.0f。B每个元素赋值为0.01f (constantInit就是实现赋值的过程)

    // Allocate device memory
    CUdeviceptr d_A, d_B, d_C;

    char *ptx, *kernel_file;
    size_t ptxSize;

定义一些后续需要的变量

kernel_file = sdkFindFilePath("matrixMul_kernel.cu", argv[0]);
compileFileToPTX(kernel_file, 0, NULL, &ptx, &ptxSize);
CUmodule module = loadPTX(ptx, argc, argv);
找到cu文件的位置。cu文件是C语言语法的，就是后缀不同，这里面主要是实现算法。随后调用
compile将cu文件编译成GPU可以理解执行的代码，然后通loadPTX执行加载函数。
就是将cu文件编译成GPU可以理解的东西。相当于翻译的过程。然后加载进来。
这里需要argc和argv是可能在argv中指定使用某个特殊的设备。比如我有几张显卡。可能要选择这样。否则就按照默认配置来选择。
其次关于 compileFileToptX和loadPtx。是拿底层的SDK进行了基础的封装。你也可以调用底层SDK实现。具体的函数在nvrtc_helper.h文件里面。
// Allocate host matrix C
dim3 dimsC(dimsB.x, dimsA.y, 1);
unsigned int mem_size_C = dimsC.x * dimsC.y * sizeof(float);
float *h_C = (float *) malloc(mem_size_C);
定义计算后保存的结果，不细说了。

checkCudaErrors(cuMemAlloc(&d_A, mem_size_A));
checkCudaErrors(cuMemAlloc(&d_B, mem_size_B));
checkCudaErrors(cuMemAlloc(&d_C, mem_size_C));
// copy host memory to device
checkCudaErrors(cuMemcpyHtoD(d_A, h_A, mem_size_A));
checkCudaErrors(cuMemcpyHtoD(d_B, h_B, mem_size_B));

cuMemAlloc这次是在显存中分配内存了。分配了矩阵A B 和结果C。所以这里要注意。千万不要内存不够了。因此尽量关闭不需要用到的显存。
接下来吧矩阵A B 的内存数据（h_A h_B）拷贝到显存(d_A d_B)当中去。
// Setup execution parameters
dim3 threads(block_size, block_size);
dim3 grid(dimsB.x / threads.x, dimsA.y / threads.y);

定义了执行参数3维 x y z
threads (x=block_size ，y=block_size ，z=1) threads(32,32,1)
grid (x=矩阵B的x/threads.x，y=矩阵A的y/thread.Y，z=1) grid (640/32=20 , 320/32=10,1 )=（20,10,1）

CUfunction kernel_addr;
if (block_size == 16)
{
checkCudaErrors(cuModuleGetFunction(&kernel_addr, module, "matrixMulCUDA_block16"));
}
else
{
checkCudaErrors(cuModuleGetFunction(&kernel_addr, module, "matrixMulCUDA_block32"));
}

通过cuModuleGetFunction得到cu模块中函数的地址，放在kernel_addr中。
void matrixMulCUDA_block32(float *C, float *A, float *B, int wA, int wB)
这个就是cu文件中实际的执行函数了，这里先看一下函数，有个映像。

void *arr[] = { (void *)&d_C, (void *)&d_A, (void *)&d_B, (void *)&dimsA.x, (void *)&dimsB.x };

接着定义了函数需要的参数。
int nIter = 300;

for (int j = 0; j < nIter; j++)
{
checkCudaErrors(cuLaunchKernel(kernel_addr,
grid.x, grid.y, grid.z, /* grid dim */
threads.x, threads.y, threads.z, /* block dim */
0,0, /* shared mem, stream */
&arr[0], /* arguments */
0));

checkCudaErrors(cuCtxSynchronize());
}
执行计算操作。
这里重点介绍cuLaunchKernel函数
CUresult CUDAAPI cuLaunchKernel(CUfunction f,
unsigned int gridDimX,
unsigned int gridDimY,
unsigned int gridDimZ,
unsigned int blockDimX,
unsigned int blockDimY,
unsigned int blockDimZ,
unsigned int sharedMemBytes,
CUstream hStream,
void **kernelParams,
void **extra);
调用核心执行f函数。处理gridDimX gridDimY gridDimZ大小的数据块。每个数据块包含 blockDimX blockDImY blockDimZ线程
sharedMemBytes指定了每个数据块可以共享的动态内存。

f函数的参数可以有两种形式
1 通过kernelParams参数指定。如果f有N个参数。那么kernelParams就是一个N大小的参数数组的指针。从kernelParams[0]到
kernelParams[N-1]，每个参数必须指向核心将要拷贝的一块内存，这里意思是核心需要的是地址，而不是值。
比如f(int x) 需要的是一个int的x。那么kernelParams[0]=&x; 而不能直接指定为x。这点要特别注意。

核心参数的数量，大小和偏移不需要指定，那些都是直接从核心的image里面直接得到，（这句暂未理解什么意思）。

2参数也可以通过程序打包到一个单独的buffer通过extra参数传递过去。这个就需要用对其等来处理好每个参数的大小等等。

extra的存在主要是允许 culaunchKernel函数拿到一些不通用的参数。extra指定了这些参数的名字和对应的值。必须以NULL或者CU_LAUNCH_PARAM_END结束
比如
void *config[] = {
CU_LAUNCH_PARAM_BUFFER_POINTER, argBuffer,
CU_LAUNCH_PARAM_BUFFER_SIZE, &argBufferSize,
CU_LAUNCH_PARAM_END
};
status = cuLaunchKernel(f, gx, gy, gz, bx, by, bz, sh, s, NULL, config);

并且两种方式只能选用一种。两个值都指定的话会导致函数执行错误。
该cuLaunchKernel函数等同的调用函数顺序如下
* Calling ::cuLaunchKernel() sets persistent function state that is
* the same as function state set through the following deprecated APIs:
* ::cuFuncSetBlockShape(),
* ::cuFuncSetSharedSize(),
* ::cuParamSetSize(),
* ::cuParamSeti(),
* ::cuParamSetf(),
* ::cuParamSetv().
通过cuLaunchKernel调用后，会覆盖前面设置的块大小。参数信息。共享大小等等。
* \param f - Kernel to launch
* \param gridDimX - Width of grid in blocks
* \param gridDimY - Height of grid in blocks
* \param gridDimZ - Depth of grid in blocks
* \param blockDimX - X dimension of each thread block
* \param blockDimY - Y dimension of each thread block
* \param blockDimZ - Z dimension of each thread block
* \param sharedMemBytes - Dynamic shared-memory size per thread block in bytes
* \param hStream - Stream identifier
* \param kernelParams - Array of pointers to kernel parameters
* \param extra - Extra options

// Copy result from device to host
checkCudaErrors(cuMemcpyDtoH(h_C, d_C, mem_size_C));
将结果从显存直接拷贝回内存

free(h_A);
free(h_B);
free(h_C);

checkCudaErrors(cuMemFree(d_A));
checkCudaErrors(cuMemFree(d_B));
checkCudaErrors(cuMemFree(d_C));
最后就是资源的释放