如何在FPGA ZYNQ上使用OpenCL（保姆级教学）

1.创建vivado HLS项目

命名为“vadd_OpenCL”

将顶级函数将顶级函数命名为 （vadd）后点击next再点击next

选择自己的开发版型号，我使用的是xilinx的PYNQ Z2开发板，型号为xc7z020clg400-1，选中后点击OK，再点击Finish即可完成HLS的配置。

2.编写一个简单的OpenCL内核

在左侧状态栏找到vadd_OpenCL下的Source，右击该选项，选择New File

将文件命名为vadd.cl，点击保存，就可以在主页面写OpenCL代码了

代码如下：

#include <clc.h>
    
__kernel void __attribute__ ((reqd_work_group_size(128,1,1)))
vadd( __global int *a, __global int *b, __global int *c) {
  int i = get_global_id(0);
    
  c[i] = a[i] + b[i];
}

看到报错不要慌，快捷键Ctrl + s保存一下，报错就消失了

接下来点击绿色按钮，进行合成

合成完成后，将生成报告

接下来，将硬件描述导出至IP目录中，点击下图所示按钮

在弹出的页面，保持默认点击OK即可

但，可能会出现以下问题，我导出IP核的时候经常碰到这个问题

但是呢，也不要慌，在我查了许多资料后，发现可以这样操作

我们在导出IP核时，点击Configuration

将Version改为0.0.0，点击OK，再点击OK，重新导出IP核

这个时候就会发现，我们已经成功导出了IP核

这样我们就完成了在HLS上的开发，请记住你IP核导出的位置

3.创建vivado项目

新建一个vivado项目

我把该项目命名为“ZynqOpenCL”，点击next

同样的，在这里选择自己开发版的型号，我的型号是xilinx的PYNQ Z2

点击finish即可完成项目的创建

在左侧状态栏中找到Create Block Design，直接保持默认后点击OK

点击中间的+号，输入ZYNQ，双击选择

现在我们可以点击一下上方的Run Block Automation，让模块自动化运行，按照下图步骤操作即可

得到下图所示

我们现在可以把之前导出的vadd的IP核导入至本项目中了，按下图步骤操作

点击+号后，我们要找到之前导出IP核所在的目录，找到 vadd_OpenCL ---》solution 的 “impl” 目录，点击Select

点击OK

再点击Apply再点击OK就可以成功导入vadd的IP核了

在主页面鼠标右击选择Add IP，输入vadd就可以找到我们刚刚导入的IP

双击该IP就可以导入了

双击ZYNQ Processing System图标

在左侧找到PS-PL Configuration

开启 “GP Slave AXI Interface”，点击OK即可

现在，我们将再次点击Run Connection Automation，选择下图所示按钮并点击OK

模块就会自动连线，如下图所示

在这个阶段，设计基本完成。但是，我们确实需要进入“地址编辑器”并进行一些小的调整。在地址编辑器中，某些地址范围被列为排除，我们需要将这些地址范围设置为包含。

将最后面两行用鼠标拖拽到上一栏合并，如下图所示

在顶栏中找到Validate design，点击它，如出现Validate sucessful即可

右击design_1，选择Create HDL Wrappers

根据默认点击OK

在创建好了之后，在左侧栏找到Run Synthesis，点击OK，等待运行完成

在运行后弹出的页面选择Run Implementation，然后点击OK

在运行后弹出来的界面选择Generate Bistream，点击OK

运行完毕后，我们可以点击Open Implemented Design查看一下生成的报告

在上述步骤都完成后，我们就可以将上述生成的比特流导出，步骤如下图所示，请注意记得勾选“Include bitstream”

然后点击“Launch SDK”

在左上角找到File，跟着下图步骤操作，点击New，选择Application Project

我把本次应用命名为“HelloOpenCL”

点击next 找到Hello World，选择Finish

在这里大家需要注意一下，我在网上看的教程里面是要将stdin和stdout的Value更改为ps7_uart_1的，但是我这里没找到有ps7_uart_1，所以我就保持默认，如果大家在stdin和stdout的Value找到有ps7_uart_1，请修改为ps7_uart_1。

接下来，首先在左侧单击HelloOpencl，随后在顶部栏中找到xilinx，选择其中的Generate  linker script

将其中的Heap Size更改为33554432，大概就是32MB，点击Generate

在运行vadd的代码前，我们可以先测试一下helloworld代码

快捷键 Ctrl + b 然后点击绿色运行按钮

点击第一个选项 “Launch on Hardware”

我们可以通过Mobaxtream软件，连接FPGA来输出字符

接下来可以尝试实现vadd的代码，将下述代码覆盖到原来的代码中

#include <stdlib.h>
#include "platform.h"

#include "xil_io.h"
#include "xil_mmu.h"
#include "xil_cache.h"
#include "xil_cache_l.h"
    
volatile char *control = (volatile char*)0x43C00000;
    
volatile int *wg_x   = (volatile int*)0x43C00010;
volatile int *wg_y   = (volatile int*)0x43C00018;
volatile int *wg_z   = (volatile int*)0x43C00020;
volatile int *o_x    = (volatile int*)0x43C00028;
volatile int *o_y    = (volatile int*)0x43C00030;
volatile int *o_z    = (volatile int*)0x43C00038;
    
volatile int *a_addr = (volatile int*)0x43C00040;
volatile int *b_addr = (volatile int*)0x43C00048;
volatile int *c_addr = (volatile int*)0x43C00050;
    
    
#define WG_SIZE_X 128
#define WG_SIZE_Y 1
#define WG_SIZE_Z 1
    
int main()
{
    init_platform();
    /* more initialization */
    Xil_SetTlbAttributes(0x43c00000,0x10c06);  /* non cacheable */
    
    int *a;
    int *b;
    int *c;
    int i;
    int ok = 1;
    
    a = (int*)malloc(WG_SIZE_X *sizeof(int));
    b = (int*)malloc(WG_SIZE_X *sizeof(int));
    c = (int*)malloc(WG_SIZE_X *sizeof(int));
    
    print("Generating input data: \n\r");
    for (i = 0; i < WG_SIZE_X; i ++) {
        a[i] = 1;
        b[i] = 2;
        c[i] = 0;
    }
    Xil_DCacheFlush();

    *a_addr = (unsigned int)a;
    *b_addr = (unsigned int)b;
    *c_addr = (unsigned int)c;

    /* set the workgroup identity */
    *wg_y = 0;
    *wg_z = 0;
    *wg_x = 0;

    *o_x = 0;
    *o_y = 0;
    *o_z = 0;


    print("Status of control register: \n\r");
    unsigned int con = *control;
    for (i = 0; i < 8; i ++) {
        if (con & (1  << i) ) {
            print("1");
        } else {
            print("0");
        }
    }
    print("\n\r");

    print("Starting OpenCL kernel execution\n\r");
    *control = *control | 1;  /* start */

    /* waiting for hardware to report "done" */
    while (! ((*control) & 2));
    print("DONE!\n\r");


    Xil_DCacheInvalidate();

    for (i = 0; i < WG_SIZE_X; i ++) {
        if (c[i] != 3) ok = 0;
    }

    if (ok) {
        print("Success!\n\r");
    } else {
        print("Error: Something went wrong!\n\r");
    }

    cleanup_platform();
    return 0;
}

 快捷键 Ctrl +s Ctrl + b 然后点击绿色运行按钮

就可以输出结果：

本文参考了以下链接：

Xilinx Zynq Opencl 入门指南 (svenssonjoel.github.io)https://svenssonjoel.github.io/pages/zynq_hls_opencl/index.html