如何阻止nvcc重排clock指令以获得准确的CUDA内核计时？

在CUDA 12.5环境下，nvcc编译器会自动重排clock64()计时语句，导致计时结果严重失真——比如以下示例代码中，测得的耗时仅15周期，但实际延迟接近400周期。

问题复现代码 #

#include <cuda.h>
#include <stdio.h>

__device__ int dedupe(int a) {
  constexpr auto all = -1u;
  const auto start = clock64();
  const auto dupemask = __match_any_sync(all, a);
  const auto leader = __clz(dupemask);
  const auto end = clock64();
  const auto time = int(end - start);
  printf("tid: %i, dupemask: $%x, leader: %i, time: %i\n", threadIdx.x, dupemask, leader, time);
  return leader;
}

__global__ void dostuff() {
    const auto tid = threadIdx.x;
    const auto leader = dedupe(tid);
}

int main() {
  dostuff<<<1, 32>>>();
  return cudaDeviceSynchronize();
}

编译后的汇编显示end = clock64()被提前重排到__clz指令之前：

CS2R R2, SR_CLOCKLO         // start = clock
 S2R R8, SR_TID.X            // tid = threadIdx.x
 MATCH.ANY R9, R8            // __match_any_sync
 CS2R R4, SR_CLOCKLO         // end = clock <<-- 被重排
 FLO.U32 R10, R9             // __clz
 ... 

注意：仅用volatile asm替换clock64()无法解决问题，如下代码依旧会被重排：

__device__ uint64_t myclock() {
  uint64_t result;
  asm volatile ("mov.u64 %0, %%clock64;" : "=l"(result) :: "memory");
  return result;
}

解决方法 #

方法1：增强自定义clock函数的约束 #

在asm语句中添加"cc"约束（告知编译器指令会影响执行流），配合"memory"约束，彻底阻止编译器重排clock调用：

__device__ uint64_t myclock() {
    uint64_t result;
    // 添加"cc"和"memory"约束，强制编译器不重排该指令
    asm volatile("mov.u64 %0, %%clock64;" : "=l"(result) : : "memory", "cc");
    return result;
}

方法2：插入同步屏障指令 #

在计时点前后插入__syncwarp()（同warp内同步）或__threadfence_block()（块内内存屏障），强制编译器和硬件保留指令执行顺序：

__device__ int dedupe(int a) {
    constexpr auto all = -1u;
    __syncwarp(); // 确保前置指令完成，阻止start被后移
    const auto start = clock64();
    const auto dupemask = __match_any_sync(all, a);
    const auto leader = __clz(dupemask);
    __syncwarp(); // 确保目标指令完成，阻止end被前移
    const auto end = clock64();
    const auto time = int(end - start);
    printf("tid: %i, dupemask: $%x, leader: %i, time: %i\n", threadIdx.x, dupemask, leader, time);
    return leader;
}

原理说明 #

nvcc优化器默认认为clock64()是无副作用的纯函数，因此会对其进行指令重排。通过添加"cc"约束或插入同步屏障，明确告知编译器：clock调用与周围指令存在执行顺序依赖，不能随意调整；同时同步指令也会约束硬件层面的指令调度，确保计时区间完全覆盖目标代码的执行周期。

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