如何实现opencl内核管道

Question

我正在使用 OpenCl 处理我的项目。为了提高我的算法的性能，是否可以流水线化单个内核？如果内核包含许多步骤，比如说 A、B、C，我希望 A 在完成其部分后立即接受新数据并将其传递给 B。我可以在它们之间创建通道，但我不知道该怎么做详细地。

我可以在 .cl 文件中写入 A、B、C（3 个内核）吗？但是如何排队NDRange？ 我正在使用 Altera SDK 进行 FPGA HPC 开发。 谢谢。

Answer 1

管道可以通过使用多个与通道连接的内核来实现。在所有内核同时运行的情况下，数据从一个内核传输到另一个内核：

Pipeline example from Intel FPGA OpenCL SDK Programming Guide

这种管道的非常基本的例子是：

channel int foo_bar_channel;
channel float bar_baz_channel;

__kernel void foo(__global int* in) {
  for (int i = 0; i < 1024; ++i) {
    int value = in[i];
    value = clamp(value, 0, 255);                 // do some work
    write_channel_altera(foo_bar_channel, value); // send data to the next kernel
  }
}

__kernel void bar() {
  for (int i = 0; i < 1024; ++i) {
    int value = read_channel_altera(foo_bar_channel); // take data from foo
    float fvalue = (float) value;
    write_channel_altera(bar_baz_channel, value); // send data to the next kernel
  }
}

__kernel void baz(__global int* out) {
  for (int i = 0; i < 1024; ++i) {n
    float value = read_channel_altera(bar_baz_channel);
    float s = sin(value);
    out[i] = s;                                  // write result in the end
  }
}

您可以将所有内核写入一个 .cl 文件中，或者使用不同的文件，然后将它们#include 到一个主 .cl 文件中。

我们希望我们所有的内核同时运行，这样它们就可以接受来自彼此的数据。由于只支持顺序命令队列，我们​​必须为每个内核使用不同的队列：

cl_queue foo_queue = clCreateCommandQueue(...);
cl_queue bar_queue = clCreateCommandQueue(...);
cl_queue baz_queue = clCreateCommandQueue(...);

clEnqueueTask(foo_queue, foo_kernel);
clEnqueueTask(bar_queue, bar_kernel);
clEnqueueTask(baz_queue, baz_kernel);

clFinish(baz_queue); // last kernel in our pipeline

与 GPU 的 OpenCL 编程不同，我们依赖于数据流水线，因此 NDRange 内核不会给我们带来任何好处。使用单个工作项内核而不是 NDRange 内核，因此我们使用 clEnqueueTask 函数将它们排入队列。额外的内核属性（reqd_work_group_size）可用于标记单个工作项内核，给编译器一些优化空间。

查看英特尔 FPGA SDK for OpenCL 编程指南，了解有关通道和内核属性的更多信息（具体而言，第 1.6.4 节实施面向 OpenCL 通道扩展的英特尔 FPGA SDK）：

https://www.altera.com/en_US/pdfs/literature/hb/opencl-sdk/aocl_programming_guide.pdf