使用多个 GPU 时 OpenCL 内核启动延迟？

Question

我有一个我设计的应用程序运行在带有 OpenCL 的 AMD GPU 上。昨天终于得到了应用程序运行ning 并且没有错误（哈哈），目标是单个 GPU。现在应用程序可以运行了，是时候将其扩展到多个 GPU 了。

阅读大量有关如何设置的信息。我们正在使用单上下文、多队列方法。

我拉出设备列表，选择 2 个 GPU 并创建一个包含它们的上下文，然后创建一个也包含两个设备的 BuildProgram。创建两个单独的队列。

原始可运行应用程序的伪代码，现已转换为处理 2 个 GPU：

context = clCreateContext(0, 2, device_list, NULL, NULL, &ret);
for(x = 0; x < 2; x++)
  queue[x] = clCreateCommandQueue(context, device_list[x], ENABLE_PROFILING, &ret);
clBuildProgram(program, 2, device_list, options, NULL, NULL);

create kernels..

run...
for(outer_loop = 0; outer_loop < 10; outer_loop++) {
  clEnqueueNDRangeKernel(queue[0], kernel_init, offset, &event[0]);
  clEnqueueNDRangeKernel(queue[1], kernel_init, different_offset, &event[1]);
  clFinish(queue[0]);
  clFinish(queue[1]);

  get profiling data and printf results
}

代码基本上就是这样。在循环之前设置参数并完成写入 - init 内核不依赖输入来开始工作。在它运行s 之后，它确实将其生成的数据 async_work_group_copy 发送到全局缓冲区。

现在，在我修改 2 个 GPU 的代码之前，内核运行在 27 毫秒内（对于每个循环）

在我修改代码后，如果我注释掉两个内核运行中的一个或另一个（EnqueueNDRangeKernel 和关联的 clFinish），它们都将在 27 毫秒内运行。

如果我在两个 GPU 上并行运行代码到运行，我会得到非常奇怪的行为。

循环中的第一个运行，它们分别执行大约 37-42 毫秒。我可以稍微放慢速度，因为我完成了两倍的工作。但是在第一个运行之后，一个或另一个内核将运行在排队和启动之间通常会有 4-5 秒的延迟。

这是我的 profiling/timing 的输出。所有数字均以毫秒为单位。

Q0: til sub:  8.8542  til start: 9.8594 til fin: 47.3749
Q1: til sub:  0.0132  til start: 13.4089 til fin: 39.2364

Q0: til sub:  0.0072  til start: 0.2310 til fin: 37.1187
Q1: til sub:  0.0122  til start: 4152.4638 til fin: 4727.1146

Q0: til sub:  0.0302  til start: 488.6218 til fin: 5049.7233
Q1: til sub:  0.0179  til start: 5023.9310 til fin: 5049.7762

Q0: til sub:  0.0190  til start: 2.0987 til fin: 39.4356
Q1: til sub:  0.0164  til start: 3996.2654 til fin: 4571.5866

Q0: til sub:  0.0284  til start: 488.5751 til fin: 5046.3555
Q1: til sub:  0.0176  til start: 5020.5919 til fin: 5046.4382

我运行安装此机器的机器有 5 个 GPU。无论我使用哪两个，两个 GPU 中的一个（它并不总是相同的）在启动时会有 4-5 秒的延迟。使用单个 GPU - 没有延迟。

这可能是什么原因造成的？任何的想法？我没有阻止 - clFinish 只是为了获取分析信息。即使它被阻止也不会延迟 5 秒。

另外 - 我认为内核正在执行的全局写入可能是延迟的一部分。我评论了写出来。没有。没有变化。

其实我加了一个return；作为内核的第一行 - 所以它什么都不做。 40 毫秒下降到 0.25，但 5 秒的延迟仍然存在。

Answer 1

OpenCL 驱动程序不关心内核中发生的事情。如果内核 writes/reads 或者是空内核，或者它只写入缓冲区的一部分。它关心 缓冲区参数标志 ，并确保数据在 GPU 之间是一致的，如果它们对其他内核有任何依赖性，则阻止内核。 GPU 到 GPU 的传输是透明的，而且成本很高。

当使用多个 GPU 时，必须认真对待隐藏的数据复制和同步，因为这通常是主要瓶颈。

如果您的内核可以运行并行（因为 GPU1 处理的数据与 GPU2 处理的数据不同，依此类推...），那么您应该为每个内核创建不同的缓冲区GPU。或者，如果数据相同，请正确设置类型 CL_READ_ONLY/CL_WRITE_ONLY，以确保正确的 OpenCL 行为。最少 copy/consistency 次操作。

例如，对于这些内核：

kernel Sum(read_only A, read_only B, write_only C);
kernel Sum_bad(read_write A, read_write B, write_only C);

如果您使用单个 GPU，两者的行为将完全相同，因为所有内存都驻留在同一个 GPU 中。但是使用多个 GPU 会导致可怕的问题，例如：

Queue 1/GPU 1: Sum_Bad(A,B,C);
Queue 2/GPU 2: Sum_Bad(A,D,E);

事件将发生如下：

内存 A、B 将被复制到 GPU1 内存（如果它不在那里）。在 GPU1 中分配的 C 内存。
GPU 1 将运行内核。
内存A将从GPU1复制到GPU2。内存 D 将被复制到 GPU2。已分配内存 E。
GPU2 将运行内核。

如你所见，GPU2必须等待第一个完成，并且还要等待所有参数复制回来。（可以是 5s 吗？也许吧，取决于尺寸）

但是使用正确的方法：

Queue 1/GPU 1: Sum(A,B,C);
Queue 2/GPU 2: Sum(A,D,E);

事件将发生如下：

内存 A、B 将被复制到 GPU1 内存（如果它不在那里）。在 GPU1 中分配的 C 内存。
GPU 1 将运行内核。

并行（因为没有依赖）

内存 A、D 将被复制到 GPU2（如果它不在那里）。已分配内存 E。
GPU2 将运行内核。

使用多个 GPU 时 OpenCL 内核启动延迟？

OpenCL kernel START delays when using multiple GPUs?

opencl

gpu

multiple-gpu