再谈Xilinx FPGA开发的Vivado流和Vitis流

前景提要：FPGA加速计算生态系统：从Vivado到Vitis AI的全面解析

在 AMD (Xilinx) 最新的 Vitis 统一软件平台架构下，FPGA 的开发被清晰地划分为两个世界。理解这两个世界的区别，是看懂现代 FPGA 论文和技术文档的关键。

逻辑一：自底向上的“硬件构建”逻辑 (Vivado Flow)

——“我是芯片设计师，我在造一台电脑。”

这是经典的、传统的 FPGA 开发方式。

1. 核心思维

• 关注点：时钟周期、复位信号、状态机翻转、AXI 总线位宽、时序收敛 (Timing Closure)。
• 目标：生成一个 .bit 文件（Bitstream），控制 FPGA 上的每一个查找表 (LUT) 和触发器 (FF)。
• 角色：你是上帝。你需要亲自搭建 CPU（MicroBlaze）或配置硬核（ARM），你需要亲自连线内存控制器 (DDR Controller)，你需要亲自处理 PCIe 的物理层接口。

2. HLS 在这里的角色 (Vivado HLS / Vitis HLS)

在这种逻辑下，HLS (High-Level Synthesis) 只是一个**“零件加工厂”**。

• 你用 C++ 写一个算法。
• HLS 把它综合成一个 RTL IP 核（带 AXI 接口的黑盒子）。
• 关键步骤：你必须手动打开 Vivado 的 IP Integrator (Block Design)，把这个 HLS 生成的 IP 拖进去，画线连接到总线上，分配地址空间。
• 局限：如果你想修改算法的并行度（比如从 1 路并行改成 4 路），你可能需要重新导出 IP，回到 Vivado 更新连线，重新综合整个工程。这很难自动化。

逻辑二：自顶向下的“软件加速”逻辑 (Vitis Flow)

——“我是系统架构师，我在写一个异构程序。”

这是 OpenCL 所在的领域，也是那篇论文使用的逻辑，更是 AMD 目前大力推崇的 Vitis 加速流。

1. 核心思维

• 关注点：数据吞吐量、内核 (Kernel) 延迟、主机-设备 (Host-Device) 带宽、异构调度。
• 目标：生成一个 .xclbin 文件（二进制容器）和一个主机可执行文件 (Host Exe)。
• 角色：你是应用开发者。底层的 PCIe 接口、DDR 控制器、DMA 搬运引擎，你不关心，也不需要写。这些由 Xilinx 提供的（或板卡厂商提供的）“Shell” (Target Platform) 自动完成。

2. 这种逻辑下的开发架构 (Host + Kernel)

在这种模式下，FPGA 不再被视为“电路板”，而被视为一个**“可编程的函数库”**。

• Host (主机)：CPU 运行的代码。负责读取文件、预处理数据、通过 API 指挥 FPGA 干活。
• Kernel (内核)：FPGA 内部的加速逻辑。

3. OpenCL 在这里的角色

OpenCL 是这一逻辑的“鼻祖”和“标准”。

• 它规定了 CPU 怎么给 FPGA 发数据（clEnqueueWriteBuffer）。
• 它规定了 FPGA 里的代码该长什么样（__kernel void ...）。
• 最重要的是：它实现了软硬件的解耦。你改了 FPGA 里的算法（Kernel），只要接口没变，CPU 的代码甚至不需要重新编译。

深度解析：AMD (Xilinx) 最新的开发逻辑 (Vitis Unified)

“OpenCL 写 FPGA”其实是 Vitis Flow 的早期形态（当时叫 SDAccel）。现在，AMD 已经将这套逻辑进化得更加灵活。

在最新的 Vitis 2023/2024 版本中，开发逻辑如下：

1. 硬件描述语言的解放：C++ Kernel 取代 OpenCL C

虽然 Vitis 依然完美支持 OpenCL C 语言写内核，但 AMD 现在更推荐用 C++ 来写内核（Vitis HLS 风格）。

• 以前 (OpenCL C)：你需要写很多 restrict, __global 等修饰符，语法受限。
• 现在 (C++ Kernel)：你可以直接用 C++ 写算法，配合 #pragma HLS 指令。Vitis 编译器会自动把它包装成符合 OpenCL 调用规范的内核（自动加上 AXI 接口）。
• 结论：内核是用 C++ 写的，但调用模型依然是 OpenCL 的逻辑。

2. 运行时库的进化：XRT (Xilinx Runtime)

这是连接软件和硬件的桥梁。

• 那篇论文还在用标准的 OpenCL API (clCreateContext 等)。
• 现在的 Vitis 推荐使用 XRT Native API。它比 OpenCL API 更轻量、更底层、开销更小，但做的事情是一样的：加载 .xclbin，分配显存，启动内核。

3. 如何 “通过 OpenCL 进行 DSE”？

在 Vivado Flow (逻辑一) 里，要改一个设计：

修改 RTL -> 更新 Block Design -> 重新连线 -> 综合 (1小时) -> 布局布线 (2小时) -> 失败重来。

在 Vitis Flow (逻辑二) 里，要改一个设计：

打开 C++ Kernel 代码。修改一行参数：const int PARALLEL_FACTOR = 8; (原本是4)。运行脚本 v++ --link ...。编译器自动调用 HLS 生成新 IP，自动调用 Vivado 重新布局布线，自动生成新的 .xclbin。

这就是 OpenCL/Vitis Flow 的威力。它允许研究人员像编译软件一样，通过脚本批量生成 100 种不同资源占用率的 FPGA 镜像，从而画出论文里那种漂亮的“面积 vs 吞吐量”曲线图。

总结：如何选择你的路线？

如果你是…

• 做板卡的：你需要画 PCB，调试 DDR 时序，把 PCIe 物理层跑通。
  • 选择：Vivado Flow (逻辑一)。这是基础，你要构建“Shell”。
• 做算法加速的：你是搞 AI、图像处理、金融计算的，你买了一块现成的 Alveo 加速卡（比如 U200, U50），或者用 AWS F1 云实例。
  • 选择：Vitis Flow (逻辑二)。
  • 具体做法：
    1. Kernel: 用 C++ (HLS风格) 写算法，打上 #pragma 标签优化流水线。
    2. Link: 用 v++ 编译器把 Kernel 连入 Alveo 卡的 Shell。
    3. Host: 用 C++ 调用 XRT API 或 OpenCL API 控制运行。