opencl中的时钟（）答案

【问题标题】：clock() in openclopencl中的时钟（）
【发布时间】：2012-01-13 10:57:16
【问题描述】：

我知道 CUDA 中有函数 clock() ，您可以在其中放入内核代码并查询 GPU 时间。但是我想知道OpenCL中是否存在这样的事情？有没有办法在 OpenCL 中查询 GPU 时间？（我使用的是 NVIDIA 的工具包）。

【问题讨论】：

标签： cuda opencl gpu

【解决方案1】：

没有直接查询时钟周期的 OpenCL 方法。但是，OpenCL 确实有一个分析机制，可以在计算设备上公开增量计数器。通过比较有序事件之间的差异，可以测量经过的时间。请参阅 clGetEventProfilingInfo。

【讨论】：

我认为这会给整个内核的执行时间。但我想查询特定数量的工作项或工作组的时间。有没有办法做到这一点？
不是直接的，但是您可以将具有特定数量的工作项或工作组的内核排入队列并对其进行分析。
是的，我明白你的意思。但事情有点复杂。工作项的时间根据总共有多少工作项而有所不同。例如，如果存在资源争用，那么只执行一个工作组将比同时拥有多个工作组更快。所以我猜想分析（计时）单个工作项，甚至工作组还没有在 OpenCL 中实现。

【解决方案2】：

试试这个（当然只适用于 NVidia OpenCL）：

uint clock_time()
{
    uint clock_time;
    asm("mov.u32 %0, %%clock;" : "=r"(clock_time));
    return clock_time;
}

【讨论】：

更新：由于 ptx 2.0 版也有一个 %clock64 寄存器

【解决方案3】：

NVIDIA OpenCL SDK 有一个示例 Using Inline PTX with OpenCL。时钟寄存器可通过内联 PTX 作为特殊寄存器 %clock 访问。 %clock 在PTX: Parallel Thread Execution ISA 手册中有描述。您应该可以将 %%laneid 替换为 %%clock。

我从未使用 OpenCL 测试过这个，但在 CUDA 中使用它。

请注意，编译器可能会重新排序或删除寄存器读取。

【讨论】：

是的，这是迄今为止我想说的最可能的解决方案。谢谢！

【解决方案4】：

仅供其他人向她寻求帮助：使用 OpenCL 分析内核运行时的简短介绍

启用分析模式：

cmdQueue = clCreateCommandQueue(context, *devices, CL_QUEUE_PROFILING_ENABLE, &err);

分析内核：

cl_event prof_event; 
clEnqueueNDRangeKernel(cmdQueue, kernel, 1 , 0, globalWorkSize, NULL, 0, NULL, &prof_event);

读取分析数据：

cl_ulong ev_start_time=(cl_ulong)0;     
cl_ulong ev_end_time=(cl_ulong)0;   

clFinish(cmdQueue);
err = clWaitForEvents(1, &prof_event);
err |= clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &ev_start_time, NULL);
err |= clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &ev_end_time, NULL);

计算内核执行时间：

float run_time_gpu = (float)(ev_end_time - ev_start_time)/1000; // in usec

目前尚无法对单个工作项/工作组进行分析。您可以设置 globalWorkSize = localWorkSize 进行分析。那么你就只有一个工作组了。

顺便说一句：单个工作项（一些工作项）的分析不是很有帮助。只有一些工作项，您将无法隐藏内存延迟和导致无意义测量的开销。

【讨论】：

NVIDIA Nsight Visual Studio Edition 执行上述对您的应用程序透明的任务，并在表格视图和时间轴上显示信息。该工具适用于大多数 OpenCL 平台，因为它适用于 ICD 层。 OpenCL 不支持同步 CPU 和 GPU 计时器的机制，因此时间线在非 NVIDIA 平台上可能存在同步问题（倾斜、漂移）。其他 OpenCL 供应商也有类似的工具。

【解决方案5】：

在 NVIDIA 上，您可以使用以下内容：

typedef unsigned long uint64_t; // if you haven't done so earlier
inline uint64_t n_nv_Clock()
{
    uint64_t n_clock;
    asm volatile("mov.u64 %0, %%clock64;" : "=l" (n_clock)); // make sure the compiler will not reorder this
    return n_clock;
}

volatile 关键字告诉优化器你是认真的，不希望它被移走/优化。这是在PTX 和 e.g. 中这样做的标准方式。在gcc。

请注意，这会返回时钟，而不是纳秒。您需要查询设备时钟频率（使用clGetDeviceInfo(device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(freq), &freq, 0))）。另请注意，在较旧的设备上，有两个频率（如果计算在这种情况下无关的内存频率，则为三个）：设备时钟和着色器时钟。你想要的是着色器时钟。

使用 64 位版本的寄存器，您无需担心溢出，因为它通常需要数百年时间。另一方面，32 位版本可能会经常溢出（您仍然可以恢复结果 - 除非它溢出两次）。

【讨论】：