Geekly articles weekly

FPGA加速器编程示例



不久前，我们谈到了新的Selectel服务-FPGA 加速器上的云高性能计算 。 在关于该主题的新文章中，我们考虑了一个用于构造Mandelbrot集的FPGA编程示例，Mandelbrot集是一种用于可视化分形图像的著名数学算法。 本文使用了来自Euler Project网站的材料。



而不是前言


首先，几个术语。 具有FPGA加速器的计算机系统-通常，这是一个PCIe适配器，其中FPGA芯片是x64服务器的一部分。 加速器承担了一项单独的资源密集型任务，其中涉及并行计算，并且执行速度比x64处理器快许多个数量级，从而卸载了并行计算并提高了整个计算系统的性能。 例如，具有10万次重复的计算周期可以在FPGA上仅执行一次，而不是在传统的x64处理器上顺序执行10万次。 逻辑元素，硬件资源，通信链接，FPGA芯片由用户直接针对任务本身进行编程，这使您可以将任务实现为硅算法的实现-硅算法，从而以非常适度的功耗实现高性能。

如今，即使对于初创企业来说，进入FPGA技术的门槛也很容易达到-带有FPGA加速器和所有必要软件（SDK）的服务器可以以合理的价格租借到Selectel云中（所谓的“云FPGA”），并且对FPGA中的Open CL标准的支持导致知道如何使用C的程序员能够在FPGA上准备和运行程序。

展望未来：试用FPGA


下面描述的用于构造Mandelbrot集的编程示例已经在Selectel实验室的测试服务器上实现，任何人都可以在其中评估其性能（需要注册）。


该项目以代码形式提供，并准备进行编译。 Selectel使用Intel Arria 10 FPGA加速器提供对服务器的远程访问。 在服务器端，部署了用于开发，调试和编译OpenCL的SDK和BSP工具，用于准备主机应用程序（用于服务器中央处理器的控制应用程序）的Visual Studio代码。
请注意，示例本身没有任何应用价值；选择该示例是出于演示使用并行性原理的加速方法的原因。 通过本示例，读者可以熟悉在具有FPGA的异构计算系统中设计应用程序的途径-稍后，该途径可用于开发具有并行计算的自己的应用程序。
更新 ：2018年春季，英特尔推出了具有集成Arria 10 FPGA芯片的高性能混合处理器Xeon Gold 6138P。 到2018年底，此类串行处理器有望通过英特尔合作伙伴向客户提供。 我们在Selectel期待这种芯片，并希望成为俄罗斯第一个为我们的客户提供测试该独特新产品的机会。

关于FPGA编程的OpenCL标准


OpenCL标准由Khronos集团开发，Khronos集团是全球领先的芯片和软件制造商，包括Intel，AMD，Apple，ARM，Nvidia，Sony Computer Entertainment等。它旨在编写在各种类型的处理器（包括FPGA）上使用并行计算的应用程序。 OpenCL标准包括基于C99语言版本的C编程语言（C99的最新版本为2007年11月15日的ISO / IEC 9899：1999 / Cor 3：2007）和一个应用程序编程环境。

使用OpenCL进行高性能计算的流行是基于以下事实：它是一种开放标准，并且不需要许可证即可使用。 此外，OpenCL并不将支持的设备范围限制为任何特定品牌，允许在同一软件平台上使用来自不同制造商的硬件。

关于OpenCL的其他信息： Habr上的 OpenCL 简介 。


一段历史-OpenCL标准之前存在的FPGA设计路线极其具体且耗时，而就复杂性而言，它甚至优于定制芯片设计（ASIC，专用集成电路，“专用集成电路”）。 需要对FPGA硬件结构有透彻的了解，其配置必须以低级硬件描述语言（HDL）进行。 由于极其复杂，这种设计和验证途径已经并且仍然是一门艺术，只有有限的开发人员才能使用。

英特尔针对FPGA的OpenCL支持工具包的问世部分解决了软件开发人员对FPGA编程的可访问性问题。 程序员独立选择适合并行处理的算法部分，并用C语言进行描述，然后面向FPGA的英特尔OpenCL编译器创建一个二进制配置文件，以在加速器上运行该算法片段。
使用通常的Visual Studio环境或标准的gcc编译器，可以准备一个主机应用程序（在主要x64处理器上执行的.exe类型的应用程序），而所有必需的支持库都包含在SDK中。 启动主机应用程序时，将加载FPGA固件，数据将被加载到芯片内核中，并且处理将根据所设想的算法进行。

