您的位置 首页 制造

怎么运用OpenCL轻松完成FPGA使用编程

如何使用OpenCL轻松实现FPGA应用编程-实现这一编程思想的转变,是因为 FPGA 借助 OpenCL 实现了编程,程序员只需要通过 C/C++ 添加适当的 pragma 就能实现 FPGA 编程。为了让您用 OpenCL 实现的 FPGA 应用能够有更高的性能,您需要熟悉如下介绍的硬件。另外,将会介绍编译优化选项,有助于将您的 OpenCL 应用更好的实现 RTL 的转换和映射,并部署到 FPGA 上执行。

关于一个开发人员,或许听说过 FPGA,甚至在大学课程规划中,或许拿 FPGA 做过核算机体系架构相关的验证,可是关于它的第一印象或许觉得这是硬件工程师干的事儿。

现在,跟着人工智能的鼓起,GPU 凭借深度学习,走上了前史的舞台,而且正如火如荼的跑着各式各样的事务,从 training 到 inference 都有它的身影。FPGA 也借着这股浪潮,渐渐地走向数据中心,发挥着它的优势。所以接下来就讲讲 FPGA 怎么能让程序员们更好友爱的开发,而不需求写那些烦人的 RTL 代码,不需求运用 VCS,Modelsim 这样的仿真软件,就能轻轻松松完成 unit test。

完成这一编程思维的改动,是由于 FPGA 凭借 OpenCL 完成了编程,程序员只需求经过 C/C++ 增加恰当的 pragma 就能完成 FPGA 编程。为了让您用 OpenCL 完成的 FPGA 运用可以有更高的功用,您需求了解如下介绍的硬件。别的,将会介绍编译优化选项,有助于将您的 OpenCL 运用更好的完成 RTL 的转化和映射,并布置到 FPGA 上履行。

FPGA 概览

FPGA 是高规范的集成电路,可以完成经过不断的装备和拼接,到达无限精度的函数功用,由于它不像 CPU 或许 GPU 那样,根本数据类型的位宽都是固定的,相反 FPGA 可以做的十分灵敏。在运用 FPGA 的进程中,特别合适一些 low-level 的操作,比方像 bit masking、shifting、addiTIon 这样的操作都可以十分简略的完成。

为了到达并行化核算,FPGA 内部包括了查找表(LUTs),寄存器(register),片上存储(on-chip memory)以及算术运算硬核(比方数字信号处理器 (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-wise AND):

怎么运用OpenCL轻松完成FPGA运用编程

注: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 是这样的:

module fsm_using_single_always (

clock , // clockreset , // AcTIve high, syn resetreq_0 , // Request 0req_1 , // Request 1gnt_0 , // Grant 0gnt_1

);//=============Input Ports=============================input clock,reset,req_0,req_1; //=============Output Ports===========================output gnt_0,gnt_1;//=============Input ports Data Type===================wire clock,reset,req_0,req_1;//=============Output Ports Data Type==================reg gnt_0,gnt_1;//=============Internal Constants======================parameter SIZE = 3 ;

parameter IDLE = 3‘b001,GNT0 = 3’b010,GNT1 = 3‘b100 ;//=============Internal Variables======================reg [SIZE-1:0] state ;// Seq part of the FSMreg [SIZE-1:0] next_state ;// combo part of FSM//==========Code startes Here==========================always @ (posedge clock)begin : FSMif (reset == 1’b1) begin

state 《= #1 IDLE;

gnt_0 《= 0;

gnt_1 《= 0;end else

case(state)

IDLE : if (req_0 == 1‘b1) begin

state 《= #1 GNT0;

gnt_0 《= 1; end else if (req_1 == 1’b1) begin

gnt_1 《= 1;

state 《= #1 GNT1; end else begin

state 《= #1 IDLE; end

GNT0 : if (req_0 == 1‘b1) begin

state 《= #1 GNT0; end else begin

gnt_0 《= 0;

state 《= #1 IDLE; end

GNT1 : if (req_1 == 1’b1) begin

state 《= #1 GNT1; end else begin

gnt_1 《= 0;

state 《= #1 IDLE; end

default : state 《= #1 IDLE;

endcaseendendmodule // End of Module arbiter

传统的 RTL 规划,关于程序员几乎便是噩梦啊,梦啊,啊~~~东西链彻底不同,开发思路彻底不同,还要剖析时序,一个 Clock 节拍不对,就要推翻重来,从头验证,一切都显得太底层,不是很便利。那么,这些就交给专业的 FPGAer 吧,下面介绍的 OpenCL 开发 FPGA,有点像 25 岁的 Linux 了。有了高层次的笼统。用起来天然也会愈加便利。

▍根据 OpenCL 的 FPGA 开发

OpenCL 关于 FPGA 开发,注入了新鲜的血液,一种面向异构体系的编程言语,将 FPGA 最为异构完成的一种可选设备。由 CPU Host 端操控整个程序的履行流程,FPGA Device 端则作为异构加快的一种方法。异构架构,有助于解放 CPU,将 CPU 不拿手的处理方法,下发到 Device 端处理。现在典型的异构 Device 有:GPU、Intel Phi、FPGA。

OpenCL 是一个用于异构渠道编程的结构,首要的异构设备有 CPU、GPU、DSP、FPGA 以及一些其它的硬件加快器。OpenCL 根据 C99 来开发设备端代码,而且供给了相应的 API 可以调用。OpenCL 供给了规范的并行核算的接口,以支撑使命并行和数据并行的核算方法。

OpenCL 事例剖析

这儿选用 Altera 官网的矩阵乘法事例进行剖析。可以经过如下链接下载事例:Altera OpenCL Matrix MulTIplicaTIon

代码结构如下:

。|– 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)))void matrixMult( __global float *restrict C,

__global float *A,

__global float *B,

int A_width,

int B_width)

