尧图网站建设 
CUDA 矩阵乘法实战解析cuda-samples 中 matrixMul 示例的共享内存分块内核与 Runtime API 详解【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplesmatrixMul 是 cuda-samples 仓库cpp/0_Introduction/matrixMul中位于 0_Introduction 入门目录下的经典示例它完整演示了如何使用 CUDA Runtime API 实现矩阵乘法C A * B内核侧通过共享内存shared memory分块tiling实现数据复用主机侧则展示了页锁定内存、非阻塞流、CUDA 事件计时与结果校验的完整工程范式。读完本文你将掌握共享内存分块矩阵乘法内核的每一行实现原理、主机端全套 Runtime API 调用链以及如何通过命令行参数与性能指标对矩阵乘法做量化评估。示例定位为教学清晰而写而非追求极致性能根据该示例的 README 说明matrixMul 实现的矩阵乘法与 CUDA C Programming Guide 中Shared Memory共享内存章节的第二个示例完全一致。它的写作目的是清晰展示各类 CUDA 编程原则而不是提供一个通用场景下性能最优的矩阵乘法内核通过__shared__数组将矩阵分块缓存到共享内存实现同一块数据的多次复用减少对全局内存device memory的重复访问使用二维线程块thread block与二维线程thread组织计算每个线程负责输出子矩阵中的一个元素在每次迭代之间用__syncthreads()保证共享内存读写同步避免数据竞争。与此同时为了展示 GPU 在矩阵乘法上的真实性能潜力示例文档还提到该仓库配套演示了CUDA 4.0 引入的 cuBLAS 接口——即同仓库下的cpp/4_CUDA_Libraries/matrixMulCUBLAS示例通过cublasSgemm调用经过高度优化的库实现来完成同样的乘法。因此 matrixMul 与 matrixMulCUBLAS 一教一用互为对照。该示例的核心 Key Concepts 为CUDA Runtime API与Linear Algebra线性代数。支持范围与环境前提README 明确列出的支持矩阵如下支持的 SM 架构SM 5.0、SM 5.2、SM 5.3、SM 6.0、SM 6.1、SM 7.0、SM 7.2、SM 7.5、SM 8.0、SM 8.6、SM 8.7、SM 8.9、SM 9.0支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l、aarch64。前置条件下载并安装与你的平台对应的 CUDA Toolkit 说明当前版本支持 CUDA Toolkit 13.3。构建与运行matrixMul 使用 CMake 构建其 CMakeLists.txt 的关键配置如下最低要求 CMake 3.20工程声明了C CXX CUDA三种语言通过find_package(CUDAToolkit REQUIRED)定位 CUDA 工具链默认编译架构列表为75 80 86 87 89 90 100 110 120对应 SM 7.5 ~ SM 12.0默认附加-Wno-deprecated-gpu-targets编译选项若开启ENABLE_CUDA_DEBUG则追加-G启用 cuda-gdb 片上调试会显著影响性能否则追加-lineinfo为性能调试工具提供行号信息设置CUDA_SEPARABLE_COMPILATION ON启用可分离编译通过setup_samples_install()接入仓库统一的安装逻辑。在 Linux 上按根目录 README 的通用流程即可完成构建与运行也可直接在该示例子目录下操作mkdir build cd build cmake .. make -j$(nproc) ./matrixMulWindows 用户可直接用 Visual Studio 2019 16.5 的 CMake 语言服务导入仓库或在x64 Native Tools Command Prompt for VS中执行cmake .. -G Visual Studio 16 2019 -A x64命令行参数程序入口main见 matrixMul.cu通过checkCmdLineFlag/getCmdLineArgumentInt实现在 Common/helper_string.h解析如下参数参数含义默认值-devicen指定使用的 GPU 设备 IDn ≥ 0否则自动挑选最佳设备自动选择-wAWidthA矩阵 A 的宽度320-hAHeightA矩阵 A 的高度320-wBWidthB矩阵 B 的宽度640-hBHeightB矩阵 B 的高度320-help/-?打印用法说明—默认维度由block_size 32推导dimsA(5*2*32, 5*2*32) (320, 320)dimsB(5*4*32, 5*2*32) (640, 320)。程序强制约束矩阵 A 的宽度必须等于矩阵 B 的高度dimsA.x dimsB.y否则直接报错退出因为这是矩阵乘法合法的前提。共享内存分块内核逐行解析内核函数MatrixMulCUDA以BLOCK_SIZE为模板参数matrixMul.cu通过模板实例化在编译期确定分块大小示例运行时使用 16 或 32template int BLOCK_SIZE __global__ void MatrixMulCUDA(float *C, float *A, float *B, int wA, int wB) { int bx blockIdx.x; // 块在网格中的 x 索引 int by blockIdx.y; // 块在网格中的 y 索引 int tx threadIdx.x; // 线程在块内的 x 索引 int ty threadIdx.y; // 线程在块内的 y 索引 // A 中本块负责处理的首个子矩阵起始下标 int aBegin wA * BLOCK_SIZE * by; // A 中本块处理的最后一个子矩阵下标 int aEnd aBegin wA - 1; int aStep BLOCK_SIZE; // 沿 A 行方向移动的步长 int bBegin BLOCK_SIZE * bx; // B 中本块首个子矩阵起始下标 int bStep BLOCK_SIZE * wB; // 沿 B 列方向移动的步长 float Csub 0; // 该线程负责输出的 C 元素累加器 // 沿 K 维滑动遍历所有需要相乘的子矩阵对 for (int a aBegin, b bBegin; a aEnd; a aStep, b bStep) { __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; // A 的分块缓存 __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; // B 的分块缓存 // 每个线程从全局内存各搬运一个元素进共享内存 As[ty][tx] A[a wA * ty tx]; Bs[ty][tx] B[b wB * ty tx]; __syncthreads(); // 等待整个块的数据搬运完成 #pragma unroll for (int k 0; k BLOCK_SIZE; k) { Csub As[ty][k] * Bs[k][tx]; // 块内小矩阵乘法 } __syncthreads(); // 确保本块计算结束再覆盖 As/Bs } // 将累加结果写回全局内存 C int c wB * BLOCK_SIZE * by BLOCK_SIZE * bx; C[c wB * ty tx] Csub; }这段内核体现了三个核心工程手法共享内存分块复用As/Bs两个BLOCK_SIZE×BLOCK_SIZE的__shared__数组在每次外层迭代中被整体搬运一次随后被BLOCK_SIZE²个线程反复读取BLOCK_SIZE次。相比每个线程直接读全局内存共享内存的访问延迟和带宽压力都大幅下降这正是数据复用data reuse教学点的精髓。双重__syncthreads()同步第一次同步保证全块数据已写入共享内存再开始计算第二次同步保证所有线程都已消费完当前分块再搬运下一批分块从而避免块内数据竞争。需要留意__syncthreads()是块内屏障不同 block 之间天然独立、无需同步。#pragma unroll循环展开提示编译器展开内层长度为BLOCK_SIZE的累加循环减少循环开销、提升指令级并行是这类小循环上常见的性能优化手段。主机端完整执行流程MatrixMultiply函数matrixMul.cu串起了主机侧的全套 Runtime API 调用可分六个阶段理解1. 内存分配与数据初始化使用cudaMallocHost为h_A、h_B、h_C分配页锁定pinned主机内存这是配合异步拷贝cudaMemcpyAsync的前提页锁定内存可被 DMA 直接访问无需经过可分页内存的中转通过ConstantInit将 A 全部初始化为1.0fB 初始化为0.01f——这个精心选择的常量使得参考结果可解析为dimsA.x * valB便于后续校验使用cudaMalloc在设备端分配d_A、d_B、d_C结果矩阵 C 的维度为dimsC(dimsB.x, dimsA.y, 1)即 C 是wB × hA的矩阵与数学定义一致。2. 流与事件创建cudaStream_t stream; checkCudaErrors(cudaEventCreate(start)); checkCudaErrors(cudaEventCreate(stop)); checkCudaErrors(cudaStreamCreateWithFlags(stream, cudaStreamNonBlocking));cudaStreamCreateWithFlags配合cudaStreamNonBlocking创建非阻塞流该流不隐式与默认流legacy default stream同步使内核与拷贝可以与主机代码或其他流并发执行。所有异步操作拷贝、内核、事件记录都挂载在该流上。3. 数据搬运与内核启动cudaMemcpyAsync(d_A, h_A, ..., cudaMemcpyHostToDevice, stream)异步将 A、B 拷入设备执行配置为dim3 threads(block_size, block_size)、dim3 grid(dimsB.x / threads.x, dimsA.y / threads.y)即网格形状由输出矩阵 C 的维度决定每个线程块负责 C 上的一个BLOCK_SIZE×BLOCK_SIZE子块根据block_size选择实例化MatrixMulCUDA16或MatrixMulCUDA32启动内核grid, threads, 0, stream第三个参数为 0 表示不动态分配共享内存第四个参数指定所在流。4. 计时与性能统计程序先执行一次**预热warmup**内核再通过事件对 300 次迭代nIter 300整体计时checkCudaErrors(cudaEventRecord(start, stream)); for (int j 0; j nIter; j) { /* 启动内核 */ } checkCudaErrors(cudaEventRecord(stop, stream)); checkCudaErrors(cudaEventSynchronize(stop)); checkCudaErrors(cudaEventElapsedTime(msecTotal, start, stop)); float msecPerMatrixMul msecTotal / nIter; double flopsPerMatrixMul 2.0 * (double)dimsA.x * (double)dimsA.y * (double)dimsB.x; double gigaFlops (flopsPerMatrixMul * 1.0e-9f) / (msecPerMatrixMul / 1000.0f);其中2 × wA × hA × wB是矩阵乘法(hA×wA)·(wA×wB)的浮点运算次数每次乘加计 2 次 FLOP单次平均耗时换算后得到GFlop/s指标。cudaEventSynchronize会阻塞主机直到 stop 事件完成保证计时覆盖全部迭代。预热的意义在于触发驱动与硬件的初始化开销避免其污染测时结果。5. 正确性校验结果回拷后程序按相对误差公式校验matrixMul.cu|h_C[i] - ref| / |h_C[i]| / dimsA.x eps其中 eps 1e-6参考值ref dimsA.x * valBA 全 1、B 全 0.01 时 C 的元素应为wA × 0.01。任一元素超出阈值即输出Error! Matrix[xxxxx]...并判Result FAIL全部通过则输出Result PASS。eps 1e-6相当于机器零量级用于容忍浮点舍入误差。6. 资源清理cudaFreeHost释放页锁定内存cudaFree释放设备内存cudaEventDestroy销毁事件并打印性能结果与说明The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.——即示例的计时仅用于教学演示GPU Boost 动态调频等机制会使结果浮动。涉及的 CUDA Runtime API 一览README 列出的全部 Runtime API 及在示例中的用途如下API用途cudaMalloc/cudaFree设备全局内存分配与释放cudaMallocHost/cudaFreeHost页锁定主机内存分配与释放cudaMemcpyAsync流内异步主机↔设备拷贝cudaStreamCreateWithFlags创建非阻塞 CUDA 流cudaStreamSynchronize等待流内全部操作完成cudaEventCreate/cudaEventDestroy创建/销毁计时事件cudaEventRecord向流中记录事件cudaEventSynchronize阻塞至事件完成cudaEventElapsedTime计算两个事件间的毫秒耗时cudaProfilerStart/cudaProfilerStop在main中显式控制 profiler如 nvprof/Nsight的启停便于只截取核心计算区间其中所有 API 调用都被 Common/helper_cuda.h 中的checkCudaErrors宏包裹该宏在返回非成功错误码时打印文件:行号 错误码 出错 API并退出是 cuda-samples 全仓库统一采用的健壮性惯例建议在自有项目中同样遵循。设备选择则由findCudaDevice同样位于 Common/helper_cuda.h完成——它遍历所有 CUDA 设备并优先挑选计算能力最高的那块。对照cuBLAS 高性能实现README 指出本示例的另一目的是展示 CUDA 4.0 的 cuBLAS 接口。仓库中的 matrixMulCUBLAS 示例即为此而设其核心调用链为checkCudaErrors(cublasCreate(handle)); // 创建 cuBLAS 句柄 checkCudaErrors(cublasSgemm(handle, ...)); // 单精度矩阵乘法 C alpha*A*B beta*C该示例同样输出Performance %.2f GFlop/s形式的结果matrixMulCUBLAS.cpp与手写内核形成同口径对比。值得注意的细节是cuBLAS 采用**列主序column-major**存储约定源码注释matrixMulCUBLAS.cpp明确指出由于C A * B的输入顺序问题不能直接写作cublasSgemm(A, B)需要在参数上做相应转置/顺序调整——这是把行主序数据交给 cuBLAS 时最常见的坑。通过对比两个示例可以直观看到通用教学内核重在可读性而库级实现经过大量调优寄存器分块、Tensor Core 等是生产场景的更优选择。小结matrixMul 示例虽小却是一条完整的教学闭环共享内存分块内核展示了 GPU 端数据复用的核心思想主机端则示范了从页锁定内存、非阻塞流、事件计时到误差校验、资源清理的完整 CUDA Runtime API 工程范式。建议的进阶路径是先在本示例上动手修改BLOCK_SIZE与矩阵维度观察性能与正确性变化再对照 matrixMulCUBLAS 体会手写内核与库实现的差距最后借助 Nsight 等工具示例已内置cudaProfilerStart/Stop埋点深入分析内核行为。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CUDA 矩阵乘法实战解析：cuda-samples 中 matrixMul 示例的共享内存分块内核与 Runtime API 详解
CUDA 矩阵乘法实战解析cuda-samples 中 matrixMul 示例的共享内存分块内核与 Runtime API 详解【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplesmatrixMul 是 cuda-samples 仓库cpp/0_Introduction/matrixMul中位于 0_Introduction 入门目录下的经典示例它完整演示了如何使用 CUDA Runtime API 实现矩阵乘法C A * B内核侧通过共享内存shared memory分块tiling实现数据复用主机侧则展示了页锁定内存、非阻塞流、CUDA 事件计时与结果校验的完整工程范式。读完本文你将掌握共享内存分块矩阵乘法内核的每一行实现原理、主机端全套 Runtime API 调用链以及如何通过命令行参数与性能指标对矩阵乘法做量化评估。示例定位为教学清晰而写而非追求极致性能根据该示例的 README 说明matrixMul 实现的矩阵乘法与 CUDA C Programming Guide 中Shared Memory共享内存章节的第二个示例完全一致。它的写作目的是清晰展示各类 CUDA 编程原则而不是提供一个通用场景下性能最优的矩阵乘法内核通过__shared__数组将矩阵分块缓存到共享内存实现同一块数据的多次复用减少对全局内存device memory的重复访问使用二维线程块thread block与二维线程thread组织计算每个线程负责输出子矩阵中的一个元素在每次迭代之间用__syncthreads()保证共享内存读写同步避免数据竞争。与此同时为了展示 GPU 在矩阵乘法上的真实性能潜力示例文档还提到该仓库配套演示了CUDA 4.0 引入的 cuBLAS 接口——即同仓库下的cpp/4_CUDA_Libraries/matrixMulCUBLAS示例通过cublasSgemm调用经过高度优化的库实现来完成同样的乘法。因此 matrixMul 与 matrixMulCUBLAS 一教一用互为对照。该示例的核心 Key Concepts 为CUDA Runtime API与Linear Algebra线性代数。支持范围与环境前提README 明确列出的支持矩阵如下支持的 SM 架构SM 5.0、SM 5.2、SM 5.3、SM 6.0、SM 6.1、SM 7.0、SM 7.2、SM 7.5、SM 8.0、SM 8.6、SM 8.7、SM 8.9、SM 9.0支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l、aarch64。前置条件下载并安装与你的平台对应的 CUDA Toolkit 说明当前版本支持 CUDA Toolkit 13.3。构建与运行matrixMul 使用 CMake 构建其 CMakeLists.txt 的关键配置如下最低要求 CMake 3.20工程声明了C CXX CUDA三种语言通过find_package(CUDAToolkit REQUIRED)定位 CUDA 工具链默认编译架构列表为75 80 86 87 89 90 100 110 120对应 SM 7.5 ~ SM 12.0默认附加-Wno-deprecated-gpu-targets编译选项若开启ENABLE_CUDA_DEBUG则追加-G启用 cuda-gdb 片上调试会显著影响性能否则追加-lineinfo为性能调试工具提供行号信息设置CUDA_SEPARABLE_COMPILATION ON启用可分离编译通过setup_samples_install()接入仓库统一的安装逻辑。在 Linux 上按根目录 README 的通用流程即可完成构建与运行也可直接在该示例子目录下操作mkdir build cd build cmake .. make -j$(nproc) ./matrixMulWindows 用户可直接用 Visual Studio 2019 16.5 的 CMake 语言服务导入仓库或在x64 Native Tools Command Prompt for VS中执行cmake .. -G Visual Studio 16 2019 -A x64命令行参数程序入口main见 matrixMul.cu通过checkCmdLineFlag/getCmdLineArgumentInt实现在 Common/helper_string.h解析如下参数参数含义默认值-devicen指定使用的 GPU 设备 IDn ≥ 0否则自动挑选最佳设备自动选择-wAWidthA矩阵 A 的宽度320-hAHeightA矩阵 A 的高度320-wBWidthB矩阵 B 的宽度640-hBHeightB矩阵 B 的高度320-help/-?打印用法说明—默认维度由block_size 32推导dimsA(5*2*32, 5*2*32) (320, 320)dimsB(5*4*32, 5*2*32) (640, 320)。程序强制约束矩阵 A 的宽度必须等于矩阵 B 的高度dimsA.x dimsB.y否则直接报错退出因为这是矩阵乘法合法的前提。共享内存分块内核逐行解析内核函数MatrixMulCUDA以BLOCK_SIZE为模板参数matrixMul.cu通过模板实例化在编译期确定分块大小示例运行时使用 16 或 32template int BLOCK_SIZE __global__ void MatrixMulCUDA(float *C, float *A, float *B, int wA, int wB) { int bx blockIdx.x; // 块在网格中的 x 索引 int by blockIdx.y; // 块在网格中的 y 索引 int tx threadIdx.x; // 线程在块内的 x 索引 int ty threadIdx.y; // 线程在块内的 y 索引 // A 中本块负责处理的首个子矩阵起始下标 int aBegin wA * BLOCK_SIZE * by; // A 中本块处理的最后一个子矩阵下标 int aEnd aBegin wA - 1; int aStep BLOCK_SIZE; // 沿 A 行方向移动的步长 int bBegin BLOCK_SIZE * bx; // B 中本块首个子矩阵起始下标 int bStep BLOCK_SIZE * wB; // 沿 B 列方向移动的步长 float Csub 0; // 该线程负责输出的 C 元素累加器 // 沿 K 维滑动遍历所有需要相乘的子矩阵对 for (int a aBegin, b bBegin; a aEnd; a aStep, b bStep) { __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; // A 的分块缓存 __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; // B 的分块缓存 // 每个线程从全局内存各搬运一个元素进共享内存 As[ty][tx] A[a wA * ty tx]; Bs[ty][tx] B[b wB * ty tx]; __syncthreads(); // 等待整个块的数据搬运完成 #pragma unroll for (int k 0; k BLOCK_SIZE; k) { Csub As[ty][k] * Bs[k][tx]; // 块内小矩阵乘法 } __syncthreads(); // 确保本块计算结束再覆盖 As/Bs } // 将累加结果写回全局内存 C int c wB * BLOCK_SIZE * by BLOCK_SIZE * bx; C[c wB * ty tx] Csub; }这段内核体现了三个核心工程手法共享内存分块复用As/Bs两个BLOCK_SIZE×BLOCK_SIZE的__shared__数组在每次外层迭代中被整体搬运一次随后被BLOCK_SIZE²个线程反复读取BLOCK_SIZE次。相比每个线程直接读全局内存共享内存的访问延迟和带宽压力都大幅下降这正是数据复用data reuse教学点的精髓。双重__syncthreads()同步第一次同步保证全块数据已写入共享内存再开始计算第二次同步保证所有线程都已消费完当前分块再搬运下一批分块从而避免块内数据竞争。需要留意__syncthreads()是块内屏障不同 block 之间天然独立、无需同步。#pragma unroll循环展开提示编译器展开内层长度为BLOCK_SIZE的累加循环减少循环开销、提升指令级并行是这类小循环上常见的性能优化手段。主机端完整执行流程MatrixMultiply函数matrixMul.cu串起了主机侧的全套 Runtime API 调用可分六个阶段理解1. 内存分配与数据初始化使用cudaMallocHost为h_A、h_B、h_C分配页锁定pinned主机内存这是配合异步拷贝cudaMemcpyAsync的前提页锁定内存可被 DMA 直接访问无需经过可分页内存的中转通过ConstantInit将 A 全部初始化为1.0fB 初始化为0.01f——这个精心选择的常量使得参考结果可解析为dimsA.x * valB便于后续校验使用cudaMalloc在设备端分配d_A、d_B、d_C结果矩阵 C 的维度为dimsC(dimsB.x, dimsA.y, 1)即 C 是wB × hA的矩阵与数学定义一致。2. 流与事件创建cudaStream_t stream; checkCudaErrors(cudaEventCreate(start)); checkCudaErrors(cudaEventCreate(stop)); checkCudaErrors(cudaStreamCreateWithFlags(stream, cudaStreamNonBlocking));cudaStreamCreateWithFlags配合cudaStreamNonBlocking创建非阻塞流该流不隐式与默认流legacy default stream同步使内核与拷贝可以与主机代码或其他流并发执行。所有异步操作拷贝、内核、事件记录都挂载在该流上。3. 数据搬运与内核启动cudaMemcpyAsync(d_A, h_A, ..., cudaMemcpyHostToDevice, stream)异步将 A、B 拷入设备执行配置为dim3 threads(block_size, block_size)、dim3 grid(dimsB.x / threads.x, dimsA.y / threads.y)即网格形状由输出矩阵 C 的维度决定每个线程块负责 C 上的一个BLOCK_SIZE×BLOCK_SIZE子块根据block_size选择实例化MatrixMulCUDA16或MatrixMulCUDA32启动内核grid, threads, 0, stream第三个参数为 0 表示不动态分配共享内存第四个参数指定所在流。4. 计时与性能统计程序先执行一次**预热warmup**内核再通过事件对 300 次迭代nIter 300整体计时checkCudaErrors(cudaEventRecord(start, stream)); for (int j 0; j nIter; j) { /* 启动内核 */ } checkCudaErrors(cudaEventRecord(stop, stream)); checkCudaErrors(cudaEventSynchronize(stop)); checkCudaErrors(cudaEventElapsedTime(msecTotal, start, stop)); float msecPerMatrixMul msecTotal / nIter; double flopsPerMatrixMul 2.0 * (double)dimsA.x * (double)dimsA.y * (double)dimsB.x; double gigaFlops (flopsPerMatrixMul * 1.0e-9f) / (msecPerMatrixMul / 1000.0f);其中2 × wA × hA × wB是矩阵乘法(hA×wA)·(wA×wB)的浮点运算次数每次乘加计 2 次 FLOP单次平均耗时换算后得到GFlop/s指标。cudaEventSynchronize会阻塞主机直到 stop 事件完成保证计时覆盖全部迭代。预热的意义在于触发驱动与硬件的初始化开销避免其污染测时结果。5. 正确性校验结果回拷后程序按相对误差公式校验matrixMul.cu|h_C[i] - ref| / |h_C[i]| / dimsA.x eps其中 eps 1e-6参考值ref dimsA.x * valBA 全 1、B 全 0.01 时 C 的元素应为wA × 0.01。任一元素超出阈值即输出Error! Matrix[xxxxx]...并判Result FAIL全部通过则输出Result PASS。eps 1e-6相当于机器零量级用于容忍浮点舍入误差。6. 资源清理cudaFreeHost释放页锁定内存cudaFree释放设备内存cudaEventDestroy销毁事件并打印性能结果与说明The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.——即示例的计时仅用于教学演示GPU Boost 动态调频等机制会使结果浮动。涉及的 CUDA Runtime API 一览README 列出的全部 Runtime API 及在示例中的用途如下API用途cudaMalloc/cudaFree设备全局内存分配与释放cudaMallocHost/cudaFreeHost页锁定主机内存分配与释放cudaMemcpyAsync流内异步主机↔设备拷贝cudaStreamCreateWithFlags创建非阻塞 CUDA 流cudaStreamSynchronize等待流内全部操作完成cudaEventCreate/cudaEventDestroy创建/销毁计时事件cudaEventRecord向流中记录事件cudaEventSynchronize阻塞至事件完成cudaEventElapsedTime计算两个事件间的毫秒耗时cudaProfilerStart/cudaProfilerStop在main中显式控制 profiler如 nvprof/Nsight的启停便于只截取核心计算区间其中所有 API 调用都被 Common/helper_cuda.h 中的checkCudaErrors宏包裹该宏在返回非成功错误码时打印文件:行号 错误码 出错 API并退出是 cuda-samples 全仓库统一采用的健壮性惯例建议在自有项目中同样遵循。设备选择则由findCudaDevice同样位于 Common/helper_cuda.h完成——它遍历所有 CUDA 设备并优先挑选计算能力最高的那块。对照cuBLAS 高性能实现README 指出本示例的另一目的是展示 CUDA 4.0 的 cuBLAS 接口。仓库中的 matrixMulCUBLAS 示例即为此而设其核心调用链为checkCudaErrors(cublasCreate(handle)); // 创建 cuBLAS 句柄 checkCudaErrors(cublasSgemm(handle, ...)); // 单精度矩阵乘法 C alpha*A*B beta*C该示例同样输出Performance %.2f GFlop/s形式的结果matrixMulCUBLAS.cpp与手写内核形成同口径对比。值得注意的细节是cuBLAS 采用**列主序column-major**存储约定源码注释matrixMulCUBLAS.cpp明确指出由于C A * B的输入顺序问题不能直接写作cublasSgemm(A, B)需要在参数上做相应转置/顺序调整——这是把行主序数据交给 cuBLAS 时最常见的坑。通过对比两个示例可以直观看到通用教学内核重在可读性而库级实现经过大量调优寄存器分块、Tensor Core 等是生产场景的更优选择。小结matrixMul 示例虽小却是一条完整的教学闭环共享内存分块内核展示了 GPU 端数据复用的核心思想主机端则示范了从页锁定内存、非阻塞流、事件计时到误差校验、资源清理的完整 CUDA Runtime API 工程范式。建议的进阶路径是先在本示例上动手修改BLOCK_SIZE与矩阵维度观察性能与正确性变化再对照 matrixMulCUBLAS 体会手写内核与库实现的差距最后借助 Nsight 等工具示例已内置cudaProfilerStart/Stop埋点深入分析内核行为。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考