定时 CUDA 操作答案

【问题标题】：Timing CUDA operations定时 CUDA 操作
【发布时间】：2011-12-14 03:59:50
【问题描述】：

我需要为 CUDA 内核执行计时。最佳实践指南说我们可以在 Windows 中使用事件或标准计时函数，如 clock()。我的问题是使用这两个函数会给我一个完全不同的结果。事实上，事件给出的结果与实践中的实际速度相比，似乎是巨大的。

我真正需要的是能够通过首先在较小的数据集上运行计算的简化版本来预测计算的运行时间。不幸的是，这个基准测试的结果完全不切实际，要么过于乐观（clock()），要么过于悲观（事件）。

【问题讨论】：

你在启动内核之后和之前的时间（结束）与时钟同步吗？
你的意思是我之前和之后是否有 cudaThreadSynchronize() 调用？是的，我愿意。
是的，我就是这个意思
顺便说一句。 CUDA 计时以毫秒（如果您使用视觉分析器，则为微秒）返回。以防万一。

标签： c cuda benchmarking

【解决方案1】：

如果您想测量 GPU 时间，您几乎必须使用事件。在nvidia forums here 上有一个关于应用程序计时的注意事项的精彩讨论帖。

【讨论】：

【解决方案2】：

您可以使用非常适合您的目的的计算视觉分析器。它测量每个 cuda 函数的时间并告诉你调用了多少次。

【讨论】：

谢谢，但我需要以编程方式进行这些测量。
@Programmer：分析器还完全序列化 API，并增加延迟，因为它需要额外的主机设备传输来收集配置文件计数器输出。它对很多事情都很有用，但准确的执行时间不是其中之一。
@talonmies：分析器完全序列化 API 是什么意思？ ——
@Programmer：CUDA API 自然是异步的（内核启动、流、某些类型的内存传输）。当您在分析器中运行程序时，它们都会变成串行的。如果您的代码将内存复制与内核执行重叠，则在分析时这些代码将是串行的。在 Fermi 上，多个同时内核执行也在分析期间被禁用。
如果您使用命令行分析器（而不是可视分析器），我认为您仍然可以通过编程方式进行。但正如 talonmies 所说，它序列化了 API 调用。因此，您得到的是以阻塞方式执行所有 API 调用。而且还有一点额外的开销用于读取计数器。

【解决方案3】：

你可以按照以下方式做一些事情：

#include <sys/time.h>

struct timeval t1, t2;

gettimeofday(&t1, 0);

kernel_call<<<dimGrid, dimBlock, 0>>>();

HANDLE_ERROR(cudaThreadSynchronize();)

gettimeofday(&t2, 0);

double time = (1000000.0*(t2.tv_sec-t1.tv_sec) + t2.tv_usec-t1.tv_usec)/1000.0;

printf("Time to generate:  %3.1f ms \n", time);

或：

float time;
cudaEvent_t start, stop;

HANDLE_ERROR( cudaEventCreate(&start) );
HANDLE_ERROR( cudaEventCreate(&stop) );
HANDLE_ERROR( cudaEventRecord(start, 0) );

kernel_call<<<dimGrid, dimBlock, 0>>>();

HANDLE_ERROR( cudaEventRecord(stop, 0) );
HANDLE_ERROR( cudaEventSynchronize(stop) );
HANDLE_ERROR( cudaEventElapsedTime(&time, start, stop) );

printf("Time to generate:  %3.1f ms \n", time);

【讨论】：

@Programmer 这是一个他没有定义的函数或宏，用于处理 cuda 函数调用返回的错误。您应该进行错误处理，但为简单起见，此处可以省略。
@程序员，没错，SDK中有一些用于错误处理的有用宏
@fbielejec SDK 中的错误处理功能只是为了使示例尽可能简单以用于教育。遇到错误时调用 exit() 并不是处理错误的最佳方式！
请注意，1e6 us = 1 s，所以在第一个示例中，时间以秒为单位，而不是毫秒。
关于 HANDLE_ERROR，请参阅 stackoverflow.com/q/14038589/2778484 并查看 CUDA 示例中的 helper_cuda.h，其中有一个名为 getLastCudaError 的宏。

【解决方案4】：

您的问题已经得到了满意的答复。

我已经构建了用于计时 C/C++ 以及 CUDA 操作的类，并希望与其他人分享，希望它们对下一个用户有所帮助。您只需要将下面报告的4 文件添加到您的项目中，并将#include 这两个头文件添加为

// --- Timing includes
#include "TimingCPU.h"
#include "TimingGPU.cuh"

这两个类可以如下使用。

定时 CPU 部分

TimingCPU timer_CPU;

timer_CPU.StartCounter();
CPU perations to be timed
std::cout << "CPU Timing = " << timer_CPU.GetCounter() << " ms" << std::endl;

计时 GPU 部分

TimingGPU timer_GPU;
timer_GPU.StartCounter();
GPU perations to be timed
std::cout << "GPU Timing = " << timer_GPU.GetCounter() << " ms" << std::endl;

在这两种情况下，时间都以毫秒为单位。另外，这两个类在linux或windows下都可以使用。

这里是4 文件：

TimingCPU.cpp

/**************/
/* TIMING CPU */
/**************/

#include "TimingCPU.h"

