CUDA clock 示例深度解析在 Kernel 内部精确测量线程块执行耗时cuda-samples【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读本文围绕 cuda-samples 仓库中cpp/0_Introduction/clock示例展开讲解如何利用 CUDA 内建clock()函数在 kernel 内部精确测量线程块block级别的执行耗时。由于 GPU 上的 block 是并行且乱序执行的且块与块之间没有同步机制本示例采用每块各取一次时钟的策略将计时样本写入设备内存后回拷主机端统计。读完本文你将掌握clock()在 kernel 内计时的完整写法、动态共享内存归约reduction的实现以及如何通过改变 block/thread 数量理解硬件占用率与延迟隐藏的关系。示例概述与核心思想官方 READMEcpp/0_Introduction/clock/README.md对该示例的定位只有一句话This example shows how to use the clock function to measure the performance of block of threads of a kernel accurately——即用clock()精确测量 kernel 中线程块的性能。关键点在于精确二字背后的原因。源码注释clock.cu明确说明Blocks are executed in parallel and out of order. Since theres no synchronization mechanism between blocks, we measure the clock once for each block. The clock samples are written to device memory.也就是说GPU 上的多个 block 是并行、乱序调度的块之间不存在全局同步原语因此无法用全局统一计时点来度量单个块的耗时正确做法是让每个 block 在自己执行的开始与结束时刻分别调用clock()把两个时间戳写入设备内存最后由主机端做差值统计。与其他计时方式的区别主机端事件计时cudaEvent度量的是整个 kernel 的墙钟时间无法区分单个 blockclock()设备端时钟返回的是每个 SM流多处理器内部的时钟计数器clock cycles粒度细到线程/块级别适合剖析 kernel 内部各阶段的耗时分布。本仓库中还存在姊妹示例 clock_nvrtc它用 libNVRTC 在运行时编译同一套 kernel 逻辑属于同一主题的另一种实现路径可对比学习。Kernel 实现归约 块级计时核心 kernel 为timedReductionclock.cu它同时完成两件事一是做一次标准的并行归约求最小值二是记录每个 block 完成归约所消耗的时钟数。__global__ static void timedReduction(const float *input, float *output, clock_t *timer) { // __shared__ float shared[2 * blockDim.x]; extern __shared__ float shared[]; const int tid threadIdx.x; const int bid blockIdx.x; if (tid 0) timer[bid] clock(); // Copy input. shared[tid] input[tid]; shared[tid blockDim.x] input[tid blockDim.x]; // Perform reduction to find minimum. for (int d blockDim.x; d 0; d / 2) { __syncthreads(); if (tid d) { float f0 shared[tid]; float f1 shared[tid d]; if (f1 f0) { shared[tid] f1; } } } // Write result. if (tid 0) output[bid] shared[0]; __syncthreads(); if (tid 0) timer[bid gridDim.x] clock(); }关键点拆解计时点块内tid 0的线程在归约开始前记录timer[bid] clock()归约结束后记录timer[bid gridDim.x] clock()。通过bid与bid gridDim.x两个区段区分起止时刻恰好对应主机端数组timer[NUM_BLOCKS * 2]的布局。动态共享内存声明为extern __shared__ float shared[]实际大小由启动配置的第三个参数指定见下文sizeof(float) * 2 * NUM_THREADS这是 CUDA 动态共享内存的标准用法被注释掉的__shared__ float shared[2 * blockDim.x]提示了静态版本的等价写法。共享内存归约每个线程先搬运两份数据到共享内存shared[tid]与shared[tid blockDim.x]随后执行经典的步长折半归约循环每一轮比较相距d的两个元素取较小者保留d依次除以 2最终shared[0]即为整块输入的最小值。循环内通过__syncthreads()保证线程同步避免数据竞争。结果写回output[bid] shared[0]将每块的最小值写回全局内存。注意该 kernel 假设输入数据量恰好为2 * blockDim.x即 512 个 float每个块处理一段独立的输入区间因此块与块之间天然互不依赖这也是它能独立计时的前提。线程块/线程规模与共享内存配置源码中定义了两个关键宏clock.cu#define NUM_BLOCKS 64 #define NUM_THREADS 256NUM_THREADS 256每个 block 的线程数NUM_BLOCKS 64网格中的 block 总数输入数组大小NUM_THREADS * 2 512个 float输出与计时数组分别为NUM_BLOCKS与NUM_BLOCKS * 2个元素。启动 kernel 时通过第三个配置参数传入动态共享内存字节数timedReductionNUM_BLOCKS, NUM_THREADS, sizeof(float) * 2 * NUM_THREADS(dinput, doutput, dtimer);即2 * 256 * 4 2048字节每线程 2 个 float符合 kernel 内shared[tid]与shared[tid blockDim.x]的写入需求。源码注释还给出了早期 G80 架构上的实测数据blocks → 平均时钟数blocksclocks1309683232163364324615649981其背后的硬件原理源码注释原文要点少于 16 个块时部分 SM 处于空闲状态超过 16 个块时所有 SM 都被使用但每个 SM 只有一个 block无法隐藏访存延迟超过 32 个块后耗时随块数近似线性增长说明硬件已通过多 block 并发充分隐藏了延迟。因此修改NUM_BLOCKS与NUM_THREADS是理解如何让硬件保持忙碌keep the hardware busy的最佳实验手段。主机端流程与 CUDA Runtime APImain函数clock.cu完整演示了 Runtime API 的标准四步设备选择调用findCudaDevice(argc, argv)自动挑选性能最佳的 CUDA 设备内存分配cudaMalloc分别分配输入、输出与计时缓冲数据搬运cudaMemcpy将主机端输入拷贝到设备cudaMemcpyHostToDevice结果回拷与释放cudaMemcpy将计时数组回拷到主机cudaMemcpyDeviceToHost随后cudaFree释放三段设备内存。对应 README 中列出的 CUDA Runtime APIcudaMalloc、cudaMemcpy、cudaFree见 README.md。所有调用都包在checkCudaErrors宏中定义于 Common/helper_cuda.h出错即打印错误信息并终止这是 cuda-samples 统一采用的健壮性写法。统计平均耗时的逻辑位于回拷之后long double avgElapsedClocks 0; for (int i 0; i NUM_BLOCKS; i) { avgElapsedClocks (long double)(timer[i NUM_BLOCKS] - timer[i]); } avgElapsedClocks avgElapsedClocks / NUM_BLOCKS; printf(Average clocks/block %Lf\n, avgElapsedClocks);即对每个块做结束时钟 − 开始时钟累加后除以块数得到平均耗时单位是 GPU 时钟周期数而非秒。程序正常运行结束时输出形如Average clocks/block xxx的结果并返回EXIT_SUCCESS。findCudaDevice 的设备选择策略findCudaDeviceCommon/helper_cuda.h的行为值得说明若命令行带--deviceN参数则使用gpuDeviceInit(N)初始化指定设备参数非法N 0或初始化失败都会报错退出否则调用gpuGetMaxGflopsDeviceId()自动选择计算性能Gflops/s最高的设备并通过cudaDeviceGetAttribute读取计算能力compute capability打印设备名称与主次版本号。因此运行示例时可通过./clock --device0指定 GPU。构建与运行编译配置要点示例自带独立的 CMakeLists.txt可直接单独构建依赖find_package(CUDAToolkit REQUIRED)并要求 CMake ≥ 3.20通过set(CMAKE_CUDA_ARCHITECTURES 75 80 86 87 89 90 100 110 120)声明目标 GPU 架构对应 README 支持的 SM 5.0~9.0 之外的现代架构集合具体到本仓库当前版本默认追加-lineinfo编译选项以支持调试工具的行号信息若定义ENABLE_CUDA_DEBUGTrue则改用-G开启 cuda-gdb 设备端调试注意会显著影响性能编译标准为cxx_std_17与cuda_std_17并开启CUDA_SEPARABLE_COMPILATION头文件搜索路径指向仓库根目录的 Common 目录其中包含本示例用到的 helper_cuda.h 与 helper_functions.h。Linux 构建命令依据仓库根 README.mdmkdir build cd build cmake .. make -j$(nproc)构建产物位于build/0_Introduction/clock/下运行./clock若机器有多个 GPU可显式指定./clock --device0。Windows 构建在 Visual Studio 的x64 Native Tools Command Prompt for VS中执行mkdir build cd build cmake .. -G Visual Studio 16 2019 -A x64随后用 Visual Studio 打开生成的CUDA_Samples.sln选择配置后按 F7 构建再从输出目录运行clock.exe。环境前提安装与平台匹配的 CUDA ToolkitREADME 的 Prerequisites 一节明确要求预先下载安装 CUDA Toolkit支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l支持的 SM 架构README 列表SM 5.0 / 5.2 / 5.3 / 6.0 / 6.1 / 7.0 / 7.2 / 7.5 / 8.0 / 8.6 / 8.7 / 8.9 / 9.0。注意虽然 README 声明的兼容面覆盖 SM 5.0 起但当前仓库的 CMakeLists 默认只编译较新的架构列表若在旧卡上运行需自行调整CMAKE_CUDA_ARCHITECTURES。扩展与进阶路径围绕kernel 内精确计时这一主题仓库中还提供了两条可对照的深化路径NVRTC 运行时编译版本cpp/0_Introduction/clock_nvrtc以 libNVRTC 在运行时编译clock_kernel.cu适合需要动态生成/加载 kernel 的场景主机端 GPU 计时参考cpp/0_Introduction/asyncAPI 演示了用 CUDA Event 同时完成 GPU 计时与 CPU/GPU 执行重叠见 cpp/0_Introduction/README.md可与clock()的块级计时形成互补——前者回答整个 kernel 花了多久后者回答每个块花了多久。小结clock示例虽然短小却浓缩了三个可迁移到真实性能工程的技能点块级计时范式timer[bid]起、timer[bid gridDim.x]止的双区段设备内存布局规避了块间无同步导致的计时错乱动态共享内存 归约extern __shared__声明与启动参数配合的完整写法占用率实验方法论通过调节NUM_BLOCKS/NUM_THREADS观察耗时变化理解 SM 数量、块并发度与访存延迟隐藏之间的权衡。对希望深入 CUDA 性能剖析的开发者而言把clock()内嵌到自己的 kernel 中做阶段计时是比主机端事件更细粒度的第一手剖析手段本示例的源码与注释即是最好的起点教材。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考