如何使用OpenCL轻松实现FPGA应用编程_第1页
如何使用OpenCL轻松实现FPGA应用编程_第2页
如何使用OpenCL轻松实现FPGA应用编程_第3页
如何使用OpenCL轻松实现FPGA应用编程_第4页
如何使用OpenCL轻松实现FPGA应用编程_第5页
已阅读5页,还剩9页未读 继续免费阅读

下载本文档

版权说明:本文档由用户提供并上传,收益归属内容提供方,若内容存在侵权,请进行举报或认领

文档简介

如何使用OpenCL轻松实现FPGA应用编程对于一个开发人员,可能听说过FPGA,甚至在大学课程设计中,可能拿FPGA做过计算机体系架构相关的验证,但是对于它的第一印象可能觉得这是硬件工程师干的事儿。目前,随着人工智能的兴起,GPU借助深度学习,走上了历史的舞台,并且正如火如荼的跑着各种各样的业务,从training到inference都有它的身影。FPGA也借着这股浪潮,慢慢地走向数据中心,发挥着它的优势。所以接下来就讲讲FPGA如何能让程序员们更好友好的开发,而不需要写那些烦人的RTL代码,不需要使用VCS,Modelsim这样的仿真软件,就能轻轻松松实现unittest。实现这一编程思想的转变,是因为FPGA借助OpenCL实现了编程,程序员只需要通过C/C++添加适当的pragma就能实现FPGA编程。为了让您用OpenCL实现的FPGA应用能够有更高的性能,您需要熟悉如下介绍的硬件。另外,将会介绍编译优化选项,有助于将您的OpenCL应用更好的实现RTL的转换和映射,并部署到FPGA上执行。FPGA概览FPGA是高规格的集成电路,可以实现通过不断的配置和拼接,达到无限精度的函数功能,因为它不像CPU或者GPU那样,基本数据类型的位宽都是固定的,相反FPGA能够做的非常灵活。在使用FPGA的过程中,特别适合一些low-level的操作,比如像bitmasking、shifting、addition这样的操作都可以非常容易的实现。为了达到并行化计算,FPGA内部包含了查找表(LUTs),寄存器(register),片上存储(on-chipmemory)以及算术运算硬核(比如数字信号处理器(DSP)块)。这些FPGA内部的模块通过网络连接在一起,通过编程的手段,可以对连接进行配置,从而实现特定的逻辑功能。这种网络连接可重配的特性为FPGA提供了高层次可编程的能力。(FPGA的可编程性就体现在改变各个模块和逻辑资源之间的连接方式)举个例子,查找表(LUTs)体现的FPGA可编程能力,对于程序猿来说,可以等价理解为一个存储器(RAM)。对于3-bits输入的LUT可以等价理解为一个拥有3位地址线并且8个1-bit存储单元的存储器(一个8长度的数组,数组内每个元素是1bit)。那么当需要实现3-bits数字按位与操作的时候,8长度数组存的是3-bits输入数字的按位与结果,一共是8种可能性。当需要实现3-bits按位异或的时候,8长度数组存的是3-bits输入数字的按位异或结果,一共也是8种可能性。这样,在一个时钟周期内,3-bits的按位运算就能够获取到,并且实现不同功能的按位运算,完全是可编程的(等价于修改RAM内的数值)。3-bits输入LUT实现按位与(bit-wiseAND):注:3-bits输入LUT查找表我们看到的三输入的按位与操作,如下所示,在FPGA内部,可通过LUT实现。如上展示了3输入,1输出的LUT实现。当将LUT并联,串联等方式结合起来后就可以实现更加复杂的逻辑运算了。传统FPGA开发▍传统FPGA与软件开发对比对于传统的FPGA开发与软件开发,工具链可以通过下表简单对比:注:传统FPGA与软件开发对比表重点介绍一下,编译阶段的Synthesis(综合),这部分与软件开发的编译有较大的不同。一般的处理器CPU、GPU等,都是已经生产出来的ASIC,有各自的指令集可以使用。但是对于FPGA,一切都是空白,有的只是零部件,什么都没有,但是可以自己创造任何结构形式的电路,自由度非常的高。这种自由度是FPGA的优势,也是开发过程中的劣势。写到这里,让我想起了最近《神秘的程序员们》中的一个梗:注:漫画来源《神秘的程序员们56》by西乔传统的FPGA开发就像10岁时候的Linux,想吃一个蛋糕,需要自己从原材料开始加工。FPGA正是这种状态,想要实现一个算法,需要写RTL,需要设计状态机,需要仿真正确性。▍传统FPGA开发方式复杂系统,需要使用有限状态机(FSM),一般就需要设计下图包含的三部分逻辑:组合电路,时序电路,输出逻辑。通过组合逻辑获取下一个状态是什么,时序逻辑用于存储当前状态,输出逻辑混合组合、时序电路,得到最终输出结果。然后,针对具体算法,设计逻辑在状态机中的流转过程:实现的RTL是这样的:modulefsm_using_single_always(clock,//clockreset,//Activehigh,synresetreq_0,//Request0req_1,//Request1gnt_0,//Grant0gnt_1);//=============InputPorts=============================inputclock,reset,req_0,req_1;//=============OutputPorts===========================outputgnt_0,gnt_1;//=============InputportsDataType===================wireclock,reset,req_0,req_1;//=============OutputPortsDataType==================reggnt_0,gnt_1;//=============InternalConstants======================parameterSIZE=3;parameterIDLE=3‘b001,GNT0=3’b010,GNT1=3‘b100;//=============InternalVariables======================reg[SIZE-1:0]state;//SeqpartoftheFSMreg[SIZE-1:0]next_state;//combopartofFSM//==========CodestartesHere==========================always@(posedgeclock)begin:FSMif(reset==1’b1)beginstate《=#1IDLE;gnt_0《=0;gnt_1《=0;endelsecase(state)IDLE:if(req_0==1‘b1)beginstate《=#1GNT0;gnt_0《=1;endelseif(req_1==1’b1)begingnt_1《=1;state《=#1GNT1;endelsebeginstate《=#1IDLE;endGNT0:if(req_0==1‘b1)beginstate《=#1GNT0;endelsebegingnt_0《=0;state《=#1IDLE;endGNT1:if(req_1==1’b1)beginstate《=#1GNT1;endelsebegingnt_1《=0;state《=#1IDLE;enddefault:state《=#1IDLE;endcaseendendmodule//EndofModulearbiter传统的RTL设计,对于程序员简直就是噩梦啊,梦啊,啊~~~工具链完全不同,开发思路完全不同,还要分析时序,一个Clock节拍不对,就要推翻重来,重新验证,一切都显得太底层,不是很方便。那么,这些就交给专业的FPGAer吧,下面介绍的OpenCL开发FPGA,有点像25岁的Linux了。有了高层次的抽象。用起来自然也会更加方便。▍基于OpenCL的FPGA开发OpenCL对于FPGA开发,注入了新鲜的血液,一种面向异构系统的编程语言,将FPGA最为异构实现的一种可选设备。由CPUHost端控制整个程序的执行流程,FPGADevice端则作为异构加速的一种方式。异构架构,有助于解放CPU,将CPU不擅长的处理方式,下发到Device端处理。目前典型的异构Device有:GPU、IntelPhi、FPGA。OpenCL是一个用于异构平台编程的框架,主要的异构设备有CPU、GPU、DSP、FPGA以及一些其它的硬件加速器。OpenCL基于C99来开发设备端代码,并且提供了相应的API可以调用。OpenCL提供了标准的并行计算的接口,以支持任务并行和数据并行的计算方式。OpenCL案例分析这里采用Altera官网的矩阵乘法案例进行分析。可以通过如下链接下载案例:AlteraOpenCLMatrixMultiplication代码结构如下:。|--common||--inc||`--AOCLUtils|||--aocl_utils.h|||--opencl.h|||--options.h||`--scoped_ptrs.h||--readme.css|`--src|`--AOCLUtils||--opencl.cpp|`--options.cpp`--matrix_mult|--Makefile|--README.html|--device|`--matrix_mult.cl`--host|--inc|`--matrixMult.h`--src`--main.cpp其中,和FPGA相关的代码是matrix_mult.cl,该部分代码描述了kernel函数,这部分函数会通过编译器生成RTL代码,然后map到FPGA电路中。kernel函数的定义如下:__kernel__attribute((reqd_work_group_size(BLOCK_SIZE,BLOCK_SIZE,1)))__attribute((num_simd_work_items(SIMD_WORK_ITEMS)))voidmatrixMult(__globalfloat*restrictC,__globalfloat*A,__globalfloat*B,intA_width,intB_width)模式比较固定,需要注意的是__global指明从CPU传过来的数据,存放到全局内存中,可以是FPGA片上存储资源,DDR,QDR等,这个视FPGA的OpenCLBSP驱动,会有所区别。num_simd_work_items用于指明SIMD的宽度。reqd_work_group_size指明了工作组的大小。这些概念,可以参考OpenCL的使用手册。函数实现如下://声明本地存储,暂存数组的某一个BLOCK__localfloatA_local[BLOCK_SIZE][BLOCK_SIZE];__localfloatB_local[BLOCK_SIZE][BLOCK_SIZE];//Blockindexintblock_x=get_group_id(0);intblock_y=get_group_id(1);//LocalIDindex(offsetwithinablock)intlocal_x=get_local_id(0);intlocal_y=get_local_id(1);//Computeloopboundsinta_start=A_width*BLOCK_SIZE*block_y;inta_end=a_start+A_width-1;intb_start=BLOCK_SIZE*block_x;floatrunning_sum=0.0f;for(inta=a_start,b=b_start;a《=a_end;a+=BLOCK_SIZE,b+=(BLOCK_SIZE*B_width)){//从globalmemory读取相应BLOCK数据到localmemoryA_local[local_y][local_x]=A[a+A_width*local_y+local_x];B_local[local_x][local_y]=B[b+B_width*local_y+local_x];//Waitfortheentireblocktobeloaded.barrier(CLK_LOCAL_MEM_FENCE);//计算部分,将计算单元并行展开,形成乘法加法树#pragmaunrollfor(intk=0;k《BLOCK_SIZE;++k){running_sum+=A_local[local_y][k]*B_local[local_x][k];}//Waitfortheblocktobefullyconsumedbeforeloadingthenextblock.barrier(CLK_LOCAL_MEM_FENCE);}//StoreresultinmatrixCC[get_global_id(1)*get_global_size(0)+get_global_id(0)]=running_sum;采用CPU模拟仿真FPGA对其进行仿真,不需要programer关心具体的时序是怎么走的,只需要验证逻辑功能就可以,AlteraOpenCLSDK提供了CPU仿真Device设备的功能,采用如下方式进行:#Togeneratea.aocxfilefordebuggingthattargetsaspecificacceleratorboard$aoc-march=emulatordevice/matrix_mult.cl-obin/matrix_mult.aocx--fp-relaxed--fpc--no-interleavingdefault--board《your-board》#GenerateHostexe.$make#Toruntheapplication$envCL_CONTEXT_EMULATOR_DEVICE_ALTERA=8。/bin/host-ah=512-aw=512-bw=512上述脚本中,通过-march=emulator设置创建一个可用于CPUdebug的设备可执行文件。-g添加调试flag。—board用于创建适配该设备的debugging文件。CL_CONTEXT_EMULATOR_DEVICE_ALTERA为用于CPU仿真的设备数量。当执行上述脚本后,输出如下:$envCL_CONTEXT_EMULATOR_DEVICE_ALTERA=8。/bin/host-ah=512-aw=512-bw=512Matrixsizes:A:512x512B:512x512C:512x512InitializingOpenCLPlatform:AlteraSDKforOpenCLUsing8device(s)EmulatorDevice:EmulatedDevice。..EmulatorDevice:EmulatedDeviceUsingAOCX:matrix_mult.aocxGeneratinginputmatricesLaunchingfordevice0(globalsize:512,64)。..Launchingfordevice7(globalsize:512,64)Time:5596.620msKerneltime(device0):5500.896ms。..Kerneltime(device7):5137.931msThroughput:0.05GFLOPSComputingreferenceoutputVerifyingVerification:PASS通过仿真时候设置Device=8,模拟8个设备运行(512,512)*(512,512)规模的矩阵,最终验证正确。接下来就可以将其真正编译到FPGA设备上后运行。FPGA设备上运行矩阵乘这个时候,真正要将代码下载到FPGA上执行了,这时候,只需要做一件事,那就是用OpenCLSDK提供的编译器,将*.cl代码适配到FPGA上,执行编译命令如下:$aocdevice/matrix_mult.cl-obin/matrix_mult.aocx--fp-relaxed--fpc--no-interleavingdefault--board《your-bo

温馨提示

  • 1. 本站所有资源如无特殊说明,都需要本地电脑安装OFFICE2007和PDF阅读器。图纸软件为CAD,CAXA,PROE,UG,SolidWorks等.压缩文件请下载最新的WinRAR软件解压。
  • 2. 本站的文档不包含任何第三方提供的附件图纸等,如果需要附件,请联系上传者。文件的所有权益归上传用户所有。
  • 3. 本站RAR压缩包中若带图纸,网页内容里面会有图纸预览,若没有图纸预览就没有图纸。
  • 4. 未经权益所有人同意不得将文件中的内容挪作商业或盈利用途。
  • 5. 人人文库网仅提供信息存储空间,仅对用户上传内容的表现方式做保护处理,对用户上传分享的文档内容本身不做任何修改或编辑,并不能对任何下载内容负责。
  • 6. 下载文件中如有侵权或不适当内容,请与我们联系,我们立即纠正。
  • 7. 本站不保证下载资源的准确性、安全性和完整性, 同时也不承担用户因使用这些下载资源对自己和他人造成任何形式的伤害或损失。

评论

0/150

提交评论