11

我想usleep()在 CUDA 内核中调用类似的东西。基本目标是让所有 GPU 内核休眠或忙碌等待数毫秒——这是我想为 CUDA 应用程序执行的一些健全性检查的一部分。我这样做的尝试如下：

#include <unistd.h>
#include <stdio.h>
#include <cuda.h>
#include <sys/time.h>

__global__ void gpu_uSleep(useconds_t wait_time_in_ms)
{
    usleep(wait_time_in_ms);
}

int main(void)
{
    //input parameters -- arbitrary
    //   TODO: set these exactly for full occupancy
    int m = 16;
    int n = 16;
    int block1D = 16;
    dim3 block(block1D, block1D);
    dim3 grid(m/block1D, n/block1D);

    useconds_t wait_time_in_ms = 1000;

    //execute the kernel
    gpu_uSleep<<< grid, block >>>(wait_time_in_ms);
    cudaDeviceSynchronize();

    return 0;
}

当我尝试使用 NVCC 编译它时出现以下错误：

error: calling a host function("usleep") from a __device__/__global__ 
       function("gpu_uSleep") is not allowed

显然，我不允许usleep()在内核中使用主机功能。什么是一个很好的选择呢？

4

3 回答 3

22

您可以在clock() 或clock64() 上旋转。CUDA SDK concurrentKernels 示例执行以下操作：

__global__ void clock_block(clock_t *d_o, clock_t clock_count)
{
    clock_t start_clock = clock();
    clock_t clock_offset = 0;
    while (clock_offset < clock_count)
    {
        clock_offset = clock() - start_clock;
    }
     d_o[0] = clock_offset;
}

我推荐使用clock64()。clock() 和 clock64() 是循环的，因此您必须使用 cudaDeviceProperties() 查询频率。频率可以是动态的，因此很难保证准确的自旋循环。

于 2012-06-27T00:56:06.243 回答
11

您可以使用读取为clock().

要等待至少 10,000 个时钟周期：

clock_t start = clock();
clock_t now;
for (;;) {
  now = clock();
  clock_t cycles = now > start ? now - start : now + (0xffffffff - start);
  if (cycles >= 10000) {
    break;
  }
}
// Stored "now" in global memory here to prevent the
// compiler from optimizing away the entire loop.
*global_now = now;

注意：这是未经测试的。@Pedro 从这个答案中借用了处理溢出的代码。有关如何工作的详细信息，请参阅他的答案和 CUDA C 编程指南 4.2 中的 B.10 节clock()。还有一个clock64()命令。

于 2012-06-27T00:53:57.213 回答
4

对于最新版本的 CUDA，以及具有 Compute Capability 7.0 或更高版本（Volta、Turing、Ampere 等）的设备，您可以使用__nanosleep()原语：

void __nanosleep(unsigned ns);

这消除了旧答案中建议的忙碌睡眠的需要。

于 2021-01-28T12:36:48.020 回答