MNN_POCL_VORTEX_GPGPU

MNN + POCL + Vortex GPGPU（simx)

1. 背景与目标

在标准 OpenCL 生态中，MNN 的 OpenCL Backend 通常运行在传统 GPU 驱动上。本文描述一套将 MNN OpenCL 计算路径接入 Vortex GPGPU（以 simx 作为验证后端）的工程验证方案。

目标

打通端到端执行链路：MNN -> POCL -> Vortex Runtime -> simx
保持可复现与可诊断：具备明确运行入口、日志与失败定位手段
适配无 Image 支持设备特征：在 no-image 条件下稳定运行 MNN OpenCL 子图
建立可回归基线：具备可批量运行的 smoke/tiny 测试矩阵

本工程当前阶段状态：

不追求一次性覆盖所有大模型/复杂图
不在本阶段展开完整性能优化
不在本文讨论硬件 FPGA/实体卡部署细节（以 simx 为主）

2. 总体架构

flowchart TD
    A[MNN
OpenCL Backend] --> B[POCL Runtime
OpenCL API + JIT + Cache]
    B --> C[POCL Vortex Device Plugin]
    C --> D[Vortex Runtime
libvortex + libvortex-simx]
    D --> E[simx Backend]

    B --> F[LLVM/Clang Codegen]
    F --> G[ELF Finalize + vxbin Packaging]
    G --> C

分层职责

应用层（MNN）：图执行、算子编排、OpenCL kernel 调度
运行时层（POCL）：OpenCL API、编译缓存、设备抽象、kernel 构建流程
设备层（Vortex Plugin）：设备能力暴露、内存传输、kernel 上传/启动
驱动层（Vortex Runtime）：统一 runtime API，转接到 simx
执行层（simx）：RISC-V/Vortex 指令执行与行为模拟

3. 执行流程（端到端）

sequenceDiagram
    participant U as MNN OpenCL Backend
    participant P as POCL Runtime
    participant V as Vortex Device Plugin
    participant R as Vortex Runtime
    participant S as simx

    U->>P: clBuildProgram / clEnqueueNDRangeKernel
    P->>P: LLVM/Clang 编译 OpenCL C
    P->>V: finalize_binary(input obj/bc)
    V->>V: entry 解析 + wrapper 生成 + ELF/vxbin 校验
    V->>R: vx_upload_kernel_file
    U->>V: enqueue kernel + args/buffers
    V->>R: vx_copy_to_dev / vx_start / vx_ready_wait
    R->>S: 执行 kernel
    S-->>R: 完成/状态
    R-->>V: 返回
    V-->>U: 结果可读

4. 实现原理（关键）

4.1 Vortex 设备后端执行链

Vortex 设备插件负责将 POCL 的抽象操作映射到 Vortex runtime API，核心包括：

设备打开与能力查询
buffer 分配与 host/device 数据传输
kernel 文件上传、启动与等待完成

该链路使“设备可枚举”升级为“kernel 可实际执行”。

4.2 Kernel Finalize 机制

为了避免“编译成功但运行失败/无法启动”，finalize 阶段做了三类保障：

入口符号自动解析
- 从编译产物中解析 _pocl_kernel_*_workgroup 入口
- 自动设置链接入口，避免人工硬编码
启动 Wrapper 语义适配
- 对齐运行时启动参数传递（如参数基址、线程上下文寄存器）
- 统一退出与完成语义，避免执行后悬挂
产物有效性校验
- 校验 ELF 是否具备可加载段（LOAD segment）
- 校验打包产物大小阈值，过滤空壳或损坏二进制

4.3 编译特征控制（Feature Control）

在实际排障中，存在“命令行传入特征”和“函数级 target-features 属性”不一致的问题。为保证稳定性，采用统一特征控制策略。

常用配置示例：

POCL_VORTEX_CODEGEN_FEATURES="+m,+f,+zicsr,-c"

含义：

+m：整数乘除扩展
+f：单精度浮点扩展
+zicsr：CSR 指令扩展
-c：禁用压缩指令（用于规避特定 simx 解码路径问题）

