Cchenqian1365replace remote with temp docs

GE Runtime 2.0 定位指导文档

本文档旨在帮助开发者快速定位 RT2.0 动态 Shape 执行器的问题，涵盖主要流程图、KernelTrace 定位机制、外部接口说明及关键定位技巧。

1. RT2.0 概述

1.1 设计背景

RT2.0（runtime/v2/）是 GE 的下一代动态 Shape 执行器，通过 Lowering 机制将高层 ComputeGraph 转换为可直接执行的 ExecuteGraph，使运行时只需一个极简的执行循环，以更低的开销原生支持动态 Shape 场景。

核心设计理念：将"翻译"工作从运行时移到编译期（Lowering 阶段）。

1.2 适用场景

场景	说明
动态 Shape 推理	NLP 变长序列、视觉模型动态分辨率等 shape 在推理时变化的场景
单算子执行	PyTorch 动态图场景下的单算子编译执行
控制流模型	包含 If/While 等控制流的模型

2. RT2.0 主要流程图

2.1 全链路执行流程

flowchart TD
    subgraph offline_load["离线场景：加载阶段"]
        direction TB
        A1["aclmdlLoadFromFile<br/>aclmdlLoadFromMem"]
        FD1{"AclIsOm2ModelByPath<br/>检查文件头魔数"}
        
        subgraph om2_load["OM2 路径"]
            O1["aclmdlLoadFromFileImplOm2"]
            O2["LoadOm2ExecutorFromData<br/>om2_model_executor.cc:884"]
            O3["Om2ModelExecutor::Load<br/>加载 SO 并初始化"]
        end
        
        subgraph non_om2_load["非 OM2 路径"]
            N1["aclmdlLoadFromFileImpl"]
            FD2{"IsSupportRuntimeV2<br/>IsDynamicModel"}
            N2["RuntimeV2ModelLoadCommon<br/>model.cpp:232"]
            N3["LoadExecutorFromModelData<br/>api.cc:102"]
            N4["ModelV2Executor::Load<br/>执行 Init Graph"]
        end
        
        V1["ModelLoadFromFileWithMem<br/>→ v1 TaskSink"]
    end

    A1 --> FD1
    FD1 -- "OM2(ZIP格式)" --> O1 --> O2 --> O3
    FD1 -- "非OM2" --> N1 --> FD2
    FD2 -- "动态Shape" --> N2 --> N3 --> N4
    FD2 -- "静态Shape" --> V1

    classDef acl fill:#c8e6c9,stroke:#2e7d32
    classDef detect fill:#fff9c4,stroke:#f9a825
    classDef om2 fill:#e3f2fd,stroke:#1565c0
    classDef rt2 fill:#fff3e0,stroke:#f57c00
    classDef v1 fill:#d7ccc8,stroke:#795548,stroke-dasharray:4 4

    class A1 acl
    class FD1,FD2 detect
    class O1,O2,O3 om2
    class N2,N3,N4 rt2
    class V1 v1

离线加载关键点：

OM2 和非OM2动态模型使用不同执行器：Om2ModelExecutor vs ModelV2Executor
OM2 判断基于文件头魔数（ZIP格式：0x50, 0x4B, 0x03, 0x04）
动态Shape判断：version >= MODEL_VERSION 且（model_num > 1 或 is_unknow_model == 1）

flowchart TD
    subgraph offline_exec["离线场景：执行阶段"]
        direction TB
        A2["aclmdlExecuteV2"]
        FD3{"AclIsOm2ModelById<br/>查询已加载模型ID"}
        
        subgraph om2_exec["OM2 执行路径"]
            E1["aclmdlExecuteV2ImplOm2"]
            E2["Om2ModelExecutor::Run<br/>om2_model_executor.cc:805"]
            E3["调用 SO 中 Om2ModelRun"]
        end
        
        subgraph non_om2_exec["非 OM2 执行路径"]
            E4["aclmdlExecuteV2Impl"]
            FD4{"GetExecutor(modelId)<br/>!= nullptr?"}
            E5["RuntimeV2ModelExecute<br/>model.cpp:605"]
            E6["ModelV2Executor::Execute<br/>执行 Main Graph"]
        end
        
        V2["ModelExecute<br/>→ GeExecutor → DavinciModel"]
    end

    A2 --> FD3
    FD3 -- "OM2模型" --> E1 --> E2 --> E3
    FD3 -- "非OM2" --> E4 --> FD4
    FD4 -- "RT2.0动态模型" --> E5 --> E6
    FD4 -- "v1静态模型" --> V2

    classDef acl fill:#c8e6c9,stroke:#2e7d32
    classDef detect fill:#fff9c4,stroke:#f9a825
    classDef om2 fill:#e3f2fd,stroke:#1565c0
    classDef rt2 fill:#fff3e0,stroke:#f57c00
    classDef v1 fill:#d7ccc8,stroke:#795548,stroke-dasharray:4 4

    class A2 acl
    class FD3,FD4 detect
    class E1,E2,E3 om2
    class E5,E6 rt2
    class V2 v1