#ifdef __linux__

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

    TimingCPU::TimingCPU(): cur_time_(0) { StartCounter(); }

    TimingCPU::~TimingCPU() { }

    void TimingCPU::StartCounter()
    {
        struct timeval time;
        if(gettimeofday( &time, 0 )) return;
        cur_time_ = 1000000 * time.tv_sec + time.tv_usec;
    }

    double TimingCPU::GetCounter()
    {
        struct timeval time;
        if(gettimeofday( &time, 0 )) return -1;

        long cur_time = 1000000 * time.tv_sec + time.tv_usec;
        double sec = (cur_time - cur_time_) / 1000000.0;
        if(sec < 0) sec += 86400;
        cur_time_ = cur_time;

        return 1000.*sec;
    }

#elif _WIN32 || _WIN64
    #include <windows.h>
    #include <iostream>

    struct PrivateTimingCPU {
        double  PCFreq;
        __int64 CounterStart;
    };

    // --- Default constructor
    TimingCPU::TimingCPU() { privateTimingCPU = new PrivateTimingCPU; (*privateTimingCPU).PCFreq = 0.0; (*privateTimingCPU).CounterStart = 0; }

    // --- Default destructor
    TimingCPU::~TimingCPU() { }

    // --- Starts the timing
    void TimingCPU::StartCounter()
    {
        LARGE_INTEGER li;
        if(!QueryPerformanceFrequency(&li)) std::cout << "QueryPerformanceFrequency failed!\n";

        (*privateTimingCPU).PCFreq = double(li.QuadPart)/1000.0;

        QueryPerformanceCounter(&li);
        (*privateTimingCPU).CounterStart = li.QuadPart;
    }

    // --- Gets the timing counter in ms
    double TimingCPU::GetCounter()
    {
        LARGE_INTEGER li;
        QueryPerformanceCounter(&li);
        return double(li.QuadPart-(*privateTimingCPU).CounterStart)/(*privateTimingCPU).PCFreq;
    }
#endif

TimingCPU.h

// 1 micro-second accuracy
// Returns the time in seconds

#ifndef __TIMINGCPU_H__
#define __TIMINGCPU_H__

#ifdef __linux__

    class TimingCPU {

        private:
            long cur_time_;

        public:

            TimingCPU();

            ~TimingCPU();

            void StartCounter();

            double GetCounter();
    };

#elif _WIN32 || _WIN64

    struct PrivateTimingCPU;

    class TimingCPU
    {
        private:
            PrivateTimingCPU *privateTimingCPU;

        public:

            TimingCPU();

            ~TimingCPU();

            void StartCounter();

            double GetCounter();

    }; // TimingCPU class

#endif

#endif

TimingGPU.cu

/**************/
/* TIMING GPU */
/**************/

#include "TimingGPU.cuh"

#include <cuda.h>
#include <cuda_runtime.h>

struct PrivateTimingGPU {
    cudaEvent_t     start;
    cudaEvent_t     stop;
};

// default constructor
TimingGPU::TimingGPU() { privateTimingGPU = new PrivateTimingGPU; }

// default destructor
TimingGPU::~TimingGPU() { }

void TimingGPU::StartCounter()
{
    cudaEventCreate(&((*privateTimingGPU).start));
    cudaEventCreate(&((*privateTimingGPU).stop));
    cudaEventRecord((*privateTimingGPU).start,0);
}

void TimingGPU::StartCounterFlags()
{
    int eventflags = cudaEventBlockingSync;

    cudaEventCreateWithFlags(&((*privateTimingGPU).start),eventflags);
    cudaEventCreateWithFlags(&((*privateTimingGPU).stop),eventflags);
    cudaEventRecord((*privateTimingGPU).start,0);
}

// Gets the counter in ms
float TimingGPU::GetCounter()
{
    float   time;
    cudaEventRecord((*privateTimingGPU).stop, 0);
    cudaEventSynchronize((*privateTimingGPU).stop);
    cudaEventElapsedTime(&time,(*privateTimingGPU).start,(*privateTimingGPU).stop);
    return time;
}

TimingGPU.cuh

#ifndef __TIMING_CUH__
#define __TIMING_CUH__

/**************/
/* TIMING GPU */
/**************/

// Events are a part of CUDA API and provide a system independent way to measure execution times on CUDA devices with approximately 0.5
// microsecond precision.

struct PrivateTimingGPU;

class TimingGPU
{
    private:
        PrivateTimingGPU *privateTimingGPU;

    public:

        TimingGPU();

        ~TimingGPU();

        void StartCounter();
        void StartCounterFlags();

        float GetCounter();

}; // TimingCPU class

#endif

【讨论】：

效果很好！除了上述两个包含之外，我还必须包含 #include "TimingCPU.cpp" 和 #include "TimingGPU.cu"。

【解决方案5】：

有一个开箱即用的GpuTimer 结构供使用：

#ifndef __GPU_TIMER_H__
#define __GPU_TIMER_H__

struct GpuTimer
{
      cudaEvent_t start;
      cudaEvent_t stop;

      GpuTimer()
      {
            cudaEventCreate(&start);
            cudaEventCreate(&stop);
      }

      ~GpuTimer()
      {
            cudaEventDestroy(start);
            cudaEventDestroy(stop);
      }

      void Start()
      {
            cudaEventRecord(start, 0);
      }

      void Stop()
      {
            cudaEventRecord(stop, 0);
      }

      float Elapsed()
      {
            float elapsed;
            cudaEventSynchronize(stop);
            cudaEventElapsedTime(&elapsed, start, stop);
            return elapsed;
      }
};

#endif  /* __GPU_TIMER_H__ */

【讨论】：