我如何知道内核是否同时执行？答案

【问题标题】：How do I know if the kernels are executing concurrently?我如何知道内核是否同时执行？
【发布时间】：2012-07-30 15:03:39
【问题描述】：

我有一个带有 CC 3.0 的 GPU，所以它应该支持 16 个并发内核。我通过循环 clEnqueueNDRangeKernel 10 次来启动 10 个内核。我如何知道内核是同时执行的？

我想到的一种方法是获取 NDRangeKernel 语句之前和之后的时间。我可能不得不使用事件来确保内核的执行已经完成。但我还是觉得循环会顺序启动内核。谁能帮帮我..

【问题讨论】：

如果你将所有内核放在同一个命令队列中，它们确实会按顺序执行（这就是它被称为队列的原因）。恐怕测量每个单独内核的时间和总执行时间是实际测量内核是否并行执行的唯一可能方法。
我已将所有内核放在不同的命令队列中。
您的假设是，在 Fermi 架构上支持 16 个 CUDA 流的 CUDA Compute Capability 3.0 可作为 OpenCL 功能使用？ NVida 文档中是否有任何内容支持这一假设？如果您希望使用 OpenCL Device Fission，可以使用 clGetDeviceInfo () 查询此扩展作为支持的扩展功能
我在OpenCL Programming Guide Section 3.2.2 中读到了对并发内核的支持

标签： opencl

【解决方案1】：

您可以通过使用C Framework for OpenCL 来避免其他答案中建议的所有样板代码（顺便说一句是正确的），这大大简化了此任务，并为您提供有关 OpenCL 事件（内核执行、数据传输）的详细信息等），包括专门用于重叠执行所述事件的表格和绘图。

我开发这个库是为了简化其他答案中描述的过程。可以看一个基本的用法示例here。

【讨论】：

【解决方案2】：

要确定您的内核执行是否重叠，您必须对它们进行概要分析。这需要几个步骤：

1。创建命令队列

仅当使用属性CL_QUEUE_PROFILING_ENABLE 创建命令队列时才会收集分析数据：

cl_command_queue queues[10];
for (int i = 0; i < 10; ++i) {
  queues[i] = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE,
                                   &errcode);
}

2。确保所有内核同时启动

您的假设是正确的，即 CPU 按顺序排列内核。但是，您可以创建单个用户事件并将其添加到所有内核的等待列表中。这会导致内核在用户事件完成之前不开始运行：

// Create the user event
cl_event user_event = clCreateUserEvent(context, &errcode);

// Reserve space for kernel events
cl_event kernel_events[10];

// Enqueue kernels
for (int i = 0; i < 10; ++i) {
  clEnqueueNDRangeKernel(queues[i], kernel, work_dim, global_work_offset,
                         global_work_size, 1, &user_event, &kernel_events[i]);
}

// Start all kernels by completing the user event
clSetUserEventStatus(user_event, CL_COMPLETE);

3。获取分析时间

最后，我们可以收集内核事件的时序信息：

// Block until all kernels have run to completion
clWaitForEvents(10, kernel_events);

for (int i = 0; i < 10; ++i) {
  cl_ulong start;
  clGetEventProfilingInfo(kernel_event[i], CL_PROFILING_COMMAND_START,
                          sizeof(start), &start, NULL);
  cl_ulong end;
  clGetEventProfilingInfo(kernel_event[i], CL_PROFILING_COMMAND_END,
                          sizeof(end), &end, NULL);
  printf("Event %d: start=%llu, end=%llu", i, start, end);
}

4。分析输出

现在您已经知道了所有内核运行的开始和结束时间，您可以检查重叠（手动或以编程方式）。输出单位是纳秒。但是请注意，设备计时器仅精确到某个分辨率。您可以使用以下方式查询分辨率：

size_t resolution;
clGetDeviceInfo(device, CL_DEVICE_PROFILING_TIMER_RESOLUTION,
                sizeof(resolution), &resolution, NULL);

FWIW，我在带有 CC 2.0（应该支持并发内核）的 NVIDIA 设备上进行了尝试，并观察到内核是按顺序运行的。

【讨论】：

非常感谢。我会试试这个，让你知道结果。就在昨天，当我在谷歌上搜索时，我发现有一种叫做设备裂变的东西，它使我们能够在单个设备上创建子设备，并且我们可以在不同的子设备上执行内核。 AMD 设备肯定支持设备裂变，但仍不确定 NVIDIA 设备是否支持它。我还在读它..
嗨，我最近使用这些步骤来测试我的内核是否同时运行。我正在使用我认为支持并发内核执行的 AMD A10 APU。但是分析结果告诉我，内核一个接一个地执行，而不是同时执行。我需要启用一些特定选项才能使并发内核执行工作吗？我在以下链接中发布了我的代码和输出：stackoverflow.com/questions/35341061/…
这个方法实际上是行不通的。首先，队列默认按顺序创建（如本例所示），但即使您传递了乱序标志，运行时也可能决定按顺序运行您的内核。其次，OpenCL 命令队列可能会也可能不会映射到硬件队列，具体取决于您的运行时间，因此即使使用单独的队列也可能会为您提供顺序执行。
@Mokosha 我认为您误解了该示例的目的。它显示了一种确定 if 运行时在将它们放在多个队列时同时执行一些内核的方法，我相信这是 OP 想要知道的。它从未缩进显示如何实际制作/强制运行时执行此操作。关于您提到的CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 标志：根据规范，这只影响同一队列中命令的重新排序，而不影响多个队列之间的交互。

【解决方案3】：

是的，按照您的建议，尝试使用事件，并分析所有 QUEUED、SUBMIT、START、END 值。这些应该是“设备时间”中的绝对值，您可以查看不同内核的处理（开始到结束）是否重叠。

【讨论】：