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

Question

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

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

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

原始工作应用程序的伪代码，现在转换为处理 2 gpus：

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 内核不依赖输入来开始工作。在它运行之后，它确实会 async_work_group_copy 将它生成的数据复制到一个全局缓冲区中。

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

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

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

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

这是我的分析/计时的输出。所有数字都以毫秒为单位。

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 驱动程序不关心内核中发生的事情。如果内核写入/读取或者是一个空内核，或者它只写入缓冲区的一部分。它关心 缓冲区参数标志，并确保数据在 GPU 之间保持一致，如果内核与其他内核有任何依赖关系，则阻塞内核。 GPU 到 GPU 的传输是透明进行的，而且成本可能非常高。

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

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

---

例如对于这些内核：

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);

事件将发生如下：

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

如您所见，GPU2 必须等待第一个完成，另外还要等待所有参数复制回来。 （可以是 5s 吗？也许，取决于大小）

---

但是使用正确的方法：

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

事件将发生如下：

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

并行（因为没有依赖关系）

1. 内存 A、D 将被复制到 GPU2（如果它不存在的话）。已分配内存 E。
2. GPU2 将运行内核。