使用Intel® FPGA SDK for OpenCL™ 离线编译器，不需要调整kernel代码便可以将其最佳的适应于固定的硬件设备，而是离线编译器会根据kernel的要求自适应调整硬件的结构。
通常来说，应该先优化针对单个计算单元的kernel，之后累哦通过增加计算单元数量来拓展硬件以填充FPGA其余的部分，从而提升性能。Kernel的使用面积与硬件编译所需要的时间有关，因此为了避免硬件编译时间过长，首先要专注于优化kernel在单个计算单元上的性能。
要优化kernel的性能，主要包括数据处理以及内存访问优化。
a. 通过SDK的channel 或pipe来传输数据。为了提高kernel之间的数据传输效率，在kernel程序中使用channel通道拓展。 如果想利用通道功能，又想使用其他SDK运行kernel，则使用OpenCL pipes。
b. 展开循环。
c. 优化浮点运算。对于浮点操作，可以手动引导SDK的离线编译器进行优化，从而在硬件中创建更有效的pipeline结构并减少总体硬件使用率。
d. 分配对齐的内存。再分配与FPGA进行数据传输的主机端存储器时，存储器至少是64字节对齐的。
e. 使用或不用Padding来对齐结构。
f. 保持向量元素的相似结构。如果更新了向量的一个元素，那么更新这个向量的所有元素。
g. 避免指针混淆。尽量在指针参数中插入strict关键字。
h. 避免开销大的函数/功能。有些函数在FPGA中实现开销很大，可能会减低kernel的性能，或是需要大量硬件来实现。
i. 避免依赖于work-item id的后向分支。避免在kernel中包括任何与工作项ID相关的向后分支（即，循环中发生的分支），因为这会降低性能。
1、通过SDK的channel 或pipe来传输数据
为了提高kernel之间的数据传输效率，在kernel程序中使用channel通道拓展。 如果想利用通道功能，又想使用其他SDK运行kernel，则使用OpenCL pipes。
有时，FPGA到global memory全局存储器带宽会限制内核之间的数据传输效率。 理论上FPGA到global memory全局存储器的最大带宽根据目标定制平台和板上可用的全局存储器bank的数量而变化。 要确定主板的理论最大带宽，要参考主板的文档。
实际上，kernel无法实现最大可用全局内存带宽的100％利用率。 利用率级别取决于算法的访问模式。
如果全局内存带宽是我们使用OpenCL内核的性能限制，首先尝试将算法分解为多个较小的kernel。 其次，通过在内核之间实现SDK的channel或OpenCL的pipe进行数据传输来消除一些全局内存访问。
(1) Channel与pipe的特性
a. Default Behavior
Channel默认行为是阻塞的，而pipe的默认行为是非阻塞的(nonblocking)。
b. 多个OpenCL内核的并发执行
可以同时执行多个OpenCL内核。 要启用并发执行，要修改主机代码以实例化多个命令队列。 每个同时执行的kernel内核都与一个单独的命令队列关联。
pipe的特别注意事项：Intel SDK 中OpenCL的pipe是允许在其他的OpenCL SDK上兼容的，但不能最大化kernel内核吞吐量。OpenCL 2.0中要求在进行pipe读取前先进性pipe写入，以免kernel在空pipe中读取数据，因此kernel无法同时运行。由于Intel SDK支持并发执行，可以修改主机应用程序以及kernel程序来实现并发执行，从而提高吞吐量。但不能将kernel移植到其他SDK上。
要启用并发执行包含pipe的内核，需要将内核代码中的depth属性替换为blocking属性（即__attribute __（（blocking）））。 blocking属性在read_pipe和write_pipe函数调用时引入blocking行为。 调用点将阻止内核执行，直到管道的另一端准备好为止。
如果同时将blocking属性和depth属性添加到内核，则当管道为空时，read_pipe仅调用一个块，而当管道为满时，write_pipe仅调用一个块。 blocking行为会导致内核之间的隐式同步，从而使得内核之间互锁。
c. 隐式内核Kernel同步
通过blocking channel调用以及blocking pipe调用来隐式同步kernel。
channel int c0; __kernel void producer (__global int * in_buf) { for (int i = 0; i < 10; i++) { write_channel_intel (c0, in_buf[i]); } } __kernel void consumer (__global int * ret_buf) { for (int i = 0; i < 10; i++) { ret_buf[i] = read_channel_intel(c0); } }