一、背景：当客户说“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，迁移零成本。