FPGA（FPGA）是一种用户可重新编程的大规模并行硬件结构，具有数百万个逻辑元素，数千个DSP信号块和数十兆字节的高速缓存，可进行板载计算，而无需访问服务器的主内存模块。 快速的I / O接口（10GE，40GE，100GE，PCIe Gen 3等）使您可以有效地与服务器的主处理器交换数据。


OpenCL标准是用于执行异构软件的环境。 环境由两个独立的部分组成：

  1. 主机软件-用C / C ++编写并使用OpenCL API函数集在服务器的主中央处理器上运行的应用程序。 主机服务器组织计算，提供源和接收输出数据的整个过程，并通过FPGA加速器与所有服务器系统进行交互。
  2. 加速器软件-用OpenCL C语言（受多种限制的C语言）编写的程序，已编译为可在FPGA芯片上运行。

典型的并行计算服务器是基于x64的计算机（用于运行主机应用程序），该计算机包括硬件FPGA加速器，通常通过PCI-Express总线连接。 顺便说一下，Selectel实验室就提供了这样的系统。

FPGA加速器的编程和编译序列包括两个阶段。 主机应用程序代码由标准编译器（Visual C ++，GCC）进行编译，以获取服务器操作系统中的可执行文件（例如* .exe）。 FPGA加速器的源代码（内核，内核）由AOC编译器作为SDK的一部分准备，并带有二进制文件（* .aocx）。 该文件仅用于加速器编程。

OpenCL软件编译环境架构
图 OpenCL软件编译环境架构

考虑一些用于以两种方式计算大向量的示例代码
（ PS请勿射击钢琴家-以下使用来自Euler Project网站的代码 ）：

void inc (float *a, float c, int N) { for (int i = 0; i<N; i++) a[i] = a[i] + c; } void main() { ... inc(a,c,N); ... } 

 _kernel void inc (_global float *a, float c) { int i = get_global_id(0); a[i] = a[i] + c; } void main() { ... clEnqueueNDRangeKernel(...,&N,...) ... }

开头的代码是使用标量元素顺序计算方法的C中单线程实现的外观示例。

该代码的第二版是在OpenCL上算法的可能实现，其形式为在FPGA加速器上计算的函数。 没有循环，并且计算在循环的一次迭代中进行。 向量数组的计算在执行此函数的N个副本时发生。 每个副本都有自己的索引，并在循环中替换为迭代器，并且在执行代码时从主机设置重试次数。 迭代器操作由get_global_id（）函数提供，该函数在0≤index <N范围内的索引下工作。

要点：建立分形


Mandelbrot集是复平面上的点“ c”的数组，对于它们，Z0 = 0的递归关系Zn + 1 =Zn²+ c定义了有界序列。

我们定义Zn = Zn + IYn，也定义c = p + iq。
对于每个点，将计算以下顺序：

Xn +1 =Xn²+Yn²+ p
Yn + 1 = 2XnYn + q


将点在集合中的归属在每次迭代时的计算公式如下：
Xn²+Yn²<4。

要在屏幕上显示Mandelbrot集，我们定义一个规则：

  1. 如果不等式在任何迭代中都成立，则该点进入集合并显示为黑色。
  2. 如果不等式不成立，则从某个迭代值n = N开始，则颜色由迭代次数N确定。

主机上的计算过程如下：

  • 像素窗口内每个点的迭代次数计算将分配给mandel_pixel（）函数。
  • 图像点的顺序枚举将由softwareCalculateFrame（）函数提供。 参数指定计算点的实际间隔，算法的实际步长以及指向图像大小（theWidth * theHeight）的颜色缓冲区的指针。
  • 点的颜色由SoftColorTable调整。

