Ascend C Matmul 输出随路反量化实战：基于 asc-devkit 的 int8 Matmul 反量化直调样例全解析

本文以 asc-devkit 仓库中的 matmul_quant 直调样例为主体，系统讲解 Ascend C 高级 API 中 Matmul"输出随路反量化"的工作原理：如何在 int8_t 矩阵乘完成后，利用 Cube 核将 C 矩阵从 CO1 搬出至 Global Memory 时随路执行反量化，并完整覆盖同一系数（SCALAR）与向量（TENSOR）两种反量化模式的核函数设置、Tiling 配置、编译选项、运行命令与精度验证全流程。读完本篇，你将能够独立搭建并跑通该样例，理解 SetQuantScalar/SetQuantVector 与 SetDequantType 的配套约束，并可将其移植到含权重量化的推理类算子开发中。

什么是 Matmul 输出随路反量化

Matmul 反量化场景指的是：在 Matmul 计算时，左、右矩阵输入为 int8_t（或 int4b_t）类型，输出为 half 类型。此时 Cube 核内部的矩阵乘累加结果为 int32 中间值，当输出 C 矩阵的数据从 L0C Buffer（CO1）搬出到 Global Memory 时，硬件会随路执行反量化操作，将累加结果乘以反量化系数（scale）转换为 half 输出。反量化系数有两种组织方式：

• 同一系数模式（SCALAR）：对输出矩阵的所有值采用同一个浮点系数进行反量化；
• 向量模式（TENSOR）：提供一个 shape 为 [1, N] 的参数向量，对输出矩阵的每一列采用该向量中对应列的系数进行反量化。

这与典型的量化推理流程吻合：网络权重经 int8 量化存储，Cube 完成 int8 矩阵乘后，在数据写回 GM 的流水环节由硬件自动完成 scale 还原，避免在算子逻辑中额外编写反量化代码，减少一次显式的 GM 读写开销。

API 文档中对该机制的正式描述见 SetQuantVector 与 SetDequantType。

样例规格与输入输出定义

本样例规格为 M = 1024, N = 1024, K = 1024，核函数名为 matmul_quant_custom。完整的张量规格如下（源自 README）：

样例类型(OpType)	Matmul
样例输入 a	shape [M, K]，int8_t，format ND，isTrans = false
样例输入 b	shape [K, N]，int8_t，format ND，isTrans = false
样例输入 bias	shape [1, N]，int32_t，format ND
样例输出 c	shape [M, N]，half，format ND
核函数名	matmul_quant_custom

对应地，在核函数源码 matmul_quant.asc 中这些规格被固化为编译期常量：

constexpr uint32_t M = 1024;
constexpr uint32_t N = 1024;
constexpr uint32_t K = 1024;
constexpr bool IS_TRANS_A = false;
constexpr bool IS_TRANS_B = false;
constexpr bool IS_BIAS = true;

目录结构

样例位于 matmul_quant 目录：

├── matmul_quant
│   ├── scripts
│   │   ├── gen_data.py         // 输入数据和真值数据生成脚本文件
│   │   └── verify_result.py    // 真值对比文件
│   ├── CMakeLists.txt          // 编译工程文件
│   ├── data_utils.h            // 数据读入写出函数
│   ├── matmul_quant.asc        // Ascend C样例实现 & 调用样例
│   └── README.md               // 样例说明文档

其中 matmul_quant.asc 一个文件同时包含了核函数（__global__ __cube__ 函数与 MatmulKernel 模板类）、Host 侧 Tiling 生成函数 GenerateTiling 以及 main 函数，是 Ascend C 直调类样例的典型组织方式。

Kernel 侧实现解析：反量化参数的两种设置方式

Matmul 对象的类型化构造

样例中的 MatmulKernel 模板类持有一个 AscendC::Matmul 对象，四个模板参数分别描述 A/B/C/Bias 在 GM 中的位置、格式与数据类型：

AscendC::Matmul<
    AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, AType>,
    AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BType>,
    AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, CType>,
    AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BiasType>, CFG_MDL>
    matmulObj;