离线执行关键点：

执行时通过查询已加载模型ID判断类型（不是重新判定文件格式）
OM2模型走 Om2ModelExecutor::Run → 调用 SO 中的符号
非OM2路径内部有二次分流：查询 GetExecutor(modelId) 判断是否是 RT2.0 动态模型

离线场景特点：Load / Execute / Unload 三阶段对应三个独立 ACL 接口，用户显式调用。

flowchart TD
    subgraph online["在线场景：RunGraph"]
        direction TB
        R["GeSession::RunGraph"]
        C{"首次调用?"}
        L["内部自动加载<br/>LoadExecutorFromModelData<br/>+ ModelV2Executor::Load"]
        I["执行 Init Graph"]
        E["ModelV2Executor::Execute<br/>执行 Main Graph"]
        N["直接执行 Main Graph"]
    end

    R --> C
    C -- "是" --> L --> I --> E
    C -- "否" --> N

    classDef entry fill:#c8e6c9,stroke:#2e7d32
    classDef check fill:#fff9c4,stroke:#f9a825
    classDef core fill:#fff3e0,stroke:#f57c00

    class R entry
    class C check
    class L,I,E,N core

在线场景特点：单一 RunGraph 接口，内部自动判断是否需要加载，首次调用完成 Load + Execute，后续只 Execute。

2.2 三子图生命周期

sequenceDiagram
    participant User as 用户
    participant Executor as ModelV2Executor
    participant Init as Init Graph
    participant Main as Main Graph
    participant DeInit as DeInit Graph
    participant Device as NPU Device
    
    rect rgb(200, 230, 200)
        Note over User,Device: Load 阶段
        User->>Executor: Load(ModelExecuteArg, ModelLoadArg)
        Executor->>Init: Load()
        Init->>Device: 分配内存/流/事件/通知
        Executor->>Init: SpecifyInputs + Execute()
        Init->>Device: 执行初始化节点
        Executor->>Init: UnLoad()
        Executor->>Main: Load()
    end
    
    rect rgb(255, 230, 200)
        Note over User,Device: Execute 阶段（可多次）
        User->>Executor: Execute(arg, inputs, outputs)
        Executor->>Main: SpecifyInputs + SpecifyOutputs
        Main->>Device: 执行算子节点
        Device-->>Main: 执行完成
        Main-->>Executor: 返回结果
    end
    
    rect rgb(255, 200, 200)
        Note over User,Device: UnLoad 阶段
        User->>Executor: UnLoad()
        Executor->>Main: UnLoad()
        Executor->>DeInit: Load()
        DeInit->>Device: 释放资源
        Executor->>DeInit: Execute()
        Executor->>DeInit: UnLoad()
    end

2.3 Lowering 详细流程

flowchart LR
    subgraph input_block["输入"]
        A[GeRootModel]
    end
    
    subgraph converter_block["ModelConverter"]
        B[ConvertGeModelToExecuteGraph]
        B1[解析 GeModel]
        B2[创建 ExecuteGraph]
        B3[设置三子图类型]
    end
    
    subgraph graphconv_block["GraphConverter"]
        C[ConvertComputeGraphToExecuteGraph]
        C1[CreateInitNode<br/>初始化节点]
        C2[CreateDeInitNode<br/>清理节点]
        C3[CreateMainNode<br/>内部调用 LoweringComputeGraph]
        C6[OfflineOptimizer<br/>优化]
        C7[TopologicalSorting<br/>拓扑排序]
        C8[CalcNodeExecutionPriorities<br/>计算节点优先级]
        C9[AppendGraphLevelData<br/>追加图级数据]
    end
    
    subgraph output_block["输出"]
        D[ExecuteGraph]
        D1[Init Graph]
        D2[Main Graph]
        D3[DeInit Graph]
    end
    
    A --> B --> B1 --> B2 --> B3 --> C
    C --> C1 --> C2 --> C3 --> C6 --> C7 --> C8 --> C9 --> D
    D --> D1 & D2 & D3

3. KernelTrace 定位机制

3.1 KernelTrace 概述

KernelTrace 是 RT2.0 的核心定位工具，通过 Subscriber 机制 在算子执行时打印维测信息，帮助开发者快速定位执行问题。

核心原理：

每个算子内核通过 REGISTER_KERNEL 注册 trace_printer 函数
ExecutorTracer 作为 Subscriber 监听执行事件
在 kExecuteEnd 事件触发时调用 trace_printer 打印内核信息

3.2 KernelTrace 架构

