CUDA中%%globaltimer寄存器使用及__nanosleep计时异常排查

问题描述 #

尝试在CUDA内核中使用%%globaltimer寄存器测量耗时，编写了如下代码：

#define NS_PER_S 1000000000

__global__ void sleepKernel() {
    uint64_t start, end;
    uint64_t sleepTime = 5 * NS_PER_S;     // Sleep for 5 seconds

    if (threadIdx.x == 0) {
        // Record start time
        asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(start));

        // Sleep for 5 seconds
        __nanosleep(sleepTime);

        // Record end time
        asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(end));

        // Calculate and print the elapsed time in nanoseconds and milliseconds
        uint64_t elapsedNs = end - start;
        double elapsedMs = (double)elapsedNs / 1000000.0;
        printf("Slept for %llu nanoseconds (%.3f milliseconds)\n", elapsedNs, elapsedMs);
    }
}

调用内核后输出远小于预期的5秒：

slept for 73728 nanoseconds (0.074 milliseconds)
slept for 471040 nanoseconds (0.471 milliseconds)

即使修正了整数溢出问题，现象仍未改变。

问题根源与解决方案 #

1. __nanosleep的参数并非纳秒，而是GPU时钟周期数 #

CUDA的__nanosleep函数名具有误导性，它实际接收的参数是GPU核心的时钟周期数，而非纳秒数。你传入的5 * NS_PER_S（5e9）被当作周期数处理，而主流GPU的核心频率通常在1-2GHz之间，5e9周期对应的时间本应是2.5-5秒，但你的输出异常短暂，是因为：

• 若未显式转换类型，5 * NS_PER_S会先以32位整数计算导致溢出，最终传入的是截断后的值；
• 即使修正了类型，参数含义错误仍会导致实际等待时间与预期不符。

正确做法是先获取GPU核心频率，再计算对应5秒的周期数：

// 主机端获取GPU频率
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
double gpuFreqHz = prop.clockRate * 1000.0; // clockRate单位为kHz，转换为Hz

// 计算5秒对应的时钟周期数
uint64_t sleepCycles = static_cast<uint64_t>(5.0 * gpuFreqHz);

将此sleepCycles传入内核中的__nanosleep。

2. %%globaltimer寄存器的值是时钟周期数，不是纳秒 #

%%globaltimer记录的是GPU的时钟周期计数，而非纳秒。你直接将end - start当作纳秒计算是错误的，需要将周期数转换为纳秒：

// 内核中修正计时计算
double elapsedNs = static_cast<double>(end - start) * (1000000000.0 / gpuFreqHz);
double elapsedMs = elapsedNs / 1000000.0;

注意需要将GPU频率传递到内核中，可以通过内核参数传入。

修正后的完整代码示例 #

#include <cstdio>
#include <cuda_runtime.h>

__global__ void sleepKernel(uint64_t sleepCycles, double gpuFreqHz) {
    uint64_t start, end;

    if (threadIdx.x == 0) {
        asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(start));
        __nanosleep(sleepCycles);
        asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(end));

        double elapsedNs = static_cast<double>(end - start) * (1000000000.0 / gpuFreqHz);
        double elapsedMs = elapsedNs / 1000000.0;
        printf("Slept for %.0f nanoseconds (%.3f milliseconds)\n", elapsedNs, elapsedMs);
    }
}

int main() {
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);
    double gpuFreqHz = prop.clockRate * 1000.0;
    uint64_t sleepCycles = static_cast<uint64_t>(5.0 * gpuFreqHz);

    sleepKernel<<<1, 1>>>(sleepCycles, gpuFreqHz);
    cudaDeviceSynchronize();

    return 0;
}

额外注意事项 #

• 确保编译时使用支持__nanosleep的CUDA版本和GPU架构（该函数仅在Compute Capability 3.0及以上支持）；
• %%globaltimer是全局计时器，所有SM共享，若有其他线程或内核同时运行，计时结果会包含其他任务的耗时，建议在无其他负载的环境下测试。

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