TIM-VX

1. 项目概述

1.1 定位与目标

TIM-VX（Tensor Interface Module for OpenVX）是 VeriSilicon 提供的神经网络部署软件集成模块，用于将神经网络模型高效部署到 VeriSilicon 的 ML 加速器（GPU / NPU）上。它作为上层推理框架（TensorFlow Lite、TVM、ONNX Runtime 等）与底层硬件驱动之间的桥梁层。

核心能力：

• 150+ 内置算子，支持量化与浮点格式
• 简化的 C++ API 用于创建 Tensor 和 Operation
• 动态图构建，支持形状推断和布局推断
• 内置自定义算子扩展框架
• 多设备支持与远程执行能力

1.2 框架集成生态

TIM-VX 已对接以下主流推理框架：

---

2. 系统架构

2.1 四层分层架构

┌──────────────────────────────────────┐
│            Application               │  上层框架 / 用户应用
├──────────────────────────────────────┤
│              TIM-VX                  │  C++ 封装层（本项目）
├──────────────────────────────────────┤
│              OVXLIB                  │  OpenVX C 封装库 (vsi_nn_*)
├──────────────────────────────────────┤
│          OpenVX Driver               │  VeriSilicon GPU/NPU 硬件驱动
└──────────────────────────────────────┘

• TIM-VX：面向应用和框架集成的 C++ API 层，提供类型安全的构图接口、图变换和平台抽象。
• OVXLIB：面向 OpenVX 驱动的 C 封装库，提供 vsi_nn_* 系列 API，管理图/节点/张量的底层生命周期。
• OpenVX Driver：VeriSilicon 对 Khronos OpenVX 标准的实现，含大量私有扩展（VX_*_VIV），直接驱动硬件。

2.2 全景架构

![[Pasted image 20260416162519.png]

Pasted image 20260416162600.png

全景架构展示了 TIM-VX 在整个软件栈中的位置：

左侧：上游推理框架（TF Lite、TVM、Tengine、OpenCV、PaddleLite 等）通过各自的适配器（External Delegate、BYOC、Backend、NNAdapter 等）接入 TIM-VX API。

中央绿色区域：TIM-VX 核心层，包含：

• TIM-VX API：统一的 C++ 公共接口
• Graph Transformation：图变换引擎（布局推断、算子融合）
• OpenVX VeriSilicon Extensions：对 OpenVX 标准的 VeriSilicon 私有扩展封装
• LiteExecutor：轻量执行器，通过 NBG 直接执行预编译模型

右侧：更多框架适配（XLA-NPU-JIT、Android NNAPI SupportLibrary、ONNX Runtime EP）。

2.3 双执行路径

TIM-VX 向下有两条并行的执行路径：

路径 A：Unified OpenVX SDK（支持 GPU + NPU）

OpenVX API → Compiler → Compute Kernels (GPU/NPU) → Runtime → HAL → 硬件

完整功能路径，含 JIT 编译器和 OpenVX 图处理引擎。适用于需要在线构图、动态编译的场景。

路径 B：VIP-Lite SDK（仅支持 NPU）

VIP-Lite API → Runtime → HAL → 硬件

轻量级路径，消费预编译的 NBG（Network Binary Graph）文件。设计用于 Linux、RTOS 乃至裸机环境，资源占用极小。

硬件平台层：支持 VeriSilicon IP 的各种 SoC，包括 Amlogic A311D/S905D、NXP iMX8mPlus 等，以及 x86 仿真器用于开发调试。

---

3. 目录结构

TIM-VX/
├── include/tim/                    # 公共 API 头文件
│   ├── vx/                         # 核心 API
│   │   ├── context.h               # Context 接口
│   │   ├── graph.h                 # Graph 接口
│   │   ├── tensor.h                # Tensor / TensorSpec / Quantization
│   │   ├── operation.h             # Operation 基类
│   │   ├── builtin_op.h            # 内置算子基类 (BuiltinOp / DirectMapOp)
│   │   ├── compile_option.h        # 编译选项
│   │   ├── types.h                 # 类型定义 (DataType, QuantType, DataLayout, ...)
│   │   ├── ops.h                   # 算子头文件聚合
│   │   └── ops/                    # 各算子公共头文件 + JSON 元数据
│   │       ├── conv2d.h
│   │       ├── pool2d.h
│   │       ├── custom_base.h       # 自定义 OpenCL 算子基类
│   │       └── ...                 # 150+ 算子
│   ├── transform/                  # 图变换 API
│   │   ├── layout_inference.h      # 布局推断入口
│   │   └── mean_stddev_normalize_fusion.h
│   ├── lite/                       # VIPLite 执行相关
│   └── experimental/trace/         # 实验性追踪/重放 API
│
├── src/tim/                        # 实现源码
│   ├── CMakeLists.txt              # 主构建脚本
│   ├── vx/                         # 核心实现
│   │   ├── context.cc / context_private.h
│   │   ├── graph.cc / graph_private.h
│   │   ├── tensor.cc
│   │   ├── operation.cc / op_impl.h / op_impl.cc
│   │   ├── builtin_op.cc / builtin_op_impl.cc
│   │   ├── type_utils.h / type_utils.cc   # 枚举翻译
│   │   ├── ops/                    # 各算子实现 + 单元测试
│   │   │   ├── conv2d.cc / conv2d_test.cc
│   │   │   ├── pool2d.cc
│   │   │   ├── rnn_cell.cc         # 组合算子示例
│   │   │   ├── custom_base.cc      # 自定义算子实现
│   │   │   └── ...
│   │   ├── platform/               # 平台抽象
│   │   │   ├── native.cc           # 本地设备枚举与执行
│   │   │   ├── lite/               # VIPLite 平台
│   │   │   └── grpc/               # gRPC 远程执行
│   │   └── internal/               # 内嵌 OVXLIB 源码
│   │       ├── tim_internal.cmake
│   │       ├── include/            # vsi_nn_*.h 头文件
│   │       └── src/                # OVXLIB 实现
│   │           ├── vsi_nn_graph.c / vsi_nn_node.c / vsi_nn_tensor.c
│   │           ├── ops/            # OVXLIB 算子实现
│   │           ├── kernel/         # 硬件内核 (cl/evis/vx)
│   │           ├── quantization/   # 量化工具
│   │           └── custom/ops/     # 驱动侧自定义算子
│   ├── transform/                  # 图变换实现
│   │   ├── layout_inference.cc     # 布局推断核心
│   │   ├── layout_infer_context.*  # 推断上下文
│   │   ├── permute_vector.*        # Permute 向量
│   │   └── ops/                    # 各算子的布局推断规则
│   ├── lite/                       # VIPLite 相关
│   └── utils/                      # 工具（NBG 解析等）
│
├── samples/                        # 示例应用
│   ├── lenet/                      # LeNet 推理示例
│   ├── custom_op_test/             # 自定义算子示例 (GEMM)
│   ├── multi_device/               # 多设备示例
│   ├── nbg_runner/                 # NBG 运行器
│   └── ...
│
├── prebuilt-sdk/                   # 预编译驱动 SDK
│   ├── x86_64_linux/               # x86 仿真环境
│   │   ├── include/VX/             # OpenVX 标准 + VSI 扩展头文件
│   │   └── lib/                    # 预编译库 (CLC, GAL, OpenVX, VSC, ...)
│   └── VIPLite/                    # VIPLite SDK
│
├── third_party/half/               # 半精度浮点库
├── CMakeLists.txt                  # 根构建脚本
└── docs/                           # 文档

---

4. 核心抽象与实现

TIM-VX 的核心对象模型由四个关键抽象组成：Context → Graph → Tensor / Operation。它们通过 pImpl（指针到实现）模式将公共 API 与底层 OVXLIB 实现解耦。

4.1 Context

Context 是 TIM-VX 的顶层运行环境，对应 OpenVX 的 vx_context。所有对象（图、张量、算子）都存活在 Context 域中。

公共接口 (include/tim/vx/context.h)：

class Context {
 public:
  static std::shared_ptr<Context> Create();
  virtual std::shared_ptr<Graph> CreateGraph() = 0;
  virtual std::shared_ptr<Graph> CreateGraph(const CompileOption& options) = 0;
  virtual bool isClOnly() = 0;   // 是否仅支持 OpenCL（无 EVIS 硬件加速）
  virtual bool hasSP() = 0;      // 是否支持流处理器
};

内部实现 (src/tim/vx/context.cc)：

class ContextImpl : public Context {
  vsi_nn_context_t context_;  // OVXLIB context 句柄
 public:
  ContextImpl() : context_(vsi_nn_CreateContext()) {}
  ~ContextImpl() { vsi_nn_ReleaseContext(&context_); }
  bool isClOnly() { return VSI_NN_HW_EVIS_NONE == context_->config.evis.ver; }
  bool hasSP()    { return 0 != context_->config.support_stream_processor; }
};

Context 是薄封装层，硬件能力查询直接读取 OVXLIB 的 vsi_nn_context_t 配置字段。

4.2 Graph

Graph 是 TIM-VX 的核心调度单元，表示一个有向无环计算图（DAG）。它管理张量和算子的创建，以及图的编译和执行。

公共接口 (include/tim/vx/graph.h)：

class Graph {
 public:
  // 张量创建
  virtual std::shared_ptr<Tensor> CreateTensor(const TensorSpec& spec, const void* data = nullptr);
  virtual std::shared_ptr<Tensor> CreateIOTensor(const TensorSpec& spec, void* data = nullptr);
  virtual std::shared_ptr<Tensor> CreateTensor(const TensorSpec& spec, const DmaBufferDesc& dmafd);
  virtual std::shared_ptr<Tensor> CreateTensorPlaceHolder();

  // 算子创建（模板方法）
  template <typename OpType, typename... Params>
  std::shared_ptr<OpType> CreateOperation(Params... parameters);

  // 图生命周期
  virtual bool Compile();
  virtual bool CompileToBinary(void* buf, size_t* size);
  virtual bool Run();

  // 图结构查询
  virtual const std::vector<std::shared_ptr<Tensor>> InputsTensor() const;
  virtual const std::vector<std::shared_ptr<Tensor>> OutputsTensor() const;
};

图生命周期：

CreateTensor/CreateOperation  →  Compile()  →  Run()
     (构建阶段)                (编译阶段)    (执行阶段)

内部实现关键流程 (src/tim/vx/graph.cc)：

GraphImpl 持有一个 vsi_nn_graph_t*（OVXLIB 图句柄），并在其上维护高层拓扑信息：

class GraphImpl : public Graph {
  ContextImpl* context_;
  vsi_nn_graph_t* graph_;                    // OVXLIB 图
  std::vector<std::shared_ptr<Operation>> op_vector_;  // 算子列表
  std::map<std::shared_ptr<Tensor>, std::vector<Operation*>> tensor_consumers_;
  std::map<std::shared_ptr<Tensor>, Operation*> tensor_producer_;
  std::vector<std::shared_ptr<Tensor>> inputs_tensor_;
  std::vector<std::shared_ptr<Tensor>> outputs_tensor_;
};

Setup() — 图初始化（通过 std::call_once 保证只执行一次）：

1. 设置图版本号（vsi_nn_SetGraphVersion）
2. 若 RelaxMode 开启，设置 fast 模式（vsi_nn_SetGraphFastMode，提示 float→bfloat16 优化）
3. 若多设备模式，设置设备索引（vxSetGraphAttribute(..., VX_GRAPH_DEVICE_INDEX_VIV, ...)）
4. 设置图的输入/输出张量（vsi_nn_SetGraphInputs / vsi_nn_SetGraphOutputs）
5. 执行图拓扑排序与资源分配（vsi_nn_SetupGraph）

Compile() — 图编译：

1. 检查并警告未消费的 INPUT/OUTPUT 张量
2. 调用 Setup()
3. 验证图的正确性（vsi_nn_VerifyGraph）

Run() — 图执行：

1. 调用 Compile()（幂等）
2. 执行推理（vsi_nn_RunGraph）

CompileToBinary() — 导出 NBG：

1. 调用 Setup()
2. 生成网络二进制图（vsi_nn_GenerateNBG）

常量张量缓存：当 TIM_VX_ENABLE_TENSOR_CACHE 开启时，GraphImpl 按张量的 TensorSpec + 数据摘要（小张量全量 MD5，大张量取前 512 字节）做去重缓存，避免重复创建相同的常量张量。

4.3 Tensor

Tensor 是多维数据对象，用于在算子之间传递数据。

TensorSpec — 张量规格描述：

struct TensorSpec {
  DataType datatype_;       // 数据类型 (INT8, UINT8, FLOAT16, FLOAT32, ...)
  ShapeType shape_;         // 形状 (Column-Major: WHCN 顺序)
  TensorAttribute attr_;    // 属性 (INPUT|OUTPUT|CONSTANT|TRANSIENT|VARIABLE)
  Quantization quantization_;  // 量化信息
};

TensorAttribute — 张量属性（位掩码）：

| 属性 | 含义 |

|------|------|

| CONSTANT | 常量张量，图编译前填充，之后不可变（权重、偏置等） |

| TRANSIENT | 虚拟张量，仅表示算子间连接，宿主不可访问（中间激活值） |

| VARIABLE | 可读写张量，可同时作为图的输入输出（RNN 状态等） |

| INPUT | 图输入张量，每次推理前由宿主更新 |

| OUTPUT | 图输出张量，每次推理后由宿主读取 |

Quantization — 量化描述：

class Quantization {
  QuantType type_;            // NONE / ASYMMETRIC / SYMMETRIC_PER_CHANNEL / DYNAMIC_FIXED_POINT / ...
  int32_t channel_dim_;       // Per-channel 量化的通道维度
  std::vector<float> scales_;
  std::vector<int32_t> zero_points_;
  int8_t fl_;                 // Dynamic fixed point 的小数位长度
};

内部实现关键点 (src/tim/vx/tensor.cc)：

PackTensorDtype 函数将 TIM-VX 的 TensorSpec / Quantization 打包为 OVXLIB 的 vsi_nn_dtype_t，处理了多种量化模式和 OVXLIB 版本兼容性差异（如旧版 OVXLIB 的 zero_points 需要 int32→float 转换）。

TensorImpl::Init 根据属性选择创建路径：

• TRANSIENT 张量：设置 dim_num = VSI_NN_DIM_AUTO，由 OVXLIB 自动推断形状
• INPUT/OUTPUT 且启用 ENABLE_TENSOR_HNDL：通过 vsi_nn_AddTensorFromHandle 创建（支持 DMA buffer）
• 其它：通过 vsi_nn_AddTensor 创建

数据搬运提供两条路径：

• Handle 路径：memcpy + FlushHandle / InvalidateHandle（零拷贝，适用于 DMA 场景）
• Copy 路径：vsi_nn_CopyDataToTensor / vsi_nn_CopyTensorToBuffer（通用拷贝）

4.4 Operation

Operation 是所有算子的基类，通过桥接模式将公共接口与底层实现解耦。

类层次：

Operation (公共接口)
  ├── BuiltinOp (内置算子基类, 别名 DirectMapOp)
  │     ├── Conv2d, Pool2d, FullyConnected, ...  (150+ 具体算子)
  │     └── CustomOpBase (自定义 OpenCL 算子基类)
  └── [组合算子, 如 RNNCell, 使用自定义 OpImpl]

OpImpl (内部实现接口)
  ├── BuiltinOpImpl (持有 vsi_nn_node_t*)
  ├── CustomOpBaseImpl (通过 vsi_nn_AddExternalNode)
  └── RNNCellImpl 等 (kind_=-1, 无单独 node)

Operation 公共接口：

class Operation {
 public:
  Operation& BindInput(const std::shared_ptr<Tensor>& tensor);
  Operation& BindOutput(const std::shared_ptr<Tensor>& tensor);
  Operation& BindInputs(const std::vector<std::shared_ptr<Tensor>>& tensors);
  Operation& BindOutputs(const std::vector<std::shared_ptr<Tensor>>& tensors);
  void SetRoundingPolicy(OverflowPolicy, RoundingPolicy, RoundType, uint32_t accumulator_bits);
  virtual std::shared_ptr<Operation> Clone(std::shared_ptr<Graph>& graph) const = 0;
};

绑定流程（src/tim/vx/operation.cc）：

BindInput 执行三步：

1. impl_->BindInput(tensor) — 将张量 ID 写入 OVXLIB node 的输入槽位
2. graph_->UpdateTensorConsumersMap(tensor, this) — 更新图级拓扑
3. OnBindInputPostProc(tensor, index) — 可选的后处理钩子（子类可覆盖）

---

5. 算子体系

项目中存在多个名为 ops 的目录，它们分别属于不同的软件层，各自解决不同的问题。理解这一点是理解 TIM-VX 架构的关键。

5.0 四层 ops 目录的关系

以 Conv2d 为例，一个算子的执行需要贯穿四个层次：

用户调用: graph->CreateOperation<Conv2d>(padding, stride, dilation, ...)
                          │
                          ▼
┌─────────────────────────────────────────────────────────┐
│  第1层: src/tim/vx/ops/conv2d.cc                         │
│  角色：C++ API 封装层 —— "说什么"                          │
│  ~160 个文件（含 *_test.cc）                               │
└─────────────────────────┬───────────────────────────────┘
                          │  graph->Compile() 触发布局推断
                          ▼
┌─────────────────────────────────────────────────────────┐
│  第2层: src/tim/transform/ops/conv2d_layout_inference.h  │
│  角色：布局推断规则层 —— "怎么摆"                           │
│  46 个文件                                                │
└─────────────────────────┬───────────────────────────────┘
                          │  最终走 OVXLIB 的 vsi_nn_SetupGraph
                          ▼
┌─────────────────────────────────────────────────────────┐
│  第3层: src/tim/vx/internal/src/ops/vsi_nn_op_conv2d.c   │
│  角色：OVXLIB 算子逻辑层 —— "怎么算"                       │
│  193 个文件                                               │
└─────────────────────────┬───────────────────────────────┘
                          │  kernel selector 选择硬件后端
                          ▼
┌─────────────────────────────────────────────────────────┐
│  第4层: src/tim/vx/internal/src/kernel/                   │
│  角色：硬件计算内核层 —— "用什么跑"                         │
│  ├── cl/   (80 个)  — OpenCL 通用 GPU shader              │
│  ├── evis/ (88 个)  — VeriSilicon 专有硬件加速指令          │
│  └── vx/   (27 个)  — OpenVX 标准 API 实现 (fallback)     │
└─────────────────────────────────────────────────────────┘

各层详细职责：

第1层 src/tim/vx/ops/ — C++ API 封装层（"用户怎么描述一个算子"）

每个文件对应一个算子的 C++ 封装类。以 Conv2d 为例，这一层做的事情是：

• 接收 C++ 参数（PadType padding, std::array<uint32_t,2> stride 等）
• 调用 vsi_nn_AddNode(graph, VSI_NN_OP_CONV2D, ...) 在 OVXLIB 图中创建节点
• 把参数翻译到 OVXLIB 的 C 结构体：node->nn_param.conv2d.stride[0] = stride_[0]
• 提供 Clone()、OnBindInputPostProc() 等钩子（如 FP16 bias→FP32 自动转换）

本质：从"C++ 类构造函数"到"OVXLIB C 结构体字段"的类型安全翻译。

第2层 src/tim/transform/ops/ — 布局推断规则层（"数据布局怎么对齐"）

解决的核心问题：上游框架（TensorFlow 用 NHWC/行主序）与底层硬件（OpenVX 用 WHCN/列主序）的布局差异。以 Conv2dLayoutInfer 为例：

• 检查输入 tensor 当前是 CWHN 还是 WHCN
• 决定是否需要插入 Permute 算子
• 检查 kernel weight 的布局（IWHO / OHWIn / ...），决定是否重排常量数据
• 在推断图中 clone 算子并绑定新 tensor

不是每个算子都需要专门的布局规则——没有注册的走 DefaultLayoutInfer，所以这一层只有 46 个文件，远少于其它层。

第3层 src/tim/vx/internal/src/ops/ — OVXLIB 算子逻辑层（"算子的计算逻辑如何调度"）

这是 OVXLIB 库的核心。每个算子实现三到四个回调函数：

• op_compute：将算子参数打包成 kernel_param，分发给 kernel selector 选择硬件后端
• op_setup：推导输出张量的形状（shape inference）
• op_check：检查输入数据类型/量化约束是否满足
• op_optimize：可选的图级优化

此层文件数（193）多于第1层（160），因为 OVXLIB 有一些内部算子（如 *_internal、pre_process_*）不暴露给 TIM-VX 用户，但底层执行需要。

第4层 src/tim/vx/internal/src/kernel/ — 硬件计算内核层（"在具体硬件上实际执行"）

分为三个子目录，对应三种硬件后端：

OVXLIB 的 kernel selector（vsi_nn_kernel_selector.c）在运行时根据硬件能力自动选择最优后端：优先 EVIS → 其次 CL → 最后 VX。EVIS 文件数最多，因为 VeriSilicon 的核心竞争力在于 NPU 硬件加速，EVIS 内核是针对自家硬件深度优化的。

为什么需要这么多层？

每一层解决一个独立的关注点（Separation of Concerns）：

这种分层的实际好处：新增一款芯片只需要改第4层的 kernel；新增一个框架的布局支持只需要改第2层的 transform；新增一个用户可见算子只需要改第1层。各层的变更互不影响。

---

5.1 算子扩展架构

Pasted image 20260416162703.png

上图展示了算子扩展的类层次设计。三种颜色的含义：

• 绿色：TIM-VX 公共 API（可被外部代码使用）
• 红色：可由用户在 TIM-VX 外部实现的组件
• 灰色：TIM-VX 内部私有实现

5.2 内置算子（BuiltinOp）

内置算子直接映射到 OVXLIB 的 vsi_nn_node_t，是最常用的算子类型。

实现模式（以 Conv2d 为例，src/tim/vx/ops/conv2d.cc）：

Conv2d::Conv2d(Graph* graph, PadType padding,
               std::array<uint32_t, 2> stride,
               std::array<uint32_t, 2> dilation,
               std::array<uint32_t, 4> pad,
               int32_t multiplier, DataLayout input_layout, DataLayout kernel_layout)
    : BuiltinOp(graph, VSI_NN_OP_CONV2D, 0, 0, input_layout) {
  this->impl()->node()->nn_param.conv2d.stride[0]   = stride_[0];
  this->impl()->node()->nn_param.conv2d.stride[1]   = stride_[1];
  this->impl()->node()->nn_param.conv2d.pad_type     = TranslatePadType(padding_);
  this->impl()->node()->nn_param.conv2d.group        = 1;
  this->impl()->node()->nn_param.conv2d.dilation[0]  = dilation_[0];
  this->impl()->node()->nn_param.conv2d.dilation[1]  = dilation_[1];
  this->impl()->node()->nn_param.conv2d.multiplier   = multiplier_;
  // ...
}

BuiltinOpImpl 的核心职责（src/tim/vx/builtin_op_impl.cc）：

• 构造时：调用 vsi_nn_AddNode(graph, kind, in_cnt, out_cnt, NULL) 在 OVXLIB 图中创建节点
• BindInput：将张量 ID 写入 node_->input.tensors[index]
• BindOutput：将张量 ID 写入 node_->output.tensors[index]
• SetRoundingPolicy：映射到 node_->vx_param 的 overflow/rounding 字段

类型翻译（src/tim/vx/type_utils.h）：

TIM-VX 的枚举类型与 OVXLIB / OpenVX 的枚举隔离，通过一组翻译函数映射：

• TranslateDataType — tim::vx::DataType → vsi_nn_type_e
• TranslateQuantType — tim::vx::QuantType → vsi_nn_qnt_type_e
• TranslatePadType — tim::vx::PadType → vsi_nn_pad_e
• TranslatePoolType、TranslateRoundType、TranslateResizeType 等

算子级 Workaround 示例：

Conv2d 的 OnBindInputPostProc 检测到 FP16 常量 bias 时，自动转换为 FP32 常量张量：

void Conv2d::OnBindInputPostProc(const std::shared_ptr<Tensor>& tensor, int32_t index) {
  if (tensor->GetDataType() == vx::DataType::FLOAT16 &&
      tensor->IsConstTensor() && impl_->inputs_tensor_.size() == 3) {
    float* float32_bias = tensor->ConvertTensorToFloat32Data();
    // 创建 FP32 常量张量替换原 bias
  }
}

5.3 组合算子（Composed Op）

组合算子不映射到单个 OVXLIB 节点，而是在 TIM-VX 层用多个内置算子构建子图。

实现模式（以 RNNCell 为例，src/tim/vx/ops/rnn_cell.cc）：

class RNNCellImpl : public OpImpl {
  // kind_ = -1（无单一 OVXLIB node）
  // node() 返回 nullptr

  // 内部持有多个内置算子
  std::shared_ptr<Operation> fc0_, fc1_;  // FullyConnected
  std::shared_ptr<Operation> add_;        // Add
  std::shared_ptr<Operation> tanh_;       // Tanh
  std::shared_ptr<Operation> data_convert_; // DataConvert

  void BindInput(const std::shared_ptr<Tensor>& tensor) override {
    // 收齐所有输入后，创建 TRANSIENT 中间张量并连接子图
  }
  void BindOutput(const std::shared_ptr<Tensor>& tensor) override {
    // 将最终算子的输出接到外部 tensor
  }
};

组合算子的关键设计：

• kind_ 设为 1，表示无对应的 OVXLIB 算子枚举
• 内部子图通过 TRANSIENT 张量连接
• 绑定外部输入/输出时在内部完成子图接线
• 布局推断时，由于 kind_ != -1 的过滤条件，组合算子不直接进入 HandleLayoutInfer，需要依赖子算子的布局推断

5.4 自定义 OpenCL 算子（Custom Op）

当内置算子不能满足需求时，用户可以通过 OpenCL 内核实现自定义算子。

基类接口（include/tim/vx/ops/custom_base.h）：

class CustomOpBase : public Operation {
 public:
  // 用户需要实现的纯虚函数
  virtual void SetupShapeInfor() = 0;     // 输出张量的形状推导
  virtual void SetupParams(              // 选择内核函数和编译选项
      std::vector<tim::vx::DataType> input_types,
      std::string& build_option) = 0;
  virtual void SetupEnqueue(             // 配置 global/local work size
      uint32_t& dim,
      std::vector<size_t>& global_size,
      std::vector<size_t>& local_size) = 0;

 protected:
  ParamTuple tuple_list_;       // 标量参数元组
  std::string kernel_resource_; // OpenCL 内核源码字符串
  std::string func_name_;       // 选中的内核函数名
};

标量参数机制：

算子参数分为两类：

• Tensor-like 参数：通过 BindInput / BindOutput 传递
• 标量参数：通过 std::tuple 定义，经 param_transform 模板函数打包为 std::vector<Param>

Param 是一个类型标记联合体：

struct Param {
  enum DataType { FLOAT, INT32, ... } type;
  union { float f; int32_t i; void* p; } data;
};

内部实现（src/tim/vx/ops/custom_base.cc）：

通过 vsi_nn_AddExternalNode 将自定义算子注册到 OVXLIB 图中，绑定 op_setup / op_compute 回调：

• op_setup：从 OVXLIB 张量收集输入尺寸，调用用户的 SetupShapeInfor，写回输出形状
• op_compute：创建 OpenCL kernel，调用 SetupParams 选择内核函数，打包标量参数，配置 enqueue

使用示例（samples/custom_op_test/custom_gemm.h）：

class CustomGemm : public CustomOpBase {
  using ParamTuple = std::tuple<int, int, int, int, int>;

  CustomGemm(Graph* graph, ParamTuple params, uint32_t in_num, uint32_t out_num)
      : CustomOpBase(graph, in_num, out_num, kernel_id_, kernel_name_) {
    tuple_list_.swap(params);
    param_transform(tuple_list_, param_list_);
    kernel_resource_ = "__kernel void gemm_F32toF32(...) { ... }";
  }

  void SetupShapeInfor() override { /* 根据 M/K/N 设置输出尺寸 */ }
  void SetupParams(...) override  { /* 选择 func_name_ 和 build_option */ }
  void SetupEnqueue(...) override { /* 设置 global/local work size */ }
};

---

6. 图变换与布局推断

6.1 问题背景

上游框架（TensorFlow、PyTorch、ONNX）使用 Row-Major 行主序，维度描述为 NHWC 或 NCHW。而 VeriSilicon 的 OpenVX 驱动使用 Column-Major 列主序，维度描述为 WHCN。

这不仅是维度翻转——每种算子的权重布局、padding 语义、axis 含义都会因此受影响。如果不做自动转换，框架适配者需要在每个算子的每个参数上手工处理布局差异。

6.2 实现机制

布局推断的入口是 tim::transform::LayoutInference（include/tim/transform/layout_inference.h）：

std::pair<
  std::shared_ptr<vx::Graph>,
  std::map<std::shared_ptr<vx::Tensor>, std::shared_ptr<vx::Tensor>>
>
LayoutInference(
    const std::shared_ptr<vx::Graph>& src_graph,
    std::shared_ptr<vx::Context>& ctx,
    const std::map<std::shared_ptr<vx::Tensor>,
                   std::shared_ptr<IPermuteVector>>& tensor_pv_map = {});

输入：源图 + Context + 可选的 per-tensor PermuteVector 映射

输出：推断后的新图 + 原始张量到新张量的映射

核心流程（src/tim/transform/layout_inference.cc）：

1. 初始化推断图：创建新的空图 infer_graph
2. 处理图输入：为每个 INPUT 张量在推断图中创建对应张量，设置默认或传入的 PermuteVector，加入 BFS 队列
3. 处理常量输入：从源张量复制数据到推断图的新常量张量
4. 处理图输出：在推断图中创建对应的 OUTPUT 张量
5. BFS 遍历：对队列中每个张量的 consumer 算子，若所有非常量输入都已具备 PermuteVector，则调用 HandleLayoutInfer：
   • 按 op->impl()->kind_ 通过宏 REGISTER_LAYOUT_INFERENCE 分派到各算子专属的 LayoutInfer 实现
   • 未注册的算子走 DefaultLayoutInfer
   • 返回新产出的张量，加入 BFS 队列继续传播
6. 返回结果：合并 input/output 的张量映射

6.3 算子级布局规则

每种算子在 src/tim/transform/ops/ 下有独立的布局推断实现。以 Conv2d 为例：

class Conv2dLayoutInfer {
  void OnInputs(std::vector<std::shared_ptr<vx::Tensor>> next_tensors) {
    // 根据 DataLayout / KernelDataLayout 选择 permute 表
    // 必要时对输入张量 InsertPermute
    // 对常量 weight 做 PermuteConstTensor（重排数据）
    // Clone op 到推断图并绑定新张量
  }
};

src/tim/transform/ops/ 目录下包含几十个算子的布局推断规则文件。

---

7. 平台抽象与多设备支持

7.1 设备枚举

src/tim/vx/platform/native.cc 提供本地设备的发现能力：

std::vector<std::shared_ptr<IDevice>> NativeDevice::Enumerate() {
  vsi_nn_context_t context = vsi_nn_CreateContext();
#ifdef VSI_DEVICE_SUPPORT
  vsi_nn_GetDevices(context, vsi_devices, &deviceCount);
  for (uint32_t i = 0; i < deviceCount; i++) {
    vsi_nn_GetDeviceCoreCount(vsi_devices[i], &available_core_count);
    device_v.push_back(std::make_shared<NativeDeviceImpl>(i, available_core_count));
  }
#else
  vxQueryContext(context->c, VX_CONTEXT_DEVICE_COUNT_VIV, &deviceCount, sizeof(deviceCount));
  // 兼容旧 SDK 的回退路径
#endif
  vsi_nn_ReleaseContext(&context);
  return device_v;
}

7.2 编译选项

CompileOption（include/tim/vx/compile_option.h）控制图编译行为：

• setRelaxMode(bool) — 开启后允许 float→bfloat16 优化，对应 OVXLIB 的 vsi_nn_SetGraphFastMode
• setDeviceId(device_id_t) — 指定目标设备（需 TIM_VX_ENABLE_PLATFORM）

7.3 NBG 执行路径

NBG（Network Binary Graph）是 VeriSilicon 的预编译模型格式。TIM-VX 提供两种 NBG 使用方式：

1. 导出 NBG：通过 Graph::CompileToBinary(buf, size) 将在线构建的图导出为 NBG 格式
2. 加载 NBG：通过 NativeExecutableImpl 加载 NBG 并创建 executor 执行推理

VIPLite SDK 路径专门针对 NBG 执行场景优化，适用于资源受限的嵌入式环境。

7.4 gRPC 远程执行

src/tim/vx/platform/grpc/ 实现了基于 gRPC 的远程执行能力，将编译出的 NBG + Executor 运行在远端进程。这是 TIM-VX 在编排和部署层面的扩展，不属于 OVXLIB 的核心功能。

---

8. 构建系统

8.1 CMake 构建选项

8.2 算子特性宏生成

构建系统从 ops.def / custom_ops.def 文件解析 DEF_OP(...) 宏定义，生成 -DVSI_FEAT_OP_xxx 编译标志。这确保 TIM-VX 与 OVXLIB 的算子能力表保持一致，避免编译时引用不存在的算子枚举。

8.3 外部依赖

---

9. 与底层 OVXLIB 的关系

9.1 调用链路

TIM-VX C++ API
  │
  │  Context::Create()            →  vsi_nn_CreateContext()
  │  Graph::CreateTensor()        →  vsi_nn_AddTensor() / vsi_nn_AddTensorFromHandle()
  │  Graph::CreateOperation<T>()  →  vsi_nn_AddNode(kind) + 填充 nn_param
  │  BindInput/Output()           →  node->input.tensors[i] = tensor_id
  │  Graph::Compile()             →  vsi_nn_SetupGraph() + vsi_nn_VerifyGraph()
  │  Graph::Run()                 →  vsi_nn_RunGraph()
  │  Graph::CompileToBinary()     →  vsi_nn_GenerateNBG()
  │  ~Graph()                     →  vsi_nn_ReleaseGraph()
  │  ~Context()                   →  vsi_nn_ReleaseContext()
  │
  ▼
OVXLIB (vsi_nn_* C API)
  │
  ▼
OpenVX Driver (vx_* + VIV 扩展)
  │
  ▼
GPU / NPU 硬件

少数场景 TIM-VX 直接调用 OpenVX API（绕过 OVXLIB），例如多设备属性设置和设备查询。

9.2 TIM-VX 的附加价值

---

10. 典型使用流程

以 LeNet 推理为例（samples/lenet/lenet_asymu8.cc）：

// 1. 创建 Context 和 Graph
auto ctx = tim::vx::Context::Create();
auto graph = ctx->CreateGraph();

// 2. 创建输入/输出张量
auto input = graph->CreateTensor(
    tim::vx::TensorSpec(tim::vx::DataType::UINT8, {28, 28, 1, 1},
                        tim::vx::TensorAttribute::INPUT,
                        tim::vx::Quantization(QuantType::ASYMMETRIC, 0.00390625f, 0)));

auto output = graph->CreateTensor(
    tim::vx::TensorSpec(tim::vx::DataType::FLOAT32, {10, 1},
                        tim::vx::TensorAttribute::OUTPUT));

// 3. 创建常量张量（权重/偏置），带数据指针
auto conv1_weight = graph->CreateTensor(weight_spec, weight_data);
auto conv1_bias   = graph->CreateTensor(bias_spec, bias_data);

// 4. 创建中间 TRANSIENT 张量（形状可自动推断）
auto conv1_out = graph->CreateTensor(
    tim::vx::TensorSpec(tim::vx::DataType::UINT8, {},
                        tim::vx::TensorAttribute::TRANSIENT, quant));

// 5. 创建算子并连接
auto conv1 = graph->CreateOperation<tim::vx::ops::Conv2d>(
    /*weights=*/20, PadType::VALID, /*ksize=*/{5,5}, /*stride=*/{1,1});
conv1->BindInput(input).BindInput(conv1_weight).BindInput(conv1_bias);
conv1->BindOutput(conv1_out);

// ... 继续构建 Pool、FC、Relu、Softmax 等

// 6. 编译图
graph->Compile();

// 7. 填充输入数据
input->CopyDataToTensor(input_data);

// 8. 执行推理
graph->Run();

// 9. 读取输出
output->CopyDataFromTensor(result_data);