我怎么知道内核是否同时执行？

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