该策略用于确保目标特征在函数级别一致生效，避免局部函数“回退”到不期望的指令特征。

4.4 MNN 无 Image 设备适配

Vortex 当前路径中，设备能力呈现为 no-image。为避免 MNN 在编译或运行中误触 image kernel：

后端能力探测后自动降级到 BUFFER 路径
对 OpenCL 源中 image helper kernel 做条件编译保护（仅 image 支持时编译）
测试入口提供稳定 tuning 模式，避免在 simx 路径中卡于不稳定探测流程

4.5 simx 诊断机制

为提升可诊断性，在 simx 解码异常路径输出关键信息：

pc
code
opcode/funct
wid 等上下文

作用：快速把“执行崩溃”映射回“具体 kernel + 具体指令位置”。

5. 关键模块清单（按子系统）

5.1 POCL 侧关键模块

lib/CL/devices/vortex/vortex.c：Vortex 设备主实现（执行链核心）
lib/CL/devices/vortex/vortex_runtime.h：runtime API 头
lib/CL/pocl_llvm_build.cc：OpenCL 编译参数拼装
lib/CL/pocl_llvm_wg.cc：workgroup codegen 与 target machine 路径
lib/CL/devices/common.c：设备通用编译调用入口（含特征接入）
CMake 配置：Vortex runtime include/lib 参数接入

5.2 MNN 侧关键模块

source/backend/opencl/core/OpenCLBackend.cpp：模式选择与 fallback
source/backend/opencl/core/runtime/OpenCLRuntime.*：设备能力（image 支持）探测
source/backend/opencl/execution/cl/loop.cl：loop kernel（含 image helper 保护）
source/backend/opencl/execution/cl/loop_mnn_cl.cpp：loop kernel 使用链路
pocl_test/run_mnn_opencl_model.cpp：strict tiny 验证入口
pocl_test/env_vortex_simx.sh：一键环境脚本
pocl_test/run_tiny_matrix_simx.sh：一键回归脚本

5.3 Vortex 侧关键模块

sim/simx/decode.cpp：解码异常诊断增强

6. 测试策略

6.1 测试原则

隔离 ICD（不覆盖系统默认 OpenCL ICD）
固定 VORTEX_DRIVER=simx
分层验证（从底到上）：
- 设备可见性
- 最小 kernel
- MNN strict 子图
全流程保留日志与回归脚本

6.2 关键测试阶段

阶段A：设备与最小 kernel
- clinfo -l
- vecadd
阶段B：MNN strict tiny 单项
- tiny_matmul_add
阶段C：tiny 矩阵回归
- add / relu / reshape / mul1

7. 当前测试结果

用例	状态	备注
clinfo（simx + isolated ICD）	✅	可见 Vortex 设备
vecadd	✅	正常返回
tiny_matmul_add（strict）	✅	返回正确输出
tiny_matmul_add_relu（strict）	✅	通过
tiny_matmul_add_reshape（strict）	✅	通过
tiny_matmul_add_mul1（strict）	✅	修复后通过

8. 可复现执行方式

在 MNN 仓库中：

source ./pocl_test/env_vortex_simx.sh run1
./pocl_test/run_tiny_matrix_simx.sh

该方式可一键加载关键环境并执行 tiny 回归矩阵。

9. 名词解释（必要术语）

POCL：Portable OpenCL，OpenCL 运行时实现
ICD：Installable Client Driver，OpenCL 驱动分发机制
simx：Vortex 的 C++ 仿真后端
Finalization：将编译产物整理为可被目标 runtime 启动执行的阶段
vxbin：Vortex runtime 使用的打包二进制格式
strict 模式：MNN_STRICT_OPENCL_NO_CPU_OP=1，用于验证不回退 CPU 路径
target-features：LLVM 函数级目标特征属性
压缩指令（C 扩展）：RISC-V 16-bit 指令编码扩展

10. 参考仓库

POCL 适配仓：github/cecwxf/POCL_SE_FORK
MNN 适配仓：github/cecwxf/MNN_SE_FORK
Vortex 仓：github/cecwxf/VORTEX_SE_FORK