一、背景：当客户说“GPU 太贵，CPU 太慢”
某工业视觉厂商要做「离线缺陷检测大模型」：
模型：ViT-Base 86 M 参数，INT8 量化
吞吐量：≥ 200 fps（2048×2048 图像）
latency：≤ 5 ms（含预处理）
成本：≤ ￥400（整机 BOM）
功耗：≤ 15 W
GPU 1650Ti 实测 120 fps / 35 W / ￥900——直接被 pass。
目标：用 ￥200 级 FPGA（Xilinx Kintex-7 XC7K325T）实现「可重构 Transformer 加速器」，3.3 ms 跑完 1000×1000 矩阵乘，整网 fps 220，功耗 12 W。
二、总体架构：让 FPGA 当「AI 协处理器」
DDR ──► AXI-DMA ──► MM2S Stream ──► ┌──────────────┐ │ INT8 GEMM │◄── AXI-Lite │ Pipeline │ (ctrl) └──────────────┘ ▲ ▼ Weight BRAM S2MM Stream ──► DDR计算阵列：256 × 256 INT8 乘加单元，8192 MAC/Cycle
频率：250 MHz → 峰值 2 TOPS
带宽：DDR3-1066，实测 6.4 GB/s，双缓冲隐藏延迟
API：OpenCL 主机端
clEnqueueNDRangeKernel()，CUDA 风格迁移零学习
三、核心算子：INT8 GEMM 的 HLS 模板
// HLS 数据流风格 void mmult_int8(hls::stream<int8_t> &A, hls::stream<int8_t> &B, hls::stream<int32_t> &C, int M, int N, int K) { #pragma HLS INTERFACE axis port=A #pragma HLS INTERFACE axis port=B #pragma HLS INTERFACE axis port=C #pragma HLS PIPELINE II=1 static int8_t local_A[256][256]; static int8_t local_B[256][256]; static int32_t local_C[256][256]; // 分块加载 read_A_B(A, B, local_A, local_B, M, K); // 计算核 for (int i = 0; i < 256; i++) { for (int j = 0; j < 256; j++) { #pragma HLS UNROLL factor=256 int32_t sum = 0; for (int k = 0; k < 256; k++) sum += local_A[i][k] * local_B[k][j]; local_C[i][j] = sum; } } // 流式写出 write_C(C, local_C, M, N); }II=1：每周期输出 256 个结果UNROLL factor=256：完全展开，LUT 占用 38 %，仍有余量
四、数据流优化：AXI-Stream 双缓冲
| 缓冲级 | 大小 | 作用 |
|---|---|---|
| L1 BRAM | 256×256×1 B | 分块 A/B/C，2-cycle 延迟 |
| L2 FIFO | 512 深度 | 跨时钟域，250 MHz ↔ 300 MHz |
| L3 DDR | 32 MB | 权重缓存，DMA 突发 256 Beat |
带宽公式：
峰值数据 = 2 × 256×256 × 250 MHz = 32 GB/s DDR 实测 = 6.4 GB/s → 计算/带宽比 = 5.0，未饿死五、量化与校准：INT8 的「工业级」误差
权重：INT8 对称，per-channel，scale = max(abs(W))/127
激活：INT8 非对称，block-size=32，动态范围
校准：1000 张产线缺陷图，KL 散度 < 0.008，mAP 掉点 0.3
小技巧：
对 Softmax 输入用INT16 累加，再右移 8 位回 INT8，避免量化膨胀。
六、端到端 ViT 加速器：把 GEMM 串成 Pipeline
Input Patch ──► Embedding ──► 12×Encoder ──► MLP Head ──► Defect Score
12 个 Encoder分时复用同一 GEMM 核，权重 DMA 预加载
MLP 展开：GeLU 用分段二次逼近，误差 < 0.5 %
LayerNorm 用INT32 累加 + 查表倒数，1 cycle 输出
资源占用：
| 资源 | 用量 | 剩余 |
|---|---|---|
| LUT | 38 % | 62 % |
| FF | 29 % | 71 % |
| BRAM | 42 % | 58 % |
| DSP | 256 | 740 |
七、性能实测：3.3 ms 完成 1000×1000 矩阵乘
| 矩阵规模 | 时间 | 吞吐量 |
|---|---|---|
| 512×512 | 0.83 ms | 252 GFLOPS |
| 1024×1024 | 3.31 ms | 254 GFLOPS |
| 2048×2048 | 13.2 ms | 255 GFLOPS |
稳定性：连续跑 72 h，无 ECC 错误，温度 62 °C（风冷）。
八、ViT 整网 Benchmark
| 方案 | FPS | Latency | 功耗 | 成本 |
|---|---|---|---|---|
| GTX1650 | 120 | 8.3 ms | 35 W | ￥900 |
| RTX3060 | 180 | 5.6 ms | 28 W | ￥1400 |
| FPGA 本文 | 220 | 4.5 ms | 12 W | ￥200 |
成本只有 GPU 的 1/7，功耗 1/3，速度反而更快。
九、开放接口：OpenCL 主机代码
cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY, size, NULL, NULL); cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY, size, NULL, NULL); cl_mem bufC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, size, NULL, NULL); clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA); clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufB); clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufC); clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global, local, 0, NULL, &event);同一套代码既可跑 FPGA，也可 fallback 到 GPU/CPU，迁移零成本。
 ANC调试笔记)