核函数入口处通过 MatmulKernel<int8_t, int8_t, half, int32_t> matmulKernel; 实例化，正好对应规格表中的 int8_t 输入与 half 输出——这正是触发"输出随路反量化"路径的合法类型组合。随后通过 REGIST_MATMUL_OBJ 将 Matmul 对象注册到 TPipe 并绑定系统 Workspace：

REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulKernel.matmulObj, &tiling);

设置反量化参数

反量化参数的设置在 Process() 中完成，由编译宏 CUSTOM_QUANT_VECTOR 区分两种模式（源码见 matmul_quant.asc#L108-L114）：

#if defined(CUSTOM_QUANT_VECTOR)
    matmulObj.SetQuantVector(quantGlobal);
#else
    float quantFloat = 0.1f;
    uint64_t quantValue = static_cast<uint64_t>(*reinterpret_cast<int32_t*>(&quantFloat));
    matmulObj.SetQuantScalar(quantValue);
#endif

两个要点值得注意：

1. 调用时序约束：SetQuantScalar/SetQuantVector 必须在 IterateAll/Iterate 之前调用。样例中在 SetTensorA/SetTensorB/SetBias 之前先设置量化参数，最后依次调用 IterateAll(cGlobal) 和 End()。
2. SCALAR 模式的 uint64 打包：SetQuantScalar 接收 uint64_t 入参。样例直接把 0.1f 的位模式截断放入 uint64 的低 32 位。而 SetQuantVector 接口文档给出了完整的打包公式：quantTensor 为 64 位、初始为 0，scale 按 bit 位取高 19 位截断（即 scale & 0xFFFFE000）存储于 32 位处，并将第 46 位置 1，即 quantTensor = quantTensor | (scale & 0xFFFFE000) | (1 << 46)；若存在 offset（范围截断到 [-256, 255]），再保留其 9 位存储于 37~45 位。gen_data.py 生成向量模式的量化系数时执行的 np.bitwise_and(data, 0xFFFFE000) 正是这一公式中 scale 的截断处理，Host 与 Device 两侧对同一硬件格式的约定在此闭环。

向量模式下，量化向量以 GlobalTensor<uint64_t> 绑定到 GM，长度为 tiling.N 个 uint64_t（每个 uint64_t 承载一列系数）：

quantGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ uint64_t*>(quant), tiling.N);

多核切分与尾块处理

CalcOffset 按"M 方向优先"的二维网格切分核函数（源码见 matmul_quant.asc#L134-L154）：mCoreIndex = blockIdx % mSingleBlocks、nCoreIndex = blockIdx / mSingleBlocks，各张量按 singleCoreM/singleCoreN 计算各自的元素偏移；bias 与量化向量均按 N 方向偏移，保证每个核只读取自己负责列区间的 bias 和量化系数。

Process() 中还处理了 M/N 方向不能被整除时的尾块：

int tailN = tiling.N - nCoreIndex * tiling.singleCoreN;
tailN = tailN < tiling.singleCoreN ? tailN : tiling.singleCoreN;
int tailM = tiling.M - mCoreIndex * tiling.singleCoreM;
tailM = tailM < tiling.singleCoreM ? tailM : tiling.singleCoreM;
if (tailM < tiling.singleCoreM || tailN < tiling.singleCoreN) {
    matmulObj.SetTail(tailM, tailN);
}

Tiling 侧实现解析：SetDequantType

Host 侧的 GenerateTiling 函数使用 matmul_tiling::MultiCoreMatmulTiling 生成 TCubeTiling，关键配置序列如下（源码见 matmul_quant.asc#L170-L203）：

tilingApi.SetDim(ascendcPlatform->GetCoreNumAic());
tilingApi.SetAType(
    matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_INT8, IS_TRANS_A);
tilingApi.SetBType(
    matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_INT8, IS_TRANS_B);
tilingApi.SetCType(
    matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16);
tilingApi.SetBiasType(
    matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_INT32);

#if defined(CUSTOM_QUANT_VECTOR)
    tilingApi.SetDequantType(matmul_tiling::DequantType::TENSOR); // set TENSOR quant mode
#else
    tilingApi.SetDequantType(matmul_tiling::DequantType::SCALAR); // set SCALAR quant mode
#endif

tilingApi.SetOrgShape(M, N, K);
tilingApi.SetShape(M, N, K);
tilingApi.EnableBias(IS_BIAS);
tilingApi.SetBufferSpace(-1, -1, -1);

int64_t res = tilingApi.GetTiling(tilingData);

SetDequantType 在 Tiling 结构中落位的枚举定义见 matmul_tiling_base.h#L168-L171：

enum class DequantType : int32_t {
    SCALAR = 0, ///< Scalar type, value is 0
    TENSOR = 1, ///< Tensor type, value is 1
};

从源码结构看，TCubeTiling 中的量化模式字段默认值为 SCALAR（见 matmul_tiling_base.h#L632），因此即使不显式调用 SetDequantType，Tiling 也处于同一系数模式；但按 SetDequantType 接口文档的约束说明，Tiling 侧设置的模式必须与核函数侧调用的 SetQuantScalar/SetQuantVector 保持一致，否则行为未定义。样例通过 CMake 统一注入 CUSTOM_QUANT_VECTOR 宏（两种模式共用同一套 #if 分支），从工程上规避了 Host/Device 两侧配置不一致的风险。

Host 侧启动与数据流

main 函数的完整启动链路（源码见 matmul_quant.asc#L205-L319）：

1. 根据架构选择平台实例：CPU 调试模式下 dav-2201 对应 Ascend910B1，dav-3510 对应 Ascend950PR_9589；NPU 运行模式则使用 PlatformAscendCManager::GetInstance()。
2. 计算各张量字节数：a/b 为 int8 各 1024×1024 字节，bias 为 int32 共 N 个，输出 c 按 uint16（即 half 的位宽）计 N×M 个元素，向量模式下量化向量为 N 个 uint64_t。
3. 调用 GenerateTiling 生成 TCubeTiling，并从中取出 usedCoreNum 作为启动的 numBlocks。
4. 依次读取 ./input/x1_gm.bin、./input/x2_gm.bin、./input/bias_gm.bin（以及向量模式下的 ./input/quant_vector_gm.bin）到 Host 内存，再 aclrtMemcpy 到 Device。注意向量模式的量化文件仅在定义 CUSTOM_QUANT_VECTOR 时才分配并上传，与核函数侧的宏分支严格对应。
5. 使用内核调用符 <<<>>> 启动：matmul_quant_custom<<<numBlocks, 0, stream>>>(aDevice, bDevice, biasDevice, cDevice, quantDevice, workspaceDevice, tilingDevice);，同步后将 c 写回 ./output/output.bin。

编译选项与运行流程

在样例根目录下执行如下步骤，编译并执行样例。

配置环境变量

请根据当前环境上 CANN 开发套件包的安装方式，配置环境变量：

source ${install_path}/cann/set_env.sh

说明： ${install_path} 为 CANN 包安装目录，未指定安装目录时默认安装至 /usr/local/Ascend 下。

样例执行

在样例目录下执行如下命令：

# -DQUANT_MODE=0：使能同一系数的反量化模式；
# -DQUANT_MODE=1：使能向量反量化模式；
# -m=0：使能同一系数的反量化模式；
# -m=1：使能向量反量化模式；
mkdir -p build && cd build;    # 创建并进入build目录
cmake -DQUANT_MODE=0 -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j;    # 编译工程，以使能同一系数的反量化模式为例，默认npu模式
python3 ../scripts/gen_data.py -m=0   # 生成测试输入数据，以使能同一系数的反量化模式为例
./demo                        # 执行编译生成的可执行程序，执行样例
python3 ../scripts/verify_result.py output/output.bin output/golden.bin    # 验证输出结果是否正确，确认算法逻辑正确

使用 CPU 调试或 NPU 仿真模式时，添加 -DCMAKE_ASC_RUN_MODE=cpu 或 -DCMAKE_ASC_RUN_MODE=sim 参数即可，示例：

cmake -DQUANT_MODE=0 -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式
cmake -DQUANT_MODE=0 -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式

注意： 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 rm CMakeCache.txt 后重新 cmake。

编译选项说明

选项	可选值	说明
CMAKE_ASC_RUN_MODE	npu（默认）、cpu、sim	运行模式：NPU 运行、CPU 调试、NPU 仿真
CMAKE_ASC_ARCHITECTURES	dav-2201（默认）、dav-3510	NPU 架构：dav-2201 对应 Atlas A2 系列产品和 Atlas A3 系列产品，dav-3510 对应 Ascend 950PR&950DT 系列产品

从 CMakeLists.txt 可以看到，CMAKE_ASC_RUN_MODE 与 CMAKE_ASC_ARCHITECTURES 以 CACHE 变量形式声明默认值（npu、dav-2201），QUANT_MODE=1 时通过 target_compile_definitions 注入 CUSTOM_QUANT_VECTOR 宏，CMAKE_ASC_ARCHITECTURES 最终映射为 ASC 编译器的 --npu-arch 选项。工程链接了 tiling_api、register、platform 等库，分别对应 Tiling 生成、Matmul 对象注册与平台查询能力。

执行结果

执行结果如下，说明精度对比成功：

test pass!

测试数据生成与精度验证逻辑

gen_data.py：真值如何构造

gen_data.py 用 NumPy 模拟了与 Cube 硬件完全一致的数值流程：

• 输入数据：x1_gm、x2_gm 为 [-10, 10) 的随机 int8 矩阵，bias_gm 为 [-2, 2) 的随机 int32 向量；
• 整数累加：先以 int32 计算 x1 @ x2 + bias，再转 float16，与硬件"CO1 中 int32 累加后随路反量化转 half"的路径对齐；
• 模式 0（SCALAR）：golden 整体乘以 0.1，与核函数侧 quantFloat = 0.1f 一致；
• 模式 1（TENSOR）：随机生成 [0.1, 2.0) 的 float32 向量，按硬件格式要求执行 scale & 0xFFFFE000 位截断后写入 quant_vector_gm.bin（以 uint64 位模式落盘），golden 则用截断后的向量按列相乘——这一细节保证真值计算使用的 scale 与硬件实际使用的 scale 逐位一致，避免验证误差。

verify_result.py：容差与判定

verify_result.py 以 float16 读取 output.bin 与 golden.bin，使用 np.isclose(rtol=1e-3, atol=1e-3, equal_nan=True) 逐元素比较，打印差异元素（最多 100 条，含期望值、实际值与相对误差），并统计错误比例；错误比例超过 1e-3 时打印全量数据并以非零码退出，全部通过时输出 test pass!。

仓库的 presmoke 预冒烟用例 run.sh 同样采用"构建（-DQUANT_MODE=0）→ 运行 ./demo → verify_result.py 比对"的三段式流程，可作为 CI 环境下复现本样例的参照。

支持的产品及 CANN 软件版本

产品	CANN软件版本
Ascend 950PR&950DT 系列产品	>= CANN 9.1.0
Atlas A3 系列产品	>= CANN 9.0.0
Atlas A2 系列产品	>= CANN 9.0.0

对应地，CMAKE_ASC_ARCHITECTURES 中 dav-2201 覆盖 Atlas A2/A3 系列，dav-3510 覆盖 Ascend 950PR&950DT 系列；在 950 系列上运行需使用 CANN 9.1.0 及以上版本。

小结与关键文件索引

本样例展示了 Ascend C 高级 API 下 Matmul 反量化场景的最小完整闭环：

• 核函数侧：类型组合 int8 × int8 → half 触发随路反量化；SetQuantScalar（同一系数）或 SetQuantVector（向量）在 IterateAll 前完成系数注入；
• Tiling 侧：SetDequantType(DequantType::SCALAR/TENSOR) 必须与核函数侧接口一一对应；
• 工程侧：QUANT_MODE 编译选项通过 CUSTOM_QUANT_VECTOR 宏同时驱动 Host 数据准备与 Device 参数设置，保证两侧一致。

关键文件索引：

文件	说明
matmul_quant.asc	核函数、Tiling 生成与 Host 启动主文件
CMakeLists.txt	编译工程，QUANT_MODE / 运行模式 / 架构选项
gen_data.py	输入与真值生成脚本
verify_result.py	精度验证脚本
SetQuantVector 接口文档	向量量化参数设置与 uint64 打包公式
SetDequantType 接口文档	量化/反量化模式 Tiling 接口
matmul_tiling_base.h	DequantType 枚举源码定义