如何实现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所有内核文件或使用不同的文件，然后＃将它们包括到主.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 

不同的OpenCL编程GPU，我们依靠数据流水线，所以NDRange内核不会给我们带来任何好处。使用单个工作项内核而不是NDRange内核，所以我们使用clEnqueueTask函数对它们进行排队。额外的内核属性（reqd_work_group_size）可用于标记单个工作项内核，为编译器提供一些优化空间。

检查有关渠道和内核属性（具体而言，第1.6.4节针对OpenCL渠道推广实施英特尔FPGA SDK）的更多信息，英特尔FPGA SDK针对OpenCL编程指南：

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