形式比较固定,需求留意的是 __global 指明从 CPU 传过来的数据,存放到大局内存中,可以是 FPGA 片上存储资源,DDR,QDR 等,这个视 FPGA 的 OpenCL BSP 驱动,会有所区别。num_simd_work_items 用于指明 SIMD 的宽度。reqd_work_group_size 指明晰作业组的巨细。这些概念,可以参阅 OpenCL 的运用手册。

函数完成如下:

// 声明本地存储,暂存数组的某一个 BLOCK__local float A_local[BLOCK_SIZE][BLOCK_SIZE];

__local float B_local[BLOCK_SIZE][BLOCK_SIZE];// Block indexint block_x = get_group_id(0);int block_y = get_group_id(1);// Local ID index (offset within a block)int local_x = get_local_id(0);int local_y = get_local_id(1);// Compute loop boundsint a_start = A_width * BLOCK_SIZE * block_y;int a_end = a_start + A_width – 1;int b_start = BLOCK_SIZE * block_x;float running_sum = 0.0f;for (int a = a_start, b = b_start; a 《= a_end; a += BLOCK_SIZE, b += (BLOCK_SIZE * B_width))

{ // 从 global memory 读取相应 BLOCK 数据到 local memory

A_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]; // Wait for the entire block to be loaded.

barrier(CLK_LOCAL_MEM_FENCE); // 核算部分,将核算单元并行打开,构成乘法加法树

#pragma unroll

for (int k = 0; k 《 BLOCK_SIZE; ++k)

{

running_sum += A_local[local_y][k] * B_local[local_x][k];

} // Wait for the block to be fully consumed before loading the next block.

barrier(CLK_LOCAL_MEM_FENCE);

}// Store result in matrix CC[get_global_id(1) * get_global_size(0) + get_global_id(0)] = running_sum;

选用 CPU 模仿仿真 FPGA

对其进行仿真,不需求 programer 关怀详细的时序是怎么走的,只需求验证逻辑功用就可以,Altera OpenCL SDK 供给了 CPU 仿真 Device 设备的功用,选用如下方法进行:

