企业官网建设流程全解析

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 数量理解硬件占用率与延迟隐藏的关系。

示例概述与核心思想

官方 README（cpp/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 there's 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 的墙钟时间，无法区分单个 block；
• clock()设备端时钟：返回的是每个 SM（流多处理器）内部的时钟计数器（clock cycles），粒度细到线程/块级别，适合剖析 kernel 内部各阶段的耗时分布。

本仓库中还存在姊妹示例 clock_nvrtc，它用 libNVRTC 在运行时编译同一套 kernel 逻辑，属于同一主题的另一种实现路径，可对比学习。

Kernel 实现：归约 + 块级计时

核心 kernel 为timedReduction（clock.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(); }

关键点拆解

1. 计时点：块内tid == 0的线程在归约开始前记录timer[bid] = clock()，归约结束后记录timer[bid + gridDim.x] = clock()。通过bid与bid + gridDim.x两个区段区分起止时刻，恰好对应主机端数组timer[NUM_BLOCKS * 2]的布局。
2. 动态共享内存：声明为extern __shared__ float shared[]，实际大小由启动配置的第三个参数指定（见下文sizeof(float) * 2 * NUM_THREADS），这是 CUDA 动态共享内存的标准用法；被注释掉的__shared__ float shared[2 * blockDim.x]提示了静态版本的等价写法。
3. 共享内存归约：每个线程先搬运两份数据到共享内存（shared[tid]与shared[tid + blockDim.x]），随后执行经典的"步长折半"归约循环：每一轮比较相距d的两个元素，取较小者保留，d依次除以 2，最终shared[0]即为整块输入的最小值。循环内通过__syncthreads()保证线程同步，避免数据竞争。
4. 结果写回：output[bid] = shared[0]将每块的最小值写回全局内存。

注意：该 kernel 假设输入数据量恰好为2 * blockDim.x（即 512 个 float），每个块处理一段独立的输入区间，因此块与块之间天然互不依赖，这也是它能独立计时的前提。

线程块/线程规模与共享内存配置

源码中定义了两个关键宏（clock.cu）：

#define NUM_BLOCKS 64 #define NUM_THREADS 256

• NUM_THREADS = 256：每个 block 的线程数；
• NUM_BLOCKS = 64：网格中的 block 总数；
• 输入数组大小NUM_THREADS * 2 = 512个 float；输出与计时数组分别为NUM_BLOCKS与NUM_BLOCKS * 2个元素。

启动 kernel 时通过第三个配置参数传入动态共享内存字节数：

timedReduction<<<NUM_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 → 平均时钟数）：

其背后的硬件原理（源码注释原文要点）：

• 少于 16 个块时，部分 SM 处于空闲状态；
• 超过 16 个块时所有 SM 都被使用，但每个 SM 只有一个 block，无法隐藏访存延迟；
• 超过 32 个块后，耗时随块数近似线性增长，说明硬件已通过多 block 并发充分隐藏了延迟。

因此，修改NUM_BLOCKS与NUM_THREADS是理解"如何让硬件保持忙碌"（keep the hardware busy）的最佳实验手段。

主机端流程与 CUDA Runtime API

main函数（clock.cu）完整演示了 Runtime API 的标准四步：

1. 设备选择：调用findCudaDevice(argc, argv)自动挑选性能最佳的 CUDA 设备；
2. 内存分配：cudaMalloc分别分配输入、输出与计时缓冲；
3. 数据搬运：cudaMemcpy将主机端输入拷贝到设备（cudaMemcpyHostToDevice）；
4. 结果回拷与释放：cudaMemcpy将计时数组回拷到主机（cudaMemcpyDeviceToHost），随后cudaFree释放三段设备内存。

对应 README 中列出的 CUDA Runtime API：cudaMalloc、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 的设备选择策略

findCudaDevice（Common/helper_cuda.h）的行为值得说明：

• 若命令行带--device=N参数，则使用gpuDeviceInit(N)初始化指定设备，参数非法（N < 0）或初始化失败都会报错退出；
• 否则调用gpuGetMaxGflopsDeviceId()自动选择计算性能（Gflops/s）最高的设备，并通过cudaDeviceGetAttribute读取计算能力（compute capability），打印设备名称与主次版本号。

因此运行示例时可通过./clock --device=0指定 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_DEBUG=True则改用-G开启 cuda-gdb 设备端调试（注意会显著影响性能）；
• 编译标准为cxx_std_17与cuda_std_17，并开启CUDA_SEPARABLE_COMPILATION；
• 头文件搜索路径指向仓库根目录的 Common 目录，其中包含本示例用到的 helper_cuda.h 与 helper_functions.h。

Linux 构建命令（依据仓库根 README.md）

mkdir build && cd build cmake .. make -j$(nproc)

构建产物位于build/0_Introduction/clock/下，运行：

./clock

若机器有多个 GPU，可显式指定：./clock --device=0。

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 Toolkit（README 的 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示例虽然短小，却浓缩了三个可迁移到真实性能工程的技能点：

1. 块级计时范式：timer[bid]起、timer[bid + gridDim.x]止的双区段设备内存布局，规避了块间无同步导致的计时错乱；
2. 动态共享内存 + 归约：extern __shared__声明与启动参数配合的完整写法；
3. 占用率实验方法论：通过调节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