CUDA Driver API 矩阵乘法实战：从 fatbin 模块加载到 cuLaunchKernel 的完整驱动式编程指南

CUDA Driver API 矩阵乘法实战从 fatbin 模块加载到 cuLaunchKernel 的完整驱动式编程指南【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples本篇文章以 NVIDIA CUDA Samples 仓库中的matrixMulDrv位于 cpp/0_Introduction/matrixMulDrv示例为主体系统讲解如何脱离熟悉的 Runtime API、直接使用 CUDA Driver API 完成矩阵乘法从设备枚举、上下文创建、fatbin 模块加载、基于占用的线程块尺寸选择到 CUDA 4.0 风格的cuLaunchKernel内核启动与结果校验。读完本文你将掌握 Driver API 编程的完整调用链并能理解cuLaunchKernel两种参数传递方式的差异为后续阅读驱动式Drv / nvrtc / JIT示例打下基础。一、示例概览与设计定位根据 matrixMulDrv/README.md本示例实现的是经典矩阵乘法C A × B其核心价值在于使用 CUDA 4.0 引入的 kernel launch Driver API即cuLaunchKernel来驱动内核执行。README 同时明确指出它的定位为清晰阐释各种 CUDA 编程原理而编写并非以提供最高性能的通用矩阵乘法内核为目标如果追求高性能矩阵乘法应当使用 CUBLAS。示例内核对矩阵尺寸A 为 4×block_size 宽、6×block_size 高B 为 4×block_size 宽见 matrixMul.h做了简化假设以便把注意力集中在 Driver API 的编程模式上。示例的关键概念Key Concepts为两项CUDA Driver API与Matrix Multiply整份代码围绕这两点展开。二、运行环境与支持范围2.1 支持的 SM 架构README 列出的支持范围覆盖 Maxwell 到 Blackwell/Ada 之后的各代架构SM 版本代表架构代次SM 5.0 / 5.2 / 5.3MaxwellSM 6.0 / 6.1PascalSM 7.0 / 7.2 / 7.5Volta / TuringSM 8.0 / 8.6 / 8.7 / 8.9Ampere / AdaSM 9.0Hopper 及之后从 CMakeLists.txt 可以看到构建时默认的 CUDA 架构列表为75 80 86 87 89 90 100 110 120并会为每个架构生成对应的 gencode 指令见下文构建章节因此只要你的 GPU 计算能力Compute Capability落在上述范围内即可运行。2.2 支持的操作系统与 CPU 架构操作系统Linux、WindowsCPU 架构x86_64、armv7l、aarch642.3 前置条件按照 README 的 Prerequisites你需要先为对应平台下载并安装 CUDA Toolkit。由于示例通过 Driver API 直接调用驱动层接口cuInit、cuDeviceGet*、cuModuleLoadData等运行时还需要系统已安装对应的 NVIDIA GPU 驱动。三、Driver API 全景示例用到的完整接口清单与 Runtime APIcudaMalloc、cudaMemcpy等不同Driver API 的函数以前缀cu开头且所有资源句柄设备CUdevice、上下文CUcontext、模块CUmodule、函数CUfunction、指针CUdeviceptr都需要显式创建与销毁。README 列出的本示例涉及的 Driver API 如下接口在本示例中的作用cuCtxCreate在指定设备上创建 CUDA 上下文initCUDA内调用cuCtxDestroy程序结束时销毁上下文cuDeviceGetAttribute查询设备计算能力major/minor等属性cuDeviceGetName获取设备名称字符串cuDeviceTotalMem查询设备全局内存总量cuModuleLoadData从内存中的 fatbin 二进制数据创建模块cuModuleGetFunction从模块中按名字取出内核函数句柄CUfunctioncuOccupancyMaxPotentialBlockSize基于占用率计算最优线程块大小cuMemAlloc/cuMemFree在设备上分配 / 释放线性内存返回CUdeviceptrcuMemcpyHtoD主机内存 → 设备内存拷贝cuMemcpyDtoH设备内存 → 主机内存拷贝cuLaunchKernelCUDA 4.0 起的统一内核启动接口支持简单与高级两种参数传递此外示例还大量使用Common/helper_cuda_drvapi.h提供的辅助函数checkCudaErrors、findCudaDeviceDRV、findFatbinPath等这部分见下文。四、源码结构与执行流程示例由以下关键文件构成均在 cpp/0_Introduction/matrixMulDrv 目录下matrixMulDrv.cpp主机端程序主体包含main→runTest→initCUDA的调用链matrixMul.h矩阵尺寸宏定义matrixMul_kernel.cu设备端模板内核及三个extern C导出内核CMakeLists.txt构建脚本负责生成 fatbin 并链接驱动库。主程序运行流程对应 matrixMulDrv.cpp 中的runTest可以概括为六步initCUDA完成驱动初始化、设备选择、上下文创建、模块加载与内核函数解析在主机侧初始化矩阵 A全 1.0与 B全 0.01并分配设备内存cuMemcpyHtoD把 A、B 拷贝到设备通过cuLaunchKernel启动矩阵乘法内核cuMemcpyDtoH把结果 C 拷回主机逐元素与参考值WA * valB比较误差阈值1e-5释放设备内存、卸载模块、销毁上下文。其中第 4 步是本文重点代码中提供了两种cuLaunchKernel的参数传递方式下文第六节详述。五、驱动初始化、模块加载与线程块尺寸选择5.1 设备选择与上下文创建initCUDAmatrixMulDrv.cpp首先调用findCudaDeviceDRV选择设备。该辅助函数定义在 Common/helper_cuda_drvapi.h若命令行指定了-deviceid则调用gpuDeviceInitDRV使用指定设备并检查其计算模式是否被禁止CU_COMPUTEMODE_PROHIBITED否则调用gpuGetMaxGflopsDeviceIdDRVhelper_cuda_drvapi.h通过CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT、CU_DEVICE_ATTRIBUTE_CLOCK_RATE、计算能力等属性估算各设备的计算性能多处理器数 × 每 SM 核心数 × 时钟频率自动选择性能最高的 GPU。随后代码通过cuDeviceGetAttribute获取计算能力打印 GPU Device has SM X.Y compute capability通过cuDeviceTotalMem打印全局内存总量最后调用cuCtxCreate创建上下文——这是 Driver API 与 Runtime API 最显著的区别之一上下文CUcontext必须显式创建并最终显式销毁cuCtxDestroy。5.2 从 fatbin 加载模块cuModuleLoadDataDriver API 中内核代码不再由nvcc在链接期直接嵌入可执行文件而是以fatbinFAT binary的形式在运行时加载。本示例的流程是用findFatbinPathhelper_cuda_drvapi.h定位matrixMul_kernel64.fatbin文件宏FATBIN_FILE默认值为该文件名并把整个文件内容读入内存流调用cuModuleLoadData(cuModule, fatbin.str().c_str())从内存中的二进制数据创建CUmodule遍历内核名列表{matrixMul_bs32_64bit, matrixMul_bs16_64bit, matrixMul_bs8_64bit}用cuModuleGetFunction取出对应的CUfunction句柄。5.3 基于占用率自动选择线程块尺寸这是示例中体现现代 CUDA 实践的重要一环matrixMulDrv.cpp默认block_size 32依次尝试三个内核对每个内核调用cuOccupancyMaxPotentialBlockSize(blocksPerGrid, threadsPerBlock, cuFunction, 0, 2 * block_size * block_size * sizeof(float), 0)其中动态共享内存大小2 * block_size * block_size * sizeof(float)对应共享内存中 A、B 两个block_size × block_size的 float 数组若block_size * block_size threadsPerBlock即该块尺寸能被占用率计算接受则选定该内核否则block_size / 2继续尝试下一个较小尺寸的内核。也就是说示例在运行时根据当前 GPU 的占用率约束在32×32、16×16、8×8三种块配置之间自动降级选择并把最终选定的block_size回传给调用方用于网格配置与共享内存计算。六、cuLaunchKernel 的两种参数传递方式示例核心代码在runTest中通过一个if (1)分支展示了cuLaunchKernel的两种调用形式这是本示例相对同类矩阵乘法示例最独特的地方。6.1 网格与块配置dim3 block(block_size, block_size, 1); dim3 grid(WC / block_size, HC / block_size, 1);网格大小为(WC/block_size) × (HC/block_size)其中WC WB 4 * block_size、HC HA 6 * block_size见 matrixMul.h因此网格恰为4 × 6的二维块阵每个块内block_size × block_size个线程——每个线程计算输出矩阵 C 的一个元素。6.2 方式一简单方法参数数组 取地址size_t Matrix_Width_A (size_t)WA; size_t Matrix_Width_B (size_t)WB; void *args[5] {d_C, d_A, d_B, Matrix_Width_A, Matrix_Width_B}; checkCudaErrors(cuLaunchKernel(matrixMul, grid.x, grid.y, grid.z, // 网格维度 block.x, block.y, block.z, // 块维度 2 * block_size * block_size * sizeof(float), // 动态共享内存 NULL, // 流 args, // 内核参数数组 NULL)); // 额外选项简单方法的要点把每个内核参数d_C、d_A、d_B、宽度Matrix_Width_A、Matrix_Width_B的指针放入void *args[]数组cuLaunchKernel会自动按内核签名解引用这些指针取值。这正是 CUDA 4.0 简化后的推荐写法——参数按声明顺序排列直观且不易出错。6.3 方式二高级方法CU_LAUNCH_PARAM_* 缓冲区代码中else分支展示了传统的高级参数传递方式把内核参数按值序列化进一块字节缓冲区再用cuLaunchKernel的extras参数以CU_LAUNCH_PARAM_BUFFER_POINTER / CU_LAUNCH_PARAM_BUFFER_SIZE / CU_LAUNCH_PARAM_END三元组描述这块缓冲区int offset 0; char argBuffer[256]; // 注意这里并非解引用 CUdeviceptr而是把参数值本身写入缓冲区 *(reinterpret_castCUdeviceptr *(argBuffer[offset])) d_C; offset sizeof(d_C); *(reinterpret_castCUdeviceptr *(argBuffer[offset])) d_A; offset sizeof(d_A); *(reinterpret_castCUdeviceptr *(argBuffer[offset])) d_B; offset sizeof(d_B); size_t Matrix_Width_A (size_t)WA; size_t Matrix_Width_B (size_t)WB; *(reinterpret_castsize_t *(argBuffer[offset])) Matrix_Width_A; offset sizeof(Matrix_Width_A); *(reinterpret_castsize_t *(argBuffer[offset])) Matrix_Width_B; offset sizeof(Matrix_Width_B); void *kernel_launch_config[5] { CU_LAUNCH_PARAM_BUFFER_POINTER, argBuffer, CU_LAUNCH_PARAM_BUFFER_SIZE, offset, CU_LAUNCH_PARAM_END}; cuLaunchKernel(matrixMul, grid.x, grid.y, grid.z, block.x, block.y, block.z, 2 * block_size * block_size * sizeof(float), NULL, NULL, reinterpret_castvoid **(kernel_launch_config));两种方式在功能上等价简单方法适合参数较少、签名固定的场景高级方法则更接近底层驱动接口的数据布局适合需要动态构造参数列表的场景。理解这两种形式对阅读其他 Driver API 示例如matrixMulDynlinkJIT、ptxjit很有帮助。6.4 内核启动后的数据流启动完成后示例用cuMemcpyDtoH把结果拷回主机并逐元素校验for (int i 0; i (int)(WC * HC); i) { if (fabs(h_C[i] - (WA * valB)) 1e-5) { printf(Error! Matrix[%05d]%.8f, ref%.8f error term is 1e-5\n, ...); correct false; } } printf(%s\n, correct ? Result PASS : Result FAIL);由于 A 全为 1.0、B 全为 0.01矩阵乘法结果理论上每个元素都等于WA × 0.01因此无需额外实现参考内核即可完成正确性校验——这也是示例选择常量初始化矩阵的巧妙之处。程序最后依次执行cuMemFree、cuModuleUnload、cuCtxDestroy完成资源清理。七、设备端内核共享内存分块实现内核实现位于 matrixMul_kernel.cu是一个模板化的分块tiling内核template int block_size, typename size_type __device__ void matrixMul(float *C, float *A, float *B, size_type wA, size_type wB)核心思路是每个线程块负责输出矩阵 C 的一个block_size × block_size子块通过循环把 A、B 的对应子矩阵加载进共享内存__shared__ float As[block_size][block_size]与Bs[...]每次迭代用两条__syncthreads()保证“加载完成”与“计算完成”的屏障同步最后累加得到Csub并写回全局内存。循环内还使用了#pragma unroll提示编译器展开内层 k 循环。文件末尾通过extern C导出三个不同块尺寸的全局内核extern C __global__ void matrixMul_bs8_64bit(float *C, float *A, float *B, size_t wA, size_t wB) { matrixMul8, size_t(C, A, B, wA, wB); } extern C __global__ void matrixMul_bs16_64bit(...) { matrixMul16, size_t(...); } extern C __global__ void matrixMul_bs32_64bit(...) { matrixMul32, size_t(...); }这三个名字正是主机端cuModuleGetFunction按字符串查找的目标且参数类型size_t64 位与主机端Matrix_Width_A/B的size_t严格对应。注意主机端并没有直接#include这个.cu文件内核代码完全通过 fatbin 在运行时加载这正是 Driver API 编程与 Runtime API 编程在构建与链接模型上的根本差异。八、构建方式如何在 CMake 中生成 fatbinCMakeLists.txt 展示了 Driver API 示例特有的构建流程核心是把设备代码单独编译为 fatbin再链接主机程序声明 CUDA 架构并构造 gencode 参数set(CMAKE_CUDA_ARCHITECTURES 75 80 86 87 89 90 100 110 120) # 对每个架构生成 -gencodearchcompute_XX,codesm_XX用自定义命令把 matrixMul_kernel.cu 编译为matrixMul_kernel64.fatbinadd_custom_command( OUTPUT ${CUDA_FATBIN_FILE} COMMAND ${CMAKE_CUDA_COMPILER} ${INCLUDES} ${ALL_CCFLAGS} -Wno-deprecated-gpu-targets ${GENCODE_FLAGS} -o ${CUDA_FATBIN_FILE} -fatbin ${CUDA_KERNEL_SOURCE} DEPENDS ${CUDA_KERNEL_SOURCE} ...)建立依赖关系确保可执行文件在 fatbin 生成之后才链接并通过add_dependencies绑定add_executable(matrixMulDrv matrixMulDrv.cpp) add_custom_target(generate_fatbin_matmulDrv ALL DEPENDS ${CUDA_FATBIN_FILE}) add_dependencies(matrixMulDrv generate_fatbin_matmulDrv)链接驱动库target_link_libraries(matrixMulDrv PUBLIC CUDA::cuda_driver)同时 CMake 还开启了CUDA_SEPARABLE_COMPILATION、cxx_std_17 / cuda_std_17并在非调试构建中追加-lineinfo以便调试工具使用。整体构建可通过仓库根目录的 CMake 体系cpp/0_Introduction/CMakeLists.txt→ 顶层 CMakeLists.txt递归完成。九、运行与验证构建成功后运行./matrixMulDrvLinux或matrixMulDrv.exeWindows典型输出流程为 Using CUDA Device [0]: GPU 名称 GPU Device has SM X.Y compute capability Total amount of global memory: xxxx bytes initCUDA loading module: matrixMul_kernel64.fatbin 路径 findModulePath found file at ... 32 block size selected Processing time: 0.xx (ms) Checking computed result for correctness: Result PASS NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.几点使用提示可通过-deviceid指定 GPU 设备不指定时自动选择估算性能最高的设备Result PASS表示逐元素误差均在1e-5阈值内若出现Result FAIL会打印第一个出错元素的位置、实际值与参考值README 与程序输出均提醒示例仅用于教学演示而非性能基准测试实际耗时可能因 GPU Boost 等机制而波动若运行时提示找不到 fatbin请确认matrixMul_kernel64.fatbin已由构建步骤生成且位于可执行文件可搜索到的路径下findFatbinPath依赖sdkFindFilePath进行路径搜索。十、总结matrixMulDrv是一个把经典矩阵乘法与CUDA Driver API结合的入门级示例它完整展示了驱动式编程的六要素设备枚举与选择findCudaDeviceDRV、上下文管理cuCtxCreate/cuCtxDestroy、模块加载cuModuleLoadData/cuModuleGetFunction、占用率驱动的块尺寸选择cuOccupancyMaxPotentialBlockSize、统一内核启动cuLaunchKernel的简单与高级两种参数传递以及显式内存管理cuMemAlloc/cuMemcpyHtoD/cuMemcpyDtoH/cuMemFree。作为同一目录下的对照示例vectorAddDrv 用 Driver API 实现向量加法simpleDrvRuntime 展示驱动与运行时 API 混用而 matrixMul_nvrtc 与 matrixMulDynlinkJIT 则进一步把内核代码改为运行时编译/链接。理解本示例的调用链之后再阅读这些进阶示例会顺畅得多。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考