# To generate a .aocx file for debugging that targets a specific accelerator board$ aoc -march=emulator device/matrix_mult.cl -o bin/matrix_mult.aocx –fp-relaxed —fpc –no-interleaving default –board 《your-board》# Generate Host exe.$ make# To run the application$ env CL_CONTEXT_EMULATOR_DEVICE_ALTERA=8 。/bin/host -ah=512 -aw=512 -bw=512

上述脚本中,经过 -march=emulator 设置创立一个可用于 CPU debug 的设备可履行文件。-g 增加调试 flag。—board 用于创立适配该设备的 debugging 文件。CL_CONTEXT_EMULATOR_DEVICE_ALTERA 为用于 CPU 仿真的设备数量。

当履行上述脚本后,输出如下:

$ env CL_CONTEXT_EMULATOR_DEVICE_ALTERA=8 。/bin/host -ah=512 -aw=512 -bw=512Matrix sizes:

A: 512 x 512

B: 512 x 512

C: 512 x 512Initializing OpenCL

Platform: Altera SDK for OpenCL

Using 8 device(s)

EmulatorDevice : Emulated Device

。..

EmulatorDevice : Emulated Device

Using AOCX: matrix_mult.aocx

Generating input matrices

Launching for device 0 (global size: 512, 64)

。..

Launching for device 7 (global size: 512, 64)

Time: 5596.620 ms

Kernel time (device 0): 5500.896 ms

。..

Kernel time (device 7): 5137.931 ms

Throughput: 0.05 GFLOPS

Computing reference output

Verifying

Verification: PASS

经过仿真时分设置 Device = 8,模仿 8 个设备运转 (512, 512) * (512, 512) 规划的矩阵,终究验证正确。接下来就可以将其真实编译到 FPGA 设备上后运转。

FPGA 设备上运转矩阵乘

这个时分,真实要将代码下载到 FPGA 上履行了,这时分,只需求做一件事,那便是用 OpenCL SDK 供给的编译器,将 *.cl 代码适配到 FPGA 上,履行编译指令如下:

$ aoc device/matrix_mult.cl -o bin/matrix_mult.aocx –fp-relaxed –fpc –no-interleaving default –board 《your-board》

这个进程比较慢,一般需求几个小时到 10 几个小时,视 FPGA 上资源巨细而定。(现在这部分时刻太长暂时无法处理,由于这儿的编译,其实是在行程一个可以正常作业的电路,软件会进行布局布线等作业)

等候编译完成后,将生成的 matrix_mult.aocx 文件烧写到 FPGA 上就 ok 啦。

烧写的指令如下:

$ aocl program 《your-board》 matrix_mult.aocx

这时分,功德圆满,可以运转 host 端程序了:

$ 。/host -ah=512 -aw=512 -bw=512Matrix sizes:

A: 512 x 512

B: 512 x 512

C: 512 x 512Initializing OpenCL

Platform: Altera SDK for OpenCL

Using 1 device(s)

《your-board》 : Altera OpenCL QPI FPGA

Using AOCX: matrix_mult.aocx

Generating input matrices

Launching for device 0 (global size: 512, 512)

Time: 2.253 ms

Kernel time (device 0): 2.191 ms

Throughput: 119.13 GFLOPS

Computing reference output

Verifying

Verification: PASS

可以看到,矩阵乘法可以在 FPGA 上正常运转,吞吐大概在 119GFlops 左右。

小结

从上述的开发流程,OpenCL 大大的解放了 FPGAer 的开发周期,而且关于软件开发者,也比较简略上手。这是他的优势,可是现在开发进程中,仍是存在一些问题,如:编译器优化缺乏,比较 RTL 写的功用存在距离;编译到 Device 端时刻太长。不过这些跟着职业的开展,一定会渐渐的前进。

责任编辑:gt

声明:本文内容来自网络转载或用户投稿,文章版权归原作者和原出处所有。文中观点,不代表本站立场。若有侵权请联系本站删除(kf@86ic.com)https://www.86ic.net/bandaoti/zhizao/96708.html

为您推荐

联系我们

联系我们

在线咨询: QQ交谈

邮箱: kf@86ic.com

关注微信
微信扫一扫关注我们

微信扫一扫关注我们

返回顶部