如何获取NVIDIA GPU指令时钟周期？测试异常问题咨询

为什么1000次mul指令仅测得6个周期？ #

你的测试代码没有设置数据依赖，GPU的超标量执行单元和流水线调度会把这些独立的mul指令并行执行，甚至在多个周期内同时发射多条指令，完全隐藏了单指令的延迟。Turing架构（2080Ti）的SM每个周期可以发射多条整数运算指令，循环展开后1000次独立的mul操作会被拆解到多个执行单元并行处理，实际占用的周期数远小于指令数量，所以clock64返回的是整个并行序列的总周期，而非单指令延迟的累加。

要测量单指令的真实延迟，必须让每条指令的输入依赖前一条指令的输出，强制指令串行执行。比如修改测试代码：

uint32_t a = 0x12345678, res = 1;
uint64_t start = clock64();
#pragma unroll
for (int i = 0; i < 1000; i++) {
    res = __mullo_u32(a, res); // 每次mul的输入依赖上一次的结果
}
uint64_t end = clock64();
uint64_t cycles = end - start;
// 单指令延迟 = cycles / 1000

这样GPU线程必须等待前一条mul完成才能执行下一条，测得的周期数除以循环次数才是单条mul.lo.u32的真实延迟。

其他获取GPU指令时钟周期的方法 #

• CUDA NSight Systems Profiler：用nsys profile命令启动程序，选择指令级分析（Instruction-level profiling），可以直接查看每条PTX/SASS指令的执行周期、吞吐量，以及SM资源占用情况。
• 编写带数据依赖的微基准测试：针对不同指令（add、mul、ld/st）分别构造串行依赖的测试循环，通过clock64测量总周期后计算单指令延迟，注意要在测量前后添加__syncthreads()同步warp内线程，避免误差。
• 查看官方架构文档：直接从官方文档获取指令的标称延迟和吞吐量，无需自行测试。

相关官方文档 #

• CUDA C++ Programming Guide：附录G（GPU Features and Specifications）列出了各架构的指令延迟、吞吐量，比如Turing架构的整数乘法延迟为3-4个周期，吞吐量为每个SM周期最多执行32次整数乘法（全warp）。
• NVIDIA GPU架构白皮书：比如《Turing Architecture Whitepaper》，详细说明SM的执行单元结构、指令调度机制，以及各类指令的性能参数。
• CUDA Profiler User Guide：指导如何使用NSight Systems/Compute进行指令级性能分析，定位指令执行的瓶颈。

内容的提问来源于stack exchange，提问作者sorfkc