opencl - 我如何知道内核是否正在并发执行？

Question

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

我想到的一种方法是获取 NDRangeKernel 语句前后的时间。我可能不得不使用事件来确保内核的执行已经完成。但我仍然觉得循环会依次启动内核。有人可以帮我吗..

Answer 1

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

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 设备上尝试了这个，并观察到内核是按顺序运行的。