让我们继续执行代码：

 inline unsigned int mandel_pixel( double x0, double y0, unsigned int maxIterations ) { // variables for the calculation double x = 0.0; double y = 0.0; double xSqr = 0.0; double ySqr = 0.0; unsigned int iterations = 0; // perform up to the maximum number of iterations to solve // the current point in the image while ( xSqr + ySqr < 4.0 &&iterations < maxIterations ) { // perform the current iteration xSqr = x*x; ySqr = y*y; y = 2*x*y + y0; x = xSqr - ySqr + x0; // increment iteration count iterations++; } // return the iteration count return iterations; } 

 int softwareCalculateFrame( double aStartX, double aStartY, double aScale, unsigned int* aFrameBuffer ) { // temporary pointer and index variables unsigned int * fb_ptr = aFrameBuffer; unsigned int j, k, pixel; // window position variables double x = aStartX; double y = aStartY; double cur_x, cur_y; double cur_step_size = aScale; // for each pixel in the y dimension window for ( j = 0, cur_y = y; j < theHeight; j++, cur_y -= cur_step_size ) { // for each pixel in the x dimension of the window for ( cur_x = x, k = 0; k< theWidth; k++, cur_x += cur_step_size ) { // set the value of the pixel in the window pixel = mandel_pixel(cur_x, cur_y, theSoftColorTableSize); if ( pixel == theSoftColorTableSize ) *fb_ptr++ = 0x0; else *fb_ptr++ = theSoftColorTable[pixel]; } } return 0; }

每个像素彼此独立地计算，因此该过程可以并行化。 在为FPGA加速器实现算法时，将创建SIMD指令以计算每个迭代像素的数量（确定调色板中的颜色代码）。 通过运行操作（theWidth * theHeight），可通过OpenCL对图像缓冲区上的两个嵌套循环的实现进行框架化。

下面清单中的内核实例称为工作项，所有实例的集合称为索引空间。 硬件功能的功能包括：

  • 函数声明以__kernel关键字开头。
  • 硬件功能的类型-返回值的类型始终为空。
  • 通过作为参数传递的缓冲区来完成返回值。
    • 前三个参数定义了材质网格，该网格的节点对应于输出图像的像素。
    • 第四个参数限制了迭代次数，从而防止了属于Mandelbrot集的点的循环。
    • 第五个参数是指向输出颜色缓冲区的指针。
    • __global关键字表示将通过其传输缓冲区的内存类型：这是加速器本身上的常规DDR（QDR）内存。
    • limit关键字使优化器禁止使用间接缓冲区引用。
    • 在第6个参数中，传递了指向调色板的指针。
    • __constant关键字通过生成具有只读属性的缓存来优化缓冲区访问。

    清单中的功能描述与x64处理器的实现非常接近。 在这里，当前内核实例的定义是通过get_global_id函数完成的，维数（0，1）作为参数传递到该函数中。

    为了更好的优化，引入了周期开始的明确指示。 在缺少有关编译时迭代次数的信息的情况下，将明确指出循环步骤的数目，因为将为其创建自己的硬件块。 使用这种编码，由于FPGA需要消耗大量周期资源，因此应该“回顾”安装在加速器上的特定芯片的容量。

     //////////////////////////////////////////////////////////////////// // mandelbrot_kernel.cl : Hardware implementation of the mandelbrot algorithm //////////////////////////////////////////////////////////////////// // Amount of loop unrolling. #ifndef UNROLL #define UNROLL 20 #endif // Define the color black as 0 #define BLACK 0x00000000 __kernel void hw_mandelbrot_frame ( const double x0, const double y0, const double stepSize, const unsigned int maxIterations, __global unsigned int *restrict framebuffer, __constant const unsigned int *restrict colorLUT, const unsigned int windowWidth) { // Work-item position const size_t windowPosX = get_global_id(0); const size_t windowPosY = get_global_id(1); const double stepPosX = x0 + (windowPosX * stepSize); const double stepPosY = y0 - (windowPosY * stepSize); // Variables for the calculation double x = 0.0; double y = 0.0; double xSqr = 0.0; double ySqr = 0.0;</code> <code>unsigned #pragma while { int iterations = 0; // Perform up to the maximum number of iterations to solve // the current work-item's position in the image // The loop unrolling factor can be adjusted based on the amount of FPGA // resources available. unroll UNROLL xSqr + ySqr < 4.0 && iterations < maxIterations ) // Perform the current iteration xSqr = x*x; ySqr = y*y; y = 2*x*y + stepPosY; x = xSqr - ySqr + stepPosX; // Increment iteration count iterations++; } // Output black if we never finished, and a color from the look up table otherwise framebuffer[windowWidth * windowPosY + windowPosX] = (iterations == maxIterations) ? BLACK : colorLUT[iterations]; }

    在编译算法的硬件实现之前，需要在主机上安装面向OpenCL的英特尔FPGA SDK实用程序包。 在预安装的软件中，您必须包括特定加速板制造商的BSP（板级支持软件包）。 在示例中，安装了Intel Quartus Prime Pro 16.1，并支持Euler线程加速器（Intel Arria 10）的OpenCL和BSP。

    以下是路径和环境变量的配置。 变量ALTERAOCLSDKROOT包含指向Intel FPGA SDK的路径，变量AOCL_BOARD_PACKAGE_ROOT包含指向加速器BSP的路径。

     set ALTERAOCLSDKROOT=C:\intelFPGA_pro\16.1\hld set AOCL_BOARD_PACKAGE_ROOT=C:\intelFPGA_pro\16.1\hld\board\euler_thread set path=%path%;C:\intelFPGA_pro\16.1\hld\bin set path=%path%;C:\intelFPGA_pro\16.1\quartus\bin64 set path=%path%;C:\intelFPGA_pro\16.1\hld\board\a10_ref\windows64\bin set path=%path%;C:\intelFPGA_pro\16.1\hld\host\windows64\bin set path=%path%;C:\intelFPGA_pro\16.1\qsys\bin set path=%path%;C:\Program Files (x86)\GnuWin32\bin\

    要进行编译，请使用SDK中的aoc编译器。

     aoc mandelbrot_kernel.cl -o mandelbrot_kernel.aocx --board thread -v -v --report

    我们解密：mandelbrot_kernel.cl-带有源文本的文件，mandelbrot_kernel.aocx-用于FPGA编程的输出对象文件，thread-BSP包中加速器的名称。 --report开关显示FPGA资源使用情况报告。 –v开关在编译期间显示诊断信息。 内核的资源消耗报告如下：

    + ------------------------------------------------- ------------------- +
    ; 估计资源使用情况摘要；
    + ---------------------------------------- + -------- ------------------- +
    ; 资源+使用情况；
    + ---------------------------------------- + -------- ------------------- +
    ; 逻辑利用率； 49％
    ; ALUT； 26％;
    ; 专用逻辑寄存器； 25％;
    ; 内存块； 21％
    ; DSP模块 16％；
    + ---------------------------------------- + -------- -------------------;

    为了编译主机应用程序，该示例使用了安装了Microsoft SDK 7.1的Microsoft Visual Studio 2010 Express程序包。 在项目设置中，选择了x64的配置。 接下来，连接用于外部头文件的文件夹，并在链接器设置中指定其他英特尔FPGA SDK库的路径。
    包含文件的其他目录= $（ALTERAOCLSDKROOT）\ host \ include;
    其他库目录= $（AOCL_BOARD_PACKAGE_ROOT）\ Windows64 \ lib;
     $(ALTERAOCLSDKROOT)\host\windows64\lib;

    在加速器上启动内核的总体行动计划如下：

    1. 获取平台列表
    2. 获取设备列表
    3. 创建环境；
    4. 将内核加载到设备中；
    5. 发送输入缓冲区到设备；
    6. 运行内核以执行；
    7. 从设备读取输出缓冲区；
    8. 自由上下文。

    考虑与内核启动直接相关的一些要点。 因此，一个内核被设计为处理图像的一个像素。 因此，您需要运行N个内核实例，其中N是图像中像素的总数。

    在下面，我们注意到服务器中有多个加速器板的情况，然后可以在它们之间分配任务。 在每个加速器中，您都需要加载内核（文件mandelbrot_kernel.aocx）。 假设加速器的数量为numDevices，并且图像行在所有加速器之间划分：

     #define MAXDEV 10 static cl_context theContext; static cl_program theProgram; static cl_kernel theKernels[MAXDEV]; //.. // Create the program object theProgram = createProgramFromBinary( theContext, "mandelbrot_kernel.aocx", theDevices, numDevices); // Create the kernels for ( unsigned i = 0; i < numDevices; ++i ) theKernels[i] = clCreateKernel( theProgram, "hw_mandelbrot_frame", &theStatus ); // Create output pixel buffers for every kernel for( unsigned i = 0; i < numDevices; ++i ) thePixelData[i] = clCreateBuffer(theContext, CL_MEM_WRITE_ONLY, thePixelDataWidth*rowsPerDevice[i]*sizeof(unsigned int), NULL, &theStatus); // Preparing and writing palette buffer to every device theHardColorTable = clCreateBuffer(theContext, CL_MEM_READ_ONLY, aColorTableSize*sizeof(unsigned int), NULL, &theStatus); for( unsigned i = 0; i < numDevices; i++ ) theStatus = clEnqueueWriteBuffer(theQueues[i], theHardColorTable, CL_TRUE, 0, aColorTableSize*sizeof(unsigned int), aColorTable, 0, NULL, NULL); // Preparing kernels and run unsigned rowOffset = 0; for ( unsigned i = 0; i < numDevices; rowOffset += rowsPerDevice[i++] ) { // Create ND range size size_t globalSize[2] = { thePixelDataWidth, rowsPerDevice[i] }; // Set the arguments unsigned argi = 0; theStatus = clSetKernelArg (theKernels[i], argi++, sizeof(cl_double), (void*) &aStartX ); const double offsetedStartY = aStartY - rowOffset * aScale; theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_double), (void*)&offsetedStartY); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_double), (void*)&aScale); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_uint), (void*)&theHardColorTableSize); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_mem), (void*)&thePixelData[i]); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_mem), (void*)&theHardColorTable); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_uint), (void*)&theWidth); // Launch kernel theStatus = clEnqueueNDRangeKernel(theQueues[i], theKernels[i], 2, NULL, globalSize, NULL, 0, NULL, NULL); } rowOffset = 0; for( unsigned i = 0; i < numDevices; rowOffset += rowsPerDevice[i++] ) { // Read the output theStatus = clEnqueueReadBuffer(theQueues[i], thePixelData[i], CL_TRUE, 0, thePixelDataWidth*rowsPerDevice[i]*sizeof(unsigned int), &aFrameBuffer[rowOffset * theWidth], 0, NULL, NULL); } / / . .

    • createProgramFromBinary函数从对象文件创建OpenCL程序对象。
    • 接下来，针对每个设备，基于程序对象创建一个内核。
    • 创建PixelData缓冲区以接收每个内核的输出。
    • 创建一个缓冲区来存储调色板并将其加载到每个加速器中。
    • 接下来，对于每个设备，使用clSetKernelArg函数设置本地应用程序参数和内核参数的绑定。
    • 参数由内核函数声明中的序列号确定，从零开始。

    下一个重点是根据根据globalSize数组的索引空间确定任务的大小。 该阵列可以是一维，二维或三维的。 对于每个维度，维度均以整数形式给出。 空间的大小将确定内核中工作项的索引顺序。

    在示例中，为每个核心指定了一个二维空间，其中一个轴是像素行元素，第二个轴是在此设备上处理的一组图像线。 在内核代码中，通过调用get_global_id（0）获得该行中的像素号，该行号为get_global_id（1）。 将globalSize变量传递到clEnqueueNDRangeKernel函数，以启动所需数量的要执行的内核实例。

    完成内核执行后，像素缓冲区将从设备读取到本地阵列。 让我们通过每秒的帧数来评估性能-结果在SelectelTechDay会议上进行的演示中可见（ 请参阅本文的开头 ）。

    结论


    用高级语言编程FPGA加速器无疑将开发人员使用该技术的门槛降低了一个数量级。 例如，对于那些刚刚掌握此工具包的人，甚至还有著名的“ Hello World”示例的FPGA实现。

    但不是那么简单。 编写（尤其是调试）明确解决实际应用问题的算法仍然需要很高的专业水平。 另一个限制是，每个FPGA芯片只能在应用程序中执行一个计算任务。 对于另一项任务，必须再次对其进行重新编程。
    顺便说一句，使用平台的模型允许您在主机上拥有多个FPGA加速器，尽管这是一个相当昂贵的解决方案。
    主机（主机应用程序）管理创建上下文（加速器的数据结构）和命令队列的过程。 即 单个主机应用程序可以在不同的加速器上加载它们，其中在FPGA上具有用于并行计算的各种子任务：
    内核1 =>加速器A
    内核2 =>加速器B

    尽管如此，掌握FPGA加速器的努力是值得的-在许多应用领域中，该技术已变得不可或缺：电信，生物技术，大数据处理，模式识别，信号和图像处理，计算数学和物理场建模。

    本文的其他信息：
    www.altera.com是英特尔FPGA核心资源。
    www.eulerproject.com是Euler项目的官方网站。
    Altera + OpenCL：我们在不了解VHDL / Verilog的情况下使用FPGA进行编程 -有关Habr的文章。

Source: https://habr.com/ru/post/zh-CN418403/

More articles:


All Articles