graph TB
    subgraph reg_block["Kernel 注册"]
        A[REGISTER_KERNEL<br/>注册 trace_printer]
    end
    
    subgraph registry_block["KernelRegistry"]
        B[FindKernelFuncs<br/>查找 trace_printer]
    end
    
    subgraph sub_block["Subscriber 系统"]
        C[ExecutorSubscribersScheduler]
        D[ExecutorTracer]
    end
    
    subgraph exec_block["执行流程"]
        E[SequentialExecuteWithCallback]
        F[kExecuteStart --> node.func --> kExecuteEnd]
    end
    
    subgraph trace_block["Trace 输出"]
        G[GetTracePrinter<br/>获取 trace_printer]
        H[OnExecuteEvent<br/>打印日志]
        I["[KernelTrace][kernel_name]..."]
    end
    
    A --> B
    C --> D
    E --> F --> C
    F --> D --> G --> B --> H --> I
    
    classDef regStyle fill:#e1bee7,stroke:#7b1fa2
    classDef registryStyle fill:#c8e6c9,stroke:#2e7d32
    classDef subStyle fill:#bbdefb,stroke:#1565c0
    classDef execStyle fill:#fff3e0,stroke:#f57c00
    classDef traceStyle fill:#ffcdd2,stroke:#c62828
    
    class A regStyle
    class B registryStyle
    class C,D subStyle
    class E,F execStyle
    class G,H,I traceStyle

3.3 关键代码位置

文件	函数/宏	作用
`runtime/v2/subscriber/tracer/executor_tracer.cc:40-79`	`OnExecuteEvent`	执行事件回调，触发 trace 打印
`runtime/v2/subscriber/tracer/executor_tracer.cc:81-89`	`GetTracePrinter`	从 KernelRegistry 获取 trace_printer
`runtime/v2/core/debug/kernel_tracing.h:16-37`	`KERNEL_TRACE_*` 宏	KernelTrace 日志打印宏
`inc/graph_metadef/register/kernel_registry.h`	`TracePrinter` 类型	trace_printer 函数签名定义

3.4 KernelTrace 输出格式

[KernelTrace][kernel_name] <具体维测信息>

示例输出：

[KernelTrace][InferShape] infer shape result: [2, 1024]
[KernelTrace][AllocMemory] [MEM]Alloc memory at stream 0, block 0x7f..., address 0x7f..., size 8192
[KernelTrace][LaunchKernelWithHandle] kernel launched, task_id=123, stream_id=0
[KernelTrace][FreeMemory] [MEM]Free memory at stream 0, address 0x7f...

3.5 如何启用 KernelTrace

KernelTrace 的启用依赖于 GE 模块日志级别。当日志级别设置为 INFO 或更低时，KernelTrace 自动生效。

方式一：设置日志级别（环境变量）

# 将 GE 模块日志级别设为 INFO（具体变量名取决于日志框架配置）
export ASCEND_GLOBAL_LOG_LEVEL=1

方式二：原理说明

// GlobalTracer::GetEnableFlags() 的实现：
// return static_cast<uint64_t>(IsLogEnable(GE_MODULE_NAME, DLOG_INFO));
// 当 GE 模块 INFO 级别日志使能时，KernelTrace 即启用

注意：GlobalTracer 没有 SetEnableFlags 方法，其启用状态完全由日志级别决定。

3.6 内核 Trace 信息类型

Trace 类型	宏定义	输出内容
内存分配	`KERNEL_TRACE_ALLOC_MEM`	流 ID、block 地址、device 地址、大小
内存释放	`KERNEL_TRACE_FREE_MEM`	流 ID、device 地址
内存迁移	`KERNEL_TRACE_WANDERING`	源流、目标流、device 地址
自定义	Kernel 注册的 `trace_printer`	内核特定的维测信息

3.7 自定义 TracePrinter 示例

注意：TracePrinter 类型为裸函数指针 std::vector<std::string> (*)(const KernelContext *)，不支持有捕获的 lambda。

// 注册内核时添加 trace_printer（必须是无捕获 lambda 或普通函数）
REGISTER_KERNEL(MyKernel)
    .RunFunc(MyKernelRun)
    .TracePrinter([](const KernelContext *context) -> std::vector<std::string> {
        std::vector<std::string> msgs;
        msgs.push_back("input shape: " + context->GetInputShape(0).ToString());
        msgs.push_back("output shape: " + context->GetOutputShape(0).ToString());
        return msgs;
    });

4. 外部接口说明

4.1 接口总览

RT2.0 有三层对外接口：ACL 公开 API → gert 内部 API → 执行器 API。

graph TB
    subgraph acl_if["ACL 公开接口 (C API)"]
        AL1[aclmdlLoadFromFile<br/>aclmdlLoadFromMem]
        AL2[aclmdlLoadFromFileWithMem<br/>aclmdlLoadFromMemWithMem]
        AE1[aclmdlExecuteV2<br/>aclmdlExecuteAsyncV2]
        AU1[aclmdlUnload]
        AG1[aclmdlGetDesc<br/>aclmdlGetDescFromFile]
    end

    subgraph om2_if["OM2 gert 接口"]
        OL1[LoadOm2ExecutorFromData]
        OL2[LoadOm2DataFromFile]
        OE1[Om2ModelExecutor::Load]
        OE2[Om2ModelExecutor::Run<br/>Om2ModelExecutor::RunAsync]
    end

    subgraph gert_if["RT2 gert 接口"]
        GL1[LoadExecutorFromFile]
        GL2[LoadExecutorFromModelData x5]
        GL3[LoadStreamExecutorFromModelData x3]
        GU1[IsDynamicModel / LoadDataFromFile]
    end

    subgraph exec_if["执行器接口"]
        VL1[ModelV2Executor::Load]
        VL2[ModelV2Executor::Execute<br/>ModelV2Executor::ExecuteSync]
        VL3[ModelV2Executor::UnLoad]
        VL4[GetModelDesc / GetIterationNum / GetSubscribers]
        SE1[StreamExecutor::GetOrCreateLoaded]
    end

    subgraph aux_if["辅助接口"]
        AX1[AllocatorFactory::Create]
        AX2[CreateExternalAllocator]
    end

    AL1 --> OL1
    AE1 --> OE2
    GL1 --> VL1
    OL1 --> GL1

    classDef acl fill:#c8e6c9,stroke:#2e7d32
    classDef om2 fill:#e3f2fd,stroke:#1565c0
    classDef gert fill:#fff3e0,stroke:#f57c00
    classDef exec fill:#f3e5f5,stroke:#7b1fa2
    classDef aux fill:#eceff1,stroke:#37474f

    class acl_if,AL1,AL2,AE1,AU1,AG1 acl
    class om2_if,OL1,OL2,OE1,OE2 om2
    class gert_if,GL1,GL2,GL3,GU1 gert
    class exec_if,VL1,VL2,VL3,VL4,SE1 exec
    class aux_if,AX1,AX2 aux

4.2 入口场景与接口对应

场景	加载接口	执行接口	适用方
ACL 在线推理	`aclmdlLoadFromFile` / `aclmdlLoadFromMem`	`aclmdlExecuteV2` / `aclmdlExecuteAsyncV2`	用户直接调用 ACL C API
GeSession 在线推理	`GeSession::RunGraph`（首次自动加载）	`GeSession::RunGraph`	GE Session 用户
Dflow / 内部直调	`gert::LoadExecutorFromModelData`	`ModelV2Executor::Execute`	Dflow、单算子等内部组件

4.3 ACL 外部接口与 RT2.0 路由

ACL 的 Load/Execute API 是统一入口，内部判定后路由：

ACL C API	路由逻辑	说明
`aclmdlLoadFromFile`	OM2 → `ImplOm2`；非OM2 → `IsDynamicModel` → `RuntimeV2ModelLoadCommon`	两种格式均可走 RT2.0
`aclmdlLoadFromMem`	同上
`aclmdlExecuteV2`	OM2 → `ImplOm2`；非OM2 → `Impl`	执行
`aclmdlUnload`	OM2 → `ImplOm2`；非OM2 → `Impl`	卸载

进入 RT2.0 的条件（非 OM2 路径）：

IsSupportRuntimeV2WithModelPath
  1. IsRuntimeV2Enable() → RuntimeV2 特性是否使能
  2. gert::IsDynamicModel() → 模型是否是动态 Shape
  两者都满足 → RuntimeV2ModelLoadCommon → gert::LoadExecutorFromModelData → RT2.0

结论：OM2 格式一定走 RT2.0；老 OM 格式中动态 Shape 模型也会走 RT2.0（通过 RuntimeV2ModelLoadCommon），静态 Shape 模型走 v1 TaskSink。

4.4 gert 加载接口

接口	参数	返回值	说明
`LoadExecutorFromFile`	`model_path, error_code`	`unique_ptr<ModelV2Executor>`	从 OM 文件加载执行器
`LoadExecutorFromModelData`	`model_data, error_code`	`unique_ptr<ModelV2Executor>`	从内存中的 ModelData 加载
`LoadExecutorFromModelData`	`model_data, ExecutorOption, error_code`	`unique_ptr<ModelV2Executor>`	带执行器选项加载
`LoadExecutorFromModelData`	`model_data, LoadExecutorArgs, error_code`	`unique_ptr<ModelV2Executor>`	带 RtSession + FileConstant 加载
`LoadExecutorFromModelDataWithRtSession`	`model_data, rt_session, error_code`	`unique_ptr<ModelV2Executor>`	绑定 RtSession 加载
`LoadStreamExecutorFromModelData`	`model_data, error_code`	`unique_ptr<StreamExecutor>`	多流场景加载

4.5 执行接口详解

ModelV2Executor 生命周期接口：

// 加载模型（执行 Init Graph，准备 Main Graph）
ge::graphStatus Load();
ge::graphStatus Load(const ModelExecuteArg &arg);
ge::graphStatus Load(const ModelExecuteArg &arg, const ModelLoadArg &load_arg);

// 异步执行模型
ge::graphStatus Execute(const ModelExecuteArg &arg, 
                        Tensor **inputs, size_t input_num,
                        Tensor **outputs, size_t output_num);

// 同步执行（内部创建 default stream 并自动同步）
ge::graphStatus ExecuteSync(Tensor **inputs, size_t input_num,
                            Tensor **outputs, size_t output_num);

// 卸载模型（执行 DeInit Graph）
ge::graphStatus UnLoad();

ModelExecuteArg 参数说明：

参数	类型	说明
`stream`	`rtStream_t`	执行流（可为空，使用默认流）
`external_allocator`	`Allocators *`	外部内存分配器
`external_stream_allocator`	`StreamAllocator *`	外部辅流分配器
`external_event_allocator`	`EventAllocator *`	外部 Event 分配器
`external_notify_allocator`	`NotifyAllocator *`	外部 Notify 分配器

ModelLoadArg 参数说明：

参数	类型	说明
`rt_session`	`RtSession *`	运行时 Session（变量管理、资源隔离）
`outer_weight_mem`	`OuterWeightMem`	外部权重内存

4.6 接口数量汇总

类别	数量
ACL 公开 C API	7 个 (Load/Execute/Unload)
OM2 gert 接口	4 个
gert:: RT2 加载接口	8 个
ModelV2Executor 接口	10+ 个
StreamExecutor 接口	2 个
辅助接口	2 个

5. 关键定位技巧

5.1 日志关键点

日志 Pattern	文件位置	定位意义
`Failed to load root model`	`api.cc:36-38`	模型加载失败入口
`Failed to lowering to execute graph`	`api.cc:65`	Lowering 转换失败
`ReadInCompileResults`	`model_converter.cc:192,198`	读取编译结果（日志：Read-in static/dynamic compiled graph）
`Root graph total stream_num`	`model_converter.cc:367`	流分配信息
`Failed to load init graph`	`model_v2_executor.cc:176`	Init Graph 加载失败
`Failed to execute init graph`	`model_v2_executor.cc:190`	Init 执行失败
`Failed to load main graph`	`model_v2_executor.cc:194`	Main Graph 加载失败
`Failed to execute model`	`model_v2_executor.cc:229`	Execute 前状态检查（未 Load 就调用 Execute）
`[KernelTrace]...`	`executor_tracer.cc:47,51,64`	Kernel 执行维测

5.2 常见问题定位流程

flowchart TD
    A[问题发生] --> B{问题类型?}
    
    B -->|加载失败| C[检查 api.cc LoadToModelV2ExecutorBuilder]
    C --> C1[检查 ModelHelper::LoadRootModel]
    C1 --> C2[检查 ModelConverter::ConvertGeModelToExecuteGraph]
    
    B -->|Lowering 失败| D[检查 graph_converter.cc]
    D --> D1[检查 ConvertComputeGraphToExecuteGraph]
    D1 --> D2[检查 LoweringComputeGraph]
    D2 --> D3[检查 NodeConverterRegistry::FindRegisterData]
    
    B -->|执行失败| E[检查 model_v2_executor.cc]
    E --> E1[检查 Load 流程]
    E1 --> E2[检查 SpecifyInputs/SpecifyOutputs]
    E2 --> E3[启用 KernelTrace 查看内核执行]
    
    B -->|内存问题| F[检查 KERNEL_TRACE_ALLOC_MEM/FREE_MEM]
    F --> F1[检查 Allocator 状态]
    F1 --> F2[检查 io_same_addr_pairs 校验]
    
    style A fill:#ffcdd2,stroke:#c62828
    style B fill:#fff3e0,stroke:#f57c00
    style C fill:#bbdefb,stroke:#1565c0
    style D fill:#c8e6c9,stroke:#2e7d32
    style E fill:#e1bee7,stroke:#7b1fa2
    style F fill:#e0f2f1,stroke:#00695c

5.3 KernelTrace 定位案例

案例：算子执行失败定位

步骤 1: 启用 KernelTrace
export GERT_KERNEL_TRACE_ENABLE=1

步骤 2: 执行模型，观察日志
[KernelTrace][InferShape] infer shape result: [2, 1024]  ← Shape 推导正常
[KernelTrace][Tiling] tiling key: 12345, workspace: 4096  ← Tiling 正常
[KernelTrace][AllocMemory] [MEM]Alloc memory... size 4096  ← 内存分配正常
[KernelTrace][LaunchKernelWithHandle] kernel launched, task_id=123  ← Kernel 启动
<无后续日志>  ← Launch 后无输出，说明 Launch 失败

步骤 3: 定位 LaunchKernelWithHandle 内核
检查 runtime/v2/engine/aicore/kernel/launch_kernel_with_handle.cc

5.4 Subscriber 系统

RT2.0 提供可扩展的事件订阅机制，可在执行过程中插入回调：

Subscriber	作用	文件位置
`ExecutorTracer`	KernelTrace 定位	`runtime/v2/subscriber/tracer/`
`CannProfilerV2`	性能分析	`runtime/v2/subscriber/profiler/`
`ExecutorDumper`	数据 Dump	`runtime/v2/subscriber/dumper/`

6. 关键文件索引

文件路径	核心函数	行号
`runtime/v2/api/api.cc`	`LoadToModelV2ExecutorBuilder`	31-70
`runtime/v2/lowering/model_converter.cc`	`ConvertGeModelToExecuteGraph`	528-581
`runtime/v2/lowering/graph_converter.cc`	`ConvertComputeGraphToExecuteGraph`	854-910
`runtime/v2/lowering/graph_converter.cc`	`LoweringComputeGraph`	778-811
`runtime/v2/core/model_v2_executor.cc`	`Load`	167-197
`runtime/v2/core/model_v2_executor.cc`	`Execute`	226-250
`runtime/v2/core/model_v2_executor.cc`	`UnLoad`	199-224
`runtime/v2/core/executor/sequential/executor/sequential_executor.c`	`SequentialExecute`	17-29
`runtime/v2/subscriber/tracer/executor_tracer.cc`	`OnExecuteEvent`	40-79
`runtime/v2/subscriber/tracer/executor_tracer.cc`	`GetTracePrinter`	81-89
`inc/framework/runtime/gert_api.h`	公共 API 定义	全文
`inc/framework/runtime/model_v2_executor.h`	ModelV2Executor 定义	全文

7. 设计约束与注意事项

7.1 RT2 运行时约束（来自 `constraints/rt2_runtime.md`）

加载时约束：

加载和执行必须使用相同的 stream 和 allocator
加载完成后需调用流同步（若 stream/allocator 不同）
RT2 构造过程中不应修改计算图
Lowering 生成的节点顺序 ≠ 执行顺序

执行时约束：

异步 H2D 拷贝必须配合 HOST_TO_DEVICE_EX 选项
涉及资源处理需考虑资源规格和生命周期

性能约束：

执行时避免动态申请内存
新增/修改 kernel 时需评估性能影响（劣化不超过 100ns）

7.2 内存 Allocator 约束

一个 allocator 仅对应唯一的 stream
流同步前，allocator 内存池内存不可归还操作系统
流同步前，allocator 不可被析构

8. 参考资料

文档	路径	内容
RT2.0 特性分析	`docs/architecture/features/unknown_shape_executor.md`	Lowering、三子图、Kernel 注册
RT2 运行时约束	`docs/architecture/constraints/rt2_runtime.md`	设计原则、约束
运行时架构	`docs/architecture/modules/runtime/runtime.md`	v1/v2 对比、执行流程
Profiling 特性	`docs/architecture/features/profiling.md`	性能分析机制
Dump 特性	`docs/architecture/features/datadump.md`	数据 Dump 机制

9. 附录：三子图节点类型

Init Graph 典型节点

节点类型	作用
`Data`	输入占位
`CreateL1Allocator`	创建 L1 分配器
`CreateL2Allocators`	创建 L2 分配器
`SelectL1Allocator`	选择 L1 分配器
`AllocMemory`	分配权重内存
`CopyH2D`	权重下沉到 Device
`SplitRtStreams`	流资源准备
`CreateGertEvents`	事件创建
`CreateNotifies`	通知创建
`InnerNetOutput`	输出

Main Graph 典型节点（一个算子展开为 6-10 个节点）

序号	节点类型	分类	必须	作用
①	`InferShape`	Shape	✅	推导输出 Shape/dtype
②	`Tiling`	Tiling	✅	计算分块策略和 workspace
③	`AllocMemory`	内存	✅	分配输出 tensor/workspace
④	`LaunchKernelWithHandle`	执行	✅	调用硬件执行算子
⑤	`FreeMemory`	内存	✅	释放临时内存
⑥	`SendEvents/WaitEvents`	同步	⚪	多流场景跨流同步
⑦	`CopyD2H/CopyH2D`	搬运	⚪	Host/Device 数据搬运
⑧	`BuildTensor`	Shape	⚪	复杂 Tensor 构造

DeInit Graph 典型节点

节点类型	作用
`Data`	输入占位
`FreeMemory` / `FreeMemHbm`	释放权重/输出内存
`FreeFixedFeatureMemory`	释放固定特征内存
`DestroyEvent`	销毁事件
`DavinciModelFinalizer`	静态子图资源清理（内部销毁流）
`InnerNetOutput`	输出

注意：不存在独立的 DestroyStream 节点类型。流的销毁由 DavinciModelFinalizer 内核内部完成，或由 StreamAllocator 生命周期管理。

10. 实战示例：AddExample 算子入图全流程

以最简单的 AddExample（逐元素加法）为例，展示一个算子需要交付哪些文件，以及这些文件在 GE RT2.0 框架中是如何被调用的。

10.1 算子交付件总览

add_example/
├── op_host/
│   ├── add_example_def.cpp          ← ① 算子定义（OpDef）
│   ├── add_example_infershape.cpp   ← ② InferShape 实现
│   └── add_example_tiling.cpp       ← ③ Tiling 实现
├── op_kernel/
│   ├── add_example_tiling_data.h    ← ④ Tiling 数据结构（Host/Device 共享）
│   ├── add_example_tiling_key.h     ← ⑤ Tiling Key 定义
│   ├── add_example.h               ← ⑥ Kernel 实现（Ascend C）
│   └── add_example.cpp             ← ⑦ Kernel 入口函数
└── (编译产物: kernel binary .o)     ← ⑧ 编译后的算子二进制

10.2 各交付件代码示例

① 算子定义 op_host/add_example_def.cpp

声明算子的输入/输出/属性/支持的硬件平台：

#include "register/op_def_registry.h"

namespace ops {
class AddExample : public OpDef {
public:
    explicit AddExample(const char* name) : OpDef(name) {
        this->Input("x1")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT, ge::DT_INT32})
            .Format({ge::FORMAT_ND, ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
        this->Input("x2")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT, ge::DT_INT32})
            .Format({ge::FORMAT_ND, ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
        this->Output("y")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT, ge::DT_INT32})
            .Format({ge::FORMAT_ND, ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});

        OpAICoreConfig aicoreConfig;
        aicoreConfig.DynamicShapeSupportFlag(true);
        this->AICore().AddConfig("ascend910b", aicoreConfig);
    }
};
OP_ADD(AddExample);
} // namespace ops

② InferShape 实现 op_host/add_example_infershape.cpp

运行时推导输出 tensor 的 shape：

#include "register/op_impl_registry.h"

namespace ops {
static ge::graphStatus InferShapeAddExample(gert::InferShapeContext* context) {
    const gert::Shape* xShape = context->GetInputShape(0);
    gert::Shape* yShape = context->GetOutputShape(0);
    // 输出 shape = 输入 shape（逐元素操作）
    yShape->SetDimNum(xShape->GetDimNum());
    for (size_t i = 0; i < xShape->GetDimNum(); i++) {
        yShape->SetDim(i, xShape->GetDim(i));
    }
    return ge::GRAPH_SUCCESS;
}
// 注册宏：将 InferShape 函数绑定到算子类型名
IMPL_OP_INFERSHAPE(AddExample).InferShape(InferShapeAddExample);
} // namespace ops

③ Tiling 实现 op_host/add_example_tiling.cpp

运行时根据实际 shape 计算分块策略：

#include "../op_kernel/add_example_tiling_data.h"

namespace optiling {
static ge::graphStatus AddExampleTilingFunc(gert::TilingContext* context) {
    // 获取实际输入 shape，计算总元素数
    auto inputX = context->GetInputShape(0);
    int64_t totalIdx = inputX->GetStorageShape().GetShapeSize();

    // 填充 TilingData（会传给 Device 侧 kernel）
    AddExampleTilingData* tiling = context->GetTilingData<AddExampleTilingData>();
    tiling->totalLength = totalIdx;
    tiling->tileNum = 8;

    context->SetBlockDim(8);       // 使用 8 个 AI Core
    context->SetTilingKey(0);      // 选择 float 分支
    return ge::GRAPH_SUCCESS;
}
// 注册宏：将 Tiling 函数绑定到算子类型名
IMPL_OP_OPTILING(AddExample).Tiling(AddExampleTilingFunc);
} // namespace optiling

④ Tiling 数据结构 op_kernel/add_example_tiling_data.h（Host/Device 共享）

struct AddExampleTilingData {
    int64_t totalLength;  // 总元素数
    int64_t tileNum;      // 分块数
};

⑦ Kernel 入口 op_kernel/add_example.cpp（Ascend C，运行在 NPU 上）

#include "add_example.h"

template <uint32_t schMode>
__global__ __aicore__ void add_example(GM_ADDR x, GM_ADDR y, GM_ADDR z,
                                        GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA_WITH_STRUCT(AddExampleTilingData, tilingData, tiling);
    NsAddExample::AddExample<float> op;
    op.Init(x, y, z, &tilingData);
    op.Process();  // CopyIn → Add → CopyOut 循环
}

10.3 GE 框架如何调用这些交付件

sequenceDiagram
    participant User as 用户代码
    participant ACL as ACL API
    participant Lowering as Lowering 阶段
    participant Init as Init Graph
    participant Main as Main Graph (执行循环)
    participant NPU as NPU Device

    Note over User,NPU: ═══ 编译期（离线） ═══
    User->>ACL: ATC 编译模型
    Note right of ACL: ① OpDef 被读取<br/>确定输入输出规格<br/>选择 AI Core 引擎

    Note over User,NPU: ═══ 加载期 ═══
    User->>ACL: aclmdlLoadFromFile()
    ACL->>Lowering: LoadExecutorFromModelData
    Note right of Lowering: 根据引擎类型查找<br/>NodeConverter<br/>(REGISTER_NODE_CONVERTER)
    Lowering->>Lowering: LoweringAiCoreNode()
    Note right of Lowering: 为该算子生成：<br/>InferShape 节点<br/>Tiling 节点<br/>AllocMemory 节点<br/>LaunchKernelWithHandle 节点<br/>FreeMemory 节点
    Lowering->>Init: 生成 Init Graph
    Init->>NPU: 分配内存/流/事件

    Note over User,NPU: ═══ 执行期（每次推理） ═══
    User->>ACL: aclmdlExecuteV2()
    ACL->>Main: ModelV2Executor::Execute()

    Main->>Main: InferShape 节点执行
    Note right of Main: 调用 ② InferShapeAddExample()<br/>推导输出 shape = [N,C,H,W]

    Main->>Main: Tiling 节点执行
    Note right of Main: 调用 ③ AddExampleTilingFunc()<br/>计算 blockDim=8, tileNum=8<br/>填充 ④ AddExampleTilingData

    Main->>Main: AllocMemory 节点执行
    Note right of Main: 根据 InferShape 结果<br/>分配输出 tensor 内存

    Main->>NPU: LaunchKernelWithHandle 节点执行
    Note right of NPU: 调用 ⑦ add_example kernel<br/>传入 TilingData<br/>NPU 上执行逐元素加法

    Main->>Main: FreeMemory 节点执行
    Note right of Main: 释放 workspace 临时内存

    NPU-->>User: 结果写入输出 tensor

10.4 交付件 → GE 注册机制 → 运行时调用对照表

交付件	注册宏	GE 内部存储位置	运行时调用入口
OpDef (`_def.cpp`)	`OP_ADD(AddExample)`	OpDef Registry	编译期确定引擎/格式/dtype
InferShape (`_infershape.cpp`)	`IMPL_OP_INFERSHAPE(AddExample)`	OpImplRegistry	`InferShape` 节点 → `FindInferShapeFunc` → 调用注册函数
Tiling (`_tiling.cpp`)	`IMPL_OP_OPTILING(AddExample)`	OpTilingRegistry	`Tiling` 节点 → `FindTilingFunc` → 调用注册函数
TilingData (`_tiling_data.h`)	无（纯数据结构）	—	Host 侧 Tiling 填充，Device 侧 Kernel 读取
Kernel binary (`.o`)	编译器自动生成	TaskDef 中的 binary	`LaunchKernelWithHandle` 节点 → 下发到 NPU

10.5 Lowering 展开细节

一个 AddExample 算子节点在 Lowering 后展开为以下 ExecuteGraph 节点序列：

ComputeGraph 中的一个节点:  AddExample(x1, x2) → y
                                    │
                                    ▼ LoweringAiCoreNode()
ExecuteGraph 中展开为 ~7 个节点:
┌─────────────────────────────────────────────────────────┐
│ FindInferShapeFunc  → 查找 InferShapeAddExample 函数指针  │
│ InferShape          → 调用它，得到输出 shape              │
│ FindTilingFunc      → 查找 AddExampleTilingFunc 函数指针  │
│ Tiling              → 调用它，得到 TilingData + blockDim  │
│ AllocMemory         → 按 shape 分配输出 tensor 内存       │
│ LaunchKernelWithHandle → 下发 kernel binary 到 NPU 执行   │
│ FreeMemory          → 释放 workspace                     │
└─────────────────────────────────────────────────────────┘

对应代码路径：

REGISTER_NODE_CONVERTER_PLACEMENT(ge::kEngineNameAiCore, ...) → LoweringAiCoreNode()（aicore_node_converter.cc:1217）
内部调用 bg::InferShape()、bg::Tiling()、bg::AllocMem()、bg::LaunchKernelWithHandle() 等 graph_builder 接口生成节点

10.6 关键结论

算子开发者只需交付 4 类文件：OpDef、InferShape、Tiling、Kernel — 不需要关心 GE 执行图的构建
GE 框架通过注册宏自动发现：OP_ADD / IMPL_OP_INFERSHAPE / IMPL_OP_OPTILING 将函数注册到全局 Registry
Lowering 阶段自动展开：NodeConverterRegistry 根据算子引擎类型（AICore/AICpu/...）找到对应的 Converter，将一个计算节点展开为多个执行节点
执行期按节点顺序调用：SequentialExecuteWithCallback 逐个执行节点的 func 指针，每个节点内部通过 KernelRegistry 找到对应的 RunFunc

文档版本: v1.0
生成日期: 2026-05-20
参考源码: GE runtime/v2/ 目录