码迷,mamicode.com
首页 > 其他好文 > 详细

OpenCL设计优化(基于Intel FPGA SDK for OpenCL)

时间:2020-07-24 21:20:09      阅读:88      评论:0      收藏:0      [点我收藏+]

标签:single   class   结构   重点   pil   span   i++   clust   例子   

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

 

OpenCL设计优化(基于Intel FPGA SDK for OpenCL)

标签:single   class   结构   重点   pil   span   i++   clust   例子   

原文地址:https://www.cnblogs.com/zhuzhudong/p/13372964.html

(0)
(0)
   
举报
评论 一句话评论(0
登录后才能评论!
© 2014 mamicode.com 版权所有  联系我们:gaon5@hotmail.com
迷上了代码!