Star596
204
代码
代码
Issues11
Pull Requests41
Actions
项目讨论
Wiki
分析
Star596
204

CV 数据通路（cv_datapath）

概述

本样例以 MatMul + ReLU 融合算子计算流程，展示 Ascend 950PR/950DT 产品在 CV（Cube-Vector）数据通路 方面的新硬件特性，介绍：

  1. Mix 核（__mix__(1,2)）下 AIC / AIV 协同编程模型；
  2. Scenario1：L0C → GM → UB → ReLU → GM（GM 中转融合）；
  3. Scenario2：L0C → UB → ReLU → GM（Ascend 950 新增 Fixpipe / L0C→UB 直通）。

支持产品与 CANN 版本

产品 CANN 软件版本
Ascend 950PR / Ascend 950DT >= CANN 9.1.0

NPU 架构：dav-3510（以下使用28核Ascend950PR平台进行展示说明）。

算子与规格

C=ReLU(A×B),ReLU(x)=max⁡(x,0)C = \mathrm{ReLU}(A \times B),\quad \mathrm{ReLU}(x)=\max(x,0)

输入shape：

  • [M, K, N] = [8192, 8192, 8192]
  • [dtypeA, dtypeB, dtypeC] = [float16, float16, float32]
  • A, B, C 数据格式均为 ND（生成数据时，会将B矩阵进行转置以实现高性能计算）

目录结构

cv_datapath/
├── CMakeLists.txt
├── README.md
├── include/
│   ├── common_utils.h
│   ├── host_runner.h
│   └── kernel_common.h          # Cube 流水与 Fixpipe 封装
├── src/
│   ├── 0_separate.asc           # Cube / Vector 分离 launch
│   ├── 1_mix_scenario1.asc      # Mix + Scenario1（L0C→GM→UB）
│   └── 2_mix_scenario2.asc      # Mix + Scenario2（L0C→UB）
└── scripts/
    ├── gen_data.py
    └── verify_result.py

样例递进

Scenario 可执行文件 说明
0 cv_datapath_0_separate 先 launch MatMul（L0C→GM），再 launch ReLU（GM→UB→GM），无融合基线
1 cv_datapath_1_mix_scenario1 单次 Mix launch；AIC Fixpipe→GM，AIV 读 GM 做 ReLU
2 cv_datapath_2_mix_scenario2 单次 Mix launch；AIC Fixpipe→UB，AIV 直接 ReLU 后写 GM

建议阅读顺序：Scenario0（理解 CV 分离方案）→ Scenario1（理解 Mix 融合方案 + GM 中转）→ Scenario2（理解 L0C→UB 收益）。

AIC↔AIV 的 CV 融合机制

在昇腾 AI Core 上，Cube（矩阵乘） 与 Vector（逐元素算子） 由不同硬件单元执行：

简称 全称 本样例职责
AIC AI Cube Core 执行 MatMul：GM→L1→L0→L0C，再经 Fixpipe 搬出
AIV AI Vector Core 执行 ReLU：在 UB 上做 max(x,0)，再写回 GM

若写成两个独立 Kernel 先后 launch（Scenario0），在实际执行时，将会在 AIC 上的MatMul计算全部结束后，才开始 AIV 上的 ReLU 计算，二者无法重叠。

为了进一步提升性能，CANN提供了融合算子的开发范式。融合算子通常由多个独立的小算子融合而成，其功能与多个小算子的功能等价，而性能方面通常优于独立的小算子。融合了Cube计算、Vector计算的算子统称为CV融合算子。

下面分三点说明本样例采用的机制（对应 Scenario1 / Scenario2）。

1. CV 融合架构：__mix__(1, 2) 核函数

使用 __global__ __mix__(1, 2) 声明核函数，含义是：

  • 该kernel声明为mix融合算子，其中 AIC 与 AIV 的配比数目为1:2；
  • 同一份 Kernel 源码会分别在 AIC 与 AIV 上执行，用宏区分路径。

以 Scenario1 为例（Scenario2 结构相同，仅搬出通路不同）：

__global__ __mix__(1, 2) void matmul_relu_scenario1(
    __gm__ uint8_t* a, __gm__ uint8_t* b, __gm__ uint8_t* c)
{
    AscendC::InitSocState();
    KernelMatmulReluScenario1 op;
    op.Init(a, b, c);
    op.Process();
    AscendC::PipeBarrier<PIPE_ALL>();
}

在 Process() 中，通过 ASCEND_IS_AIC / ASCEND_IS_AIV 拆分职责：

__aicore__ inline void Process()
{
    if ASCEND_IS_AIC {
        // AIC：MatMul 累加到 L0C → Fixpipe 搬出 → 通知 AIV
        AscendC::LocalTensor<float> cLocal(...);
        pipe.RunMatmul(cLocal);
        pipe.FixpipeToGm(cLocal);   // Scenario1；Scenario2 则为 FixpipeToUbDualM
        AscendC::CrossCoreSetFlag<0x2, PIPE_FIX>(AIC_SYNC_AIV_FLAG);
    }
    if ASCEND_IS_AIV {
        // AIV：等待就绪 → ReLU → 写回 GM
        ReluFromGm();               // Scenario1；Scenario2 则为 ReluFromUb
    }
}

代码说明：

  1. __mix__(1,2) = 「 AIC 与 AIV 的配比数目为 1:2 」；
  2. ASCEND_IS_AIC / ASCEND_IS_AIV 决定当前硬件上实际走哪段代码；
  3. 两边共享同一套 GM 指针与问题规模，但本地 Buffer（L0C / UB）彼此独立。

2. 跨核同步：AIC→AIV 数据就绪通知

AIC 把 MatMul 结果经 Fixpipe 写出之后，AIV 才能安全读取。本样例用 CrossCore Flag 相关的同步API 实现跨核同步：

// AIC 侧：Fixpipe 完成后置位（PIPE_FIX 表示与 Fixpipe 流水绑定）
AscendC::CrossCoreSetFlag<0x2, PIPE_FIX>(AIC_SYNC_AIV_FLAG);  // flag = 0x8

// AIV 侧：进入 ReLU 前等待
AscendC::CrossCoreWaitFlag(AIC_SYNC_AIV_FLAG);

时序可以理解为：

AIC:  CopyIn → Mmad → Fixpipe ──► SetFlag(0x8)
                                      │
AIV:  WaitFlag(0x8) ──► (读 GM 或直接用 UB) → Relu → 写 GM

注意：

  • Set / Wait 必须成对出现；若某条路径提前 return，仍需保证对侧不会永久等待（多核/无任务场景尤其重要）；
  • Flag 编号在工程中集中定义为 AIC_SYNC_AIV_FLAG（本样例为 0x8）；
  • Scenario0 的两次独立 launch 之间由 Host aclrtSynchronizeStream 保证顺序，不需要 CrossCore；Scenario1/2 则在核内完成同步。

3. __mix__(1, 2) 模式下的多核切分与双 AIV 分工

3.1 逻辑核 ID（多核时）

__mix__(1,2) 下，GetBlockIdx() 在 AIC 与 AIV 侧各自独立编号：

  • AIC：0 ~ numAICores-1
  • AIV：0 ~ numAIVCores-1（约为 AIC 数量的 2 倍）

因此，当将来扩展到多逻辑核时，AIV 侧通常要把 BlockIdx 映射回逻辑核 ID，才能与 AIC 使用同一套 M/N 切分：

uint32_t logicCoreId;
if ASCEND_IS_AIC {
    logicCoreId = AscendC::GetBlockIdx();
} else {
    logicCoreId = AscendC::GetBlockIdx() / 2;  // 每逻辑核对应 2 个 AIV
}
// 再由 logicCoreId 计算 mIterIdx / nIterIdx ...

本样例 numBlocks=28，MN 共 8×8=64 个 singleCore 块 / 16×16=256 个 base tile，按
blockId = logicCoreId; blockId < 64; blockId += numBlocks 轮询分配；AIV 侧用 GetBlockIdx()/2 映射回逻辑核 ID。

3.2 双 AIV 按 M 维拆分（本样例已使用）

即便只有一个逻辑核，__mix__(1,2) 仍会启动 2 个 AIV。本样例把 MatMul 输出的 M 行均分：

  • SubBlock0：处理行 [0, M/2)
  • SubBlock1：处理行 [M/2, M)

AIV 侧通过 GetSubBlockIdx() % 2 取得自己的子块索引：

uint32_t localSubIdx = AscendC::GetSubBlockIdx() % 2;
uint32_t gmOffset = localSubIdx * (MATMUL_RELU_M / 2) * MATMUL_RELU_N;

Scenario2 中，AIC 的 Fixpipe 还会设置 dualDstCtl = 0b01，把 L0C 结果按 M 维直接拆进两个 AIV 各自的 UB，与上述分工一致。

Scenario1：CV 融合 — GM 中转

对应代码：src/1_mix_scenario1.asc
定位：Mix 核入门通路；中间结果落 GM，便于调试与理解「先 Cube 后 Vector」。

数据流

GM ──(MTE2)──► L1 ──(MTE1)──► L0A/L0B ──(Cube)──► L0C
                                                    │
                                               Fixpipe (CO1→GM)
                                                    ▼
                                                   GM
                                                    │
                                               MTE2 (GM→UB)
                                                    ▼
                                                   UB ──VEC(ReLU)──► UB ──(MTE3)──► GM

与 Scenario0 的本质区别：Scenario0 是两次 Kernel launch；Scenario1 是一次 Mix launch，AIC/AIV 通过 CrossCore 衔接，Vector 侧仍从 GM 取数。

AIC 侧：Fixpipe L0C → GM

Ascend 950 上使用 FixpipeParamsArch3510，目的端为 GM（CFG_ROW_MAJOR_GM，即 isToUB=false）：

AscendC::FixpipeParamsArch3510<AscendC::CO2Layout::ROW_MAJOR> fixpipeParams;
fixpipeParams.mSize = baseM;
fixpipeParams.nSize = baseN;
fixpipeParams.srcStride = baseM;           // L0C 源侧 M 维步长
fixpipeParams.dstStride = MATMUL_RELU_N;   // GM 为 [M,N]，行宽是 N 而非 baseN
fixpipeParams.unitFlag = 3;               // 与 Cube unitFlag 收尾一致
uint32_t gmOffset = mTileIdx * baseM * MATMUL_RELU_N + nTileIdx * baseN;
AscendC::Fixpipe<float, float, CFG_ROW_MAJOR_GM>(cGM[gmOffset], cLocal, fixpipeParams);
AscendC::CrossCoreSetFlag<0x2, PIPE_FIX>(AIC_SYNC_AIV_FLAG);

AIV 侧：GM → UB → ReLU → GM

AscendC::CrossCoreWaitFlag(AIC_SYNC_AIV_FLAG);

uint32_t localSubIdx = AscendC::GetSubBlockIdx() % 2;
uint32_t gmOffset = localSubIdx * halfM * MATMUL_RELU_N;

// 1) GM → xUB
AscendC::DataCopyPad<float>(xUB, pipe.cGM[gmOffset], copyInParams, padParams);
// 2) ReLU
AscendC::Relu(yUB, xUB, halfMN);
// 3) yUB → GM
AscendC::DataCopyPad<float>(pipe.cGM[gmOffset], yUB, copyOutParams);

关键点：AIV 必须先完成 MTE2（GM→UB），才能做 Vector 计算，因此 profiling 中通常能看到明显的 aiv_mte2_time。

Scenario2：CV 融合 — UB 直通 / L0C→UB

对应代码：src/2_mix_scenario2.asc
定位：Ascend 950 新特性使能样例；省去 AIV 侧 GM→UB，是本 Story 的核心收益点。

为什么需要 L0C→UB？

Cube 累加结果默认在 L0C（CO1）。传统路径是 Fixpipe 到 GM，再由 AIV 用 MTE2 搬进 UB。
Ascend 950 的 Fixpipe 支持把 L0C 直通 UB（文档称通路 CO1→UB），于是：

  • 少一次「写 GM + 读 GM」；
  • AIV 的 aiv_mte2_time 可降至接近 0；
  • 仅 dav-3510 / Ascend 950 支持该通路（A2/A3 无此能力）。

官方接口说明见：Fixpipe API。

数据流

GM ──(MTE2)──► L1 ──(MTE1)──► L0A/L0B ──(Cube)──► L0C
                                                    │
                                               Fixpipe (CO1→UB, dualDstCtl=0b01)
                                                    ▼
                                                   UB ──VEC(ReLU)──► UB ──(MTE3)──► GM

与 Scenario1 对比，删除了「L0C→GM」以及「GM→UB」两段，其余 Cube 前半段与 ReLU→GM 后半段保持一致。

AIC 侧：Fixpipe L0C → UB（双目标按 M 拆分）

关键配置：

配置项 取值 含义
CFG_ROW_MAJOR_UB {ROW_MAJOR, true} NZ→ND，且目的为 UB（isToUB=true）
dualDstCtl 0b01 双目标模式，按 M 维拆到两个 SubBlock 的 UB
mSize DivCeilU32(baseM, 2) * 2 按 M 拆分时对齐到 2 的倍数
srcStride baseM L0C 源侧 M 维步长
dstStride baseN 每个 AIV 子块 UB 内 ND 行宽
unitFlag 3 与 Cube unitFlag 收尾一致 
AscendC::FixpipeParamsArch3510<AscendC::CO2Layout::ROW_MAJOR> fixpipeParams;
fixpipeParams.mSize = DivCeilU32(baseM, 2) * 2;  // 按 M 拆分时对齐到 2 的倍数
fixpipeParams.nSize = baseN;
fixpipeParams.srcStride = baseM;
fixpipeParams.dstStride = baseN;                 // 每个 AIV 子块 UB 内 ND 行宽
fixpipeParams.dualDstCtl = 0b01;                 // 按 M 拆到 2 个 AIV 的 UB
fixpipeParams.unitFlag = 3;
AscendC::Fixpipe<float, float, CFG_ROW_MAJOR_UB>(xUB, cLocal, fixpipeParams);
AscendC::CrossCoreSetFlag<0x2, PIPE_FIX>(AIC_SYNC_AIV_FLAG);

AIV 侧：UB → ReLU → GM（无 GM→UB）

AscendC::CrossCoreWaitFlag(AIC_SYNC_AIV_FLAG);

// xUB 已由 Fixpipe 写入本 SubBlock，无需 DataCopy GM→UB
AscendC::Relu(yUB, xUB, halfMN);

AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);

uint32_t localSubIdx = AscendC::GetSubBlockIdx() % 2;
uint32_t gmOffset = localSubIdx * halfM * MATMUL_RELU_N;
AscendC::DataCopyPad<float>(pipe.cGM[gmOffset], yUB, copyOutParams);

Scenario1 至 Scenario2 的实现代码差异可概括为两处：

  1. AIC：FixpipeToGm → FixpipeToUbDualM（CFG_ROW_MAJOR_UB + dualDstCtl）；
  2. AIV：删除 DataCopyPad GM→UB 段。

Scenario1 vs Scenario2 对照小结

对比项 Scenario1（GM 中转） Scenario2（UB 直通）
Fixpipe 目的 GM（CFG_ROW_MAJOR_GM） UB（CFG_ROW_MAJOR_UB）
AIV 取数 DataCopy GM→UB Fixpipe 已写入 UB
dualDstCtl 不需要 0b01（按 M 拆双 AIV）
典型 profiling 信号 aiv_mte2_time 可见 aiv_mte2_time ≈ 0
平台 Atlas A2 / Ascend 950PR 仅 Ascend 950PR

编译与运行

在仓库根目录先配置环境与工程（参考根目录 README.md）：

source ${install_path}/cann/set_env.sh
cmake -S . -B build -DNPU_ARCH=dav-3510

方式 A：只编译本样例（推荐）

cmake --install 默认会安装仓库内所有已配置目标；若只编译了本样例，安装会因缺少 vector_add 等其它二进制而失败。
单样例请构建后直接在 build 产物目录运行（gen_data.py / verify_result.py 已由 CMake COPYONLY 到该目录），或使用本样例的 install COMPONENT：

cmake --build build --target cv_datapath --parallel

# 选项 1：在 build 目录直接运行
cd build/Samples/1_Features/hardware_features/cv_datapath

# 选项 2：仅安装本样例到 build_out
# cmake --install build --prefix ./build_out --component cv_datapath
# cd build_out/1_Features/hardware_features/cv_datapath

方式 B：全量编译后再安装

与仓库根 README 一致：先编完全部 Target，再 cmake --install：

cmake --build build --parallel
cmake --install build --prefix ./build_out
cd build_out/1_Features/hardware_features/cv_datapath

运行与校验

在上述工作目录执行以下命令执行算子二进制文件，确认功能正常、算子精度测试通过：

python3 gen_data.py

./cv_datapath_0_separate
python3 verify_result.py output/output.bin output/golden.bin

./cv_datapath_1_mix_scenario1
python3 verify_result.py output/output.bin output/golden.bin

./cv_datapath_2_mix_scenario2
python3 verify_result.py output/output.bin output/golden.bin

成功时打印 test pass!。

性能对比与分析

采集方法

在样例运行目录（已 gen_data.py 且精度校验通过）执行：

# Scenario0 含两次 launch，需采集 matmul_only + relu_only
msprof op --launch-count=2 ./cv_datapath_0_separate
msprof op ./cv_datapath_1_mix_scenario1
msprof op ./cv_datapath_2_mix_scenario2

# 采集Scenario1和Scenario2的算子流水执行图（使用Mindstudio Insight打开获得的visualize_data.bin产物）
msprof op --aic-metrics=PipeTimeline ./cv_datapath_1_mix_scenario1 
msprof op --aic-metrics=PipeTimeline ./cv_datapath_2_mix_scenario2

关注 OpBasicInfo.csv 的 Task Duration，以及 PipeUtilization.csv 中各 AIC/AIV 的中位数（aic_cube_* / aiv_mte2_* 等）。

实测总览（Ascend 950PR，28 核）

Case Op Task Duration(μs) aic_cube_time(μs)¹ aic_cube_ratio¹ aiv_mte2_time(μs)¹ 说明
0 分离 matmul_only 3891.2 2542.0 0.980 ≈0（AIV 空闲） 仅 Cube→GM
0 分离 relu_only 302.0 0 0 161.9 独立 Vector launch
0 合计 两次 launch 4193.2 — — — 无 CV 重叠
1 Scenario1 matmul_relu_scenario1 3909.7 2542.0 0.975 57.4 GM 中转，融合覆盖 ReLU
2 Scenario2 matmul_relu_scenario2 3852.0 2542.0 0.990 ≈0.005 UB 直通，AIV 几乎无 MTE2

¹ PipeUtilization.csv 各核中位数（多数核持有 2 个 singleCore；临界路径核持有 3 个，aic_cube_max≈3813）；
平台：Ascend 950PR，Aicore Count=28，M=K=N=8192，baseM/N/K=256/256/64，singleCoreM/N/K=1024/1024/8192，numBlocks=28，Freq=1650 MHz。

排序满足：分离合计 > Scenario1 > Scenario2。相对 Scenario0 合计，Scenario1 / Scenario2 约节省 6.8% / 8.1%；Scenario2 相对 Scenario1 约再省 58 μs（≈1.5%）。

Scenario0：CV分离执行

Scenario Task Duration(μs) Block Dim Mix Block Dim aic_time(μs) aic_cube_time(μs) aic_cube_ratio aic_mte1_time(μs) aic_mte2_time(μs) aic_fixpipe_time(μs) aiv_time(μs) aiv_vec_time(μs) aiv_mte2_time(μs) aiv_mte3_time(μs)
MatMul 3891.2 28 56 2595.9 2542.0 0.980 818.4 2021.4 253.4 ≈0.46 0 ≈0 ≈0
ReLU 302.0 28 56 ≈0.8 0 0 ≈0 ≈0 ≈0 243.6 10.7 161.9 71.3

解读：

  • MatMul 侧 aic_cube_ratio≈0.98，说明 L1/L0 PingPong + 大包搬运已把 Cube 算力吃满。
  • Block Dim=28 对齐物理核；
  • ReLU 侧耗时主要在 MTE2（GM→UB）+ MTE3（UB→GM），aiv_vec_time 仅约 11 μs；分离执行时这段时间全部叠在 MatMul 之后。

Scenario1：CV融合 - GM 中转

Scenario Task Type Task Duration(μs) Block Dim Mix Block Dim aic_time(μs) aic_cube_time(μs) aic_cube_ratio aic_fixpipe_time(μs) aiv_time(μs) aiv_vec_time(μs) aiv_mte2_time(μs) aiv_mte3_time(μs) 备注
Scenario1 MIX_AIC 3909.7 28 56 2610.8 2542.0 0.975 267.3 2614.0 10.7 57.4 50.0 L0C→GM→UB

分析：

  • Scenario1总耗时3909.7us，与Scenario0的单独MatMul任务耗时（3891.2us）接近，表明此时AIV的Vector计算时间基本被AIC的Cube计算时间覆盖，体现了 CV 融合方案的性能收益。

下图展示了Scenario1场景下的算子流水执行图。CV融合算子执行的流程为，AIC侧完成MatMul计算后通过Fixpipe将结果写入GM，AIV侧从GM读取数据，完成ReLU激活函数计算。

image.png

Scenario2：CV融合 - UB 直通

Scenario Task Type Task Duration(μs) Block Dim Mix Block Dim aic_time(μs) aic_cube_time(μs) aic_cube_ratio aic_fixpipe_time(μs) aiv_time(μs) aiv_vec_time(μs) aiv_mte2_time(μs) aiv_mte3_time(μs) 备注
Scenario2 MIX_AIC 3852.0 28 56 2568.5 2542.0 0.990 211.6 2570.8 10.7 ≈0.005 80.2 L0C→UB

分析：

  • 从Scenario1改进至Scenario2，总耗时进一步减少约57us
  • 由于UB直通的方案中，数据搬运不再涉及GM通路搬运，Scenario2中aiv_mte2_time显著降低

下图展示了Scenario2场景下的算子流水执行图。改进为UB直通的CV融合算子执行后，在AIC侧完成MatMul计算后通过Fixpipe将结果直接写入UB，AIV侧在UB上进行ReLU激活函数计算。

image.png