如何在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>
#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 文档中可发现的内容。