如何在CUDA中使用%%globaltimer寄存器？

Question

这个问题的答案建议使用%%globaltimer寄存器来测量 CUDA 内核中经过的时间。我决定尝试一下：

#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);
    }
}

但是，当我调用这个内核时，输出如下：

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

两者都远小于 5 秒。我错过了什么吗？

编辑：做：

    uint64_t sleepTime = 5 * (uint64_t)NS_PER_S;     // Sleep for 5 seconds

有点帮助（防止整数溢出），但这还不够

Answer 1

我认为这里混乱的根源是由于__nanosleep(). 如果我们首先阅读文档和相关的ptx 文档，我们将观察到值得注意的事情：

1. 传递给它的参数是unsigned，它是一个 32 位数量。当量子为纳秒时，无法精确表达 5 秒，参数为unsigned。传递 64 位值尚未正式可用，并且不会改变这一点。但等等，还有更多。
2. 使用“大约”一词。在 PTX 文档中，我们看到它所基于的函数实际上提供了范围从零（！）到您指定值两倍的延迟。这可能不是大多数人对延迟函数的期望。
3. 至少在PTX中，该功能不能用于提供超过1毫秒的延迟。与 5 秒相差甚远，也远低于您可以要求的 32 位数量。

一些编译和 PTX 分析将向您展示 CUDA C++ 内在函数或多或少地直接使用 PTX 函数，因此它具有相同的奇怪特征。我无法回答“为什么会这样？” 问题或“它有什么用？” 问题。

考虑到这一点，由于您的问题的标题至少与 相关globaltimer，而不是__nanosleep，我们可以很容易地验证这globaltimer似乎与广告中的大致相同：

# cat t118.cu
#define NS_PER_S 1000000000ULL
#include <time.h>
#include <sys/time.h>
#include <iostream>
#include <cstdint>
#define USECPSEC 1000000ULL

unsigned long long dtime_usec(unsigned long long start=0){

  timeval tv;
  gettimeofday(&tv, 0);
  return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}

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

    if (threadIdx.x == 0) {
        // Record start time
        asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(start));
        end = start;
        // Sleep for 5 seconds
        while (end < (start + sleepTime))
          asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(end));
    }
}


int main(){
        sleepKernel<<<1,1>>>(100);
        cudaDeviceSynchronize();
        unsigned long long dt = dtime_usec(0);
        sleepKernel<<<1,1>>>();
        cudaDeviceSynchronize();
        dt = dtime_usec(dt);
        std::cout << "elapsed: " << dt << "us" << std::endl;
}
# nvcc -o t118 t118.cu -arch=sm_89
# ./t118
elapsed: 5000023us
#

笔记：

1. 我通常不推荐 PTX 分析来理解。然而，它在这里显示 CUDA C++ 内在函数和我们可以从中推断行为的底层 PTX 函数之间的联系是有用且足够的。

2. 我已经向 NVIDIA 提交了一个内部错误 (3608779)，以更新 CUDA C++ 内在函数 ( ) 的文档，__nanosleep()以更好地反映 PTX 文档中可发现的内容。

3. 第一次内核启动是一个“热身”，以吸收各种 CUDA 启动开销。如果您想了解这意味着什么，只需将其删除即可。