OpenCL设计优化（基于Intel FPGA SDK for OpenCL）

Posted 2020-12-18 zhuzhudong

1、首先了解Intel FPGA SDK for OpenCL实现OpenCL的设计组件，包括:

kernels, global memory interconnect, local memory, loops 以及channels

(1) Kernels

　　Loops一般是Kernel优化的重点，尤其是nested loops。

　　OpenCL系统中每个kernel是通过一系列block表示的。Block主要由三部分：输入或循环输入节点，一组指令以及一个分支节点。若block中没有分支则没有输入与分支节点，输入或循环输入节点根据分支在block中的起始位置决定变量的初始值。block的余下部分包含可停顿以及不可停顿指令以及clusters。优化好的模型应该包含最少数量的可停顿指令，例如I/O或内存访问。

　　在block中不可停顿的指令被分组成多个clusters，以减小可停顿指令的握手开销。cluster包含入口/出口节点，且cluster是不可停顿的。可以在cluster的出口节点找到出口FIFO信息，在这种情况下，分支节点会通知下一个要跳转到的block。

　　只有当Kernel中不使用get_global_id()以及get_local_id()这种内置的work-item时，SDK才会将其编译为single work-item kernel，否则将其编译为NDRange kernel。

　　而offline compiler不能将NDRange中的loops进行pineline，但这些loops可以同时接受多个work-item。一个kernel可能包括多个循环，且每个都有嵌套循环。若将每个外部loop的nested loop的迭代总数制表，kernel的吞吐量一般会因此降低。而要有效率地执行NDRange Kernel，通常需要巨大数量的线程。

(2) Global memory interconnect

　　为读写访问最大化内存带宽的能力对高性能计算非常重要。OpenCL中存在用于读写全局存储器的各种类型的模块，称为负载储存单元(load-store units, LSUs)。

　　与GPU不同，FPGA可以构建最适合应用程序编译内存访问模式的任意自定义的LSU。选择最理想的LSU类型会为应用程序显著地提高设计性能。

(3) Local memory

　　Local Memory比较复杂，不同于GPU结构具有不同级别缓存的架构，FPGA在内部使用专用的内存块实现local memory。

　　Local Memory有以下特性：

　　　　对local memory的每次读写访问都要映射到一个端口；

　　　　可以将local memory中的内容划分为一个或多个存储体bank，以使每个存储体包含该local memory的数据子集。

　　　　一个存储体由一个或多个副本组成。Bank中的副本与其他的副本包含相同的数据。创建副本是为了有效的支持local memory的多次访问。每个副本都具有一个写端口与读端口，从而可以同时访问。如果local memory是double dumped，那么每个副本有四个物理端口，可以最多有三个读端口。

　　在kernel代码中，用 local 类型来声明local memory：

local int lmem[1024];

　　Intel SDK能够自己设定local memor的width, deph, banks, replication, private copies的数量, interconnect等。离线编译器可以分析访问模式然后优化local memory最小化访问竞争。

　　实现Kernel高效工作的关键是不停歇地访问内存。离线编译器始终尝试为Kernel中所有的local memory找到不停歇访问的配置方法，但由于kernel比较复杂，离线编译器可能没有足够的信息推断内存访问是否有冲突，那么就需要local interconnect仲裁寄存器来仲裁内存访问，会降低性能。

(4) Local Memory Banks and Private Copis

　　默认local memory的存储体只在最小维度上起作用，多个存储体允许同时写入。在下面的例子中，循环中每个local memory的访问都有单独的的地址。离线编译器可以推断出访问模式，从而创建四个单独的存储体bank，四个独立bank允许对4个lmem同时访问，从而实现了无停顿的程序配置。此外，离线编译器为lmem创建了两个private copies，从而允许两个work-groups同事pipline运行。

#define BANK_SIZE 4
__attribute__(reqd_work_group_size(8, 1, 1))
kernel void bank_arb_consecutive_multidim(global int* restrict in, global int* restrict out)
{
    local int lmem[1024][BANK_SIZE];
    int gi = get_global_id(0);
    int gs = get_global_size(0);
    int li = get_local_id(0);
    int ls = get_local_size(0);
    int res = in[gi];
    #progma unroll
    for(int i = 0; i < BANK_SIZE; i++)
    {
        lmem[((li + i ) & 0x7f)][i] = res + i;
        res = res >> 1;
    }
    int rdata = 0;
    barrier(CLK_GLOBAL_MEM_FENCE);
    #progma unroll
    for(int i = 0; i < BANK_SIZE; i++)
    {
        rdata ^= lmem[(li + i) & 0x7f][i];
    }
    out[gi] = rdata;
    return;
}