Ascend C 归约算子开发实战：基于 ReduceMax 高阶 API 实现矩阵列最大值样例全解析

本文以 asc-devkit 开源仓库中的 ReduceMax 样例 为主体，完整讲解如何基于 Ascend C 归约高阶 API ReduceMax 实现"对矩阵每一列求最大值"的算子开发全过程，覆盖 Kernel 侧实现、Tiling 参数计算、Host 侧调用与三种运行模式（NPU/CPU 调试/仿真）的编译运行方法。读完本文，你将掌握 ReduceMax/ReduceMin/ReduceMean/ReduceSum 归约接口的统一调用方式、临时空间获取接口 GetReduceMaxMaxMinTmpSize 的底层原理，以及一套可直接复用的归约算子样例工程。

一、样例概述：用高阶归约 API 完成矩阵列归约

本样例位于 examples/01_simd_cpp_api/04_advanced_api/04_reduce/reducemax，其核心目标是在归约（Reduction）场景下，基于 Ascend C 高阶 API ReduceMax 对输入矩阵的每一列求最大值：

• 输入：[32, 136] 的 float 矩阵（32 行、136 列），ND 格式；
• 输出：[1, 136] 的结果行向量（等价于一维的 136 个值），float 类型，ND 格式；
• 核函数名：reduce_custom。

从样例代码（reducemax.asc）可以看到，归约模式采用 PATTERN = 1，即 AscendC::Pattern::Reduce::RA——沿**第一个维度（外层轴，即 32 行方向）**归约，对每一列的 32 个元素求最大值，最终得到 136 个结果值。这与数据生成脚本 gen_data.py 中 y_gm = np.max(x_gm, axis=0) 的真值计算逻辑完全对应。

整个计算流程是典型的 Ascend C 算子三段式：输入数据先通过 DataCopyPad 从全局内存（GM）搬运到片上存储（UB），随后调用 ReduceMax 高阶接口在片上完成归约计算，最后将结果从片上搬回全局内存。

二、归约接口家族与 ReduceMax 接口速览

2.1 一个样例，切换四种归约算子

ReduceMax 只是 Ascend C 归约高阶接口家族的一员。仓库中同一目录下还提供了 ReduceMin、ReduceMean、ReduceSum 等归约接口，它们在 include/adv_api/reduce/reduce.h 中集中声明，调用方式与 ReduceMax 完全一致。以本样例 reducemax.asc 为基础，只需两处替换即可切换归约语义：

目标接口	替换调用	替换临时空间接口
ReduceMin（求最小值）	AscendC::ReduceMax → AscendC::ReduceMin	GetReduceMaxMaxMinTmpSize → GetReduceMinMaxMinTmpSize
ReduceMean（求平均值）	AscendC::ReduceMax → AscendC::ReduceMean	GetReduceMaxMaxMinTmpSize → GetReduceMeanMaxMinTmpSize
ReduceSum（求和）	AscendC::ReduceMax → AscendC::ReduceSum	GetReduceMaxMaxMinTmpSize → GetReduceSumMaxMinTmpSize

此外，reduce.h 中还定义了 ReduceProd（求积）、ReduceAny（是否存在真值）、ReduceAll（是否全为真值）等接口；仓库 impl/adv_api/detail/reduce 目录下按算子分别组织实现，包括 reduce_max、reduce_min、reduce_sum、reduce_mean、reduce_prod、reduce_any、reduce_all、reduce_xor_sum 等子目录，每个目录按架构提供 _v220_impl.h 与 _3510_impl.h 两套实现。

2.2 ReduceMax 接口原型与参数

ReduceMax 支持两种调用形态（定义见 include/adv_api/reduce/reduce.h）：

• 通过 sharedTmpBuffer 显式传入临时空间（样例采用该方式，便于开发者自行管理、复用内存）：

  template <class T, class pattern, bool isReuseSource = false>
  __aicore__ inline void ReduceMax(const LocalTensor<T>& dstTensor, const LocalTensor<T>& srcTensor,
                                   const LocalTensor<uint8_t>& sharedTmpBuffer,
                                   const uint32_t srcShape[], bool srcInnerPad);
  

• 由接口框架自动申请临时空间（开发者只需按 GetReduceMaxMaxMinTmpSize 预留空间）：

  template <class T, class pattern, bool isReuseSource = false>
  __aicore__ inline void ReduceMax(const LocalTensor<T>& dstTensor, const LocalTensor<T>& srcTensor,
                                   const uint32_t srcShape[], bool srcInnerPad);
  

关键参数语义如下（详见仓库 API 文档 ReduceMax.md）：

参数	说明
T（模板参数）	操作数数据类型。Atlas A2/A3 系列支持 half、float；Ascend 950PR&950DT 系列还支持 int8_t/uint8_t/int16_t/uint16_t/bfloat16_t/int32_t/uint32_t/int64_t/uint64_t 等
pattern（模板参数）	归约计算轴：A 表示 Normal 轴（不归约），R 表示 Reduce 轴（归约），由与输入维度数量相同的 A/R 字母组合而成，当前只支持 AR 和 RA
isReuseSource（模板参数）	是否允许修改源操作数，默认 false。置 true 时接口复用 src 内存，可节省片上空间，但会改写输入数据
dstTensor	目的操作数，LocalTensor，TPosition 支持 VECIN/VECCALC/VECOUT
srcTensor	源操作数，数据类型需与 dstTensor 一致
sharedTmpBuffer	临时缓存，用于存储中间变量，大小由 GetReduceMaxMaxMinTmpSize 获取
srcShape	uint32_t 数组，表示源操作数 shape，维度必须与 pattern 一致（当前只支持二维）
srcInnerPad	最内层轴数据是否 32 字节对齐；Atlas A2/A3 系列当前只支持 true

2.3 支持的产品及 CANN 软件版本

本样例在归约语义和临时空间计算上与不同架构强相关，支持的硬件与软件版本如下（与 ReduceMax.md 中"产品支持情况"章节一致）：

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

需要说明的是，API 文档同时明确 Atlas 推理系列产品（AI Core/Vector Core）及 Atlas 训练系列产品不支持 ReduceMax 高阶接口。

三、目录结构与样例规格

3.1 目录结构

examples/01_simd_cpp_api/04_advanced_api/04_reduce/reducemax/
├── scripts
│   └── gen_data.py             // 输入数据和真值（golden）数据生成脚本
├── CMakeLists.txt              // 编译工程文件
├── data_utils.h                // 数据读入写出工具函数（ReadFile/WriteFile）
├── reducemax.asc               // Ascend C 样例实现 & Host 侧调用样例（含 Kernel 与 Tiling）
└── README.md                   // 样例说明文档

其中 reducemax.asc 是一个"单文件全流程"样例：既包含 __global__ __vector__ 修饰的核函数 reduce_custom，也包含 Host 侧的 Tiling 生成函数与 main 主流程，方便开发者完整理解算子各环节的衔接。

3.2 样例规格

类别	项目	取值
样例类型（OpType）	—	reducemax
样例输入	name / shape / data type / format	x / [32, 136] / float / ND
样例输出	name / shape / data type / format	y / [1, 136] / float / ND
核函数名	—	reduce_custom

四、Kernel 侧实现解析：从数据搬运到 ReduceMax 计算

4.1 核函数入口与模板实例化

核函数入口（reducemax.asc）定义为：

extern "C" __global__ __vector__ void reduce_custom(
    __gm__ uint8_t* src_gm, __gm__ uint8_t* dst_gm, __gm__ uint8_t* workspace, __gm__ uint8_t* tiling)
{
    AscendC::TPipe pipe;
    VecTiling tilingData;
    CopyTiling(&tilingData, tiling);
    typedef ReduceType<float, 4352> xType;   // 输入 4352 个 float 元素（32×136）
    typedef ReduceType<float, 136> yType;    // 输出 136 个 float 元素
    KernelReduceMax<xType, yType, 1, 1> op;  // Pattern=1(RA)，IsReuseSource=1
    op.Init(src_gm, dst_gm, tilingData.aLength, tilingData.rLength, tilingData.srcInnerPad,
            tilingData.sharedTmpBufferSize, &pipe);
    op.Process();
}

核函数模板参数 Pattern 与 IsReuseSource 均为编译期常量（样例中固定为 1, 1），因此 Compute 内部可以通过 if constexpr 在编译期展开出最优分支，避免运行期判断开销。

4.2 初始化：片上 Buffer 与对齐计算

KernelReduceMax::Init（reducemax.asc）完成以下工作：

• 通过 SetGlobalBuffer 将全局内存指针绑定为 GlobalTensor<T_TYPE>；
• pipe->InitBuffer(inQueueX, BUFFER_NUM, Align32B<X_TYPE>() * sizeof(T_TYPE)) 为输入队列分配片上 Buffer，其中 Align32B() 把元素个数按 32 字节对齐（(extent + 32/sizeof(T) - 1) / (32/sizeof(T)) * (32/sizeof(T))）；
• 当 IsReuseSource == 0（不复用输入）时，额外初始化 tmpQueue 作为共享临时 Buffer，容量取自 Tiling 参数 sharedTmpBufferSize；
• 计算 aLengthAlign、rLengthAlign 等对齐后的长度，供带 padding 的数据搬运使用。

4.3 数据搬运 CopyIn/CopyOut：处理 32B 对齐与 padding

CopyIn 使用 DataCopyPad 完成 GM→UB 搬运（reducemax.asc），并按 srcInnerPad 分两条路径：

• srcInnerPad == 0：输入本身无需 padding，使用 blockCount = 1、blockLen = 全部元素字节数 的整块拷贝；
• srcInnerPad == 1：输入需要 padding，AR 模式下按 aLength 个 block 逐行拷贝，行内 blockLen = rLength * sizeof(T)，通过 rightPadding = rLengthAlign - rLength 补齐到 32 字节对齐；RA 模式则对第一维做右 padding。

CopyOut 以对称方式把结果从 outQueueY 写回 GM。data_utils.h 中提供的 ReadFile/WriteFile 用于 Host 侧读写二进制数据文件。

4.4 Compute 核心：ReduceMax 两种归约模式

Compute（reducemax.asc）是核心计算逻辑，按 Pattern 分两条路径：

• AR 模式（Pattern == 0）：沿第二个维度（内轴）归约。srcShape = {aLength, rLength}，例如 ReduceMax<float, AscendC::Pattern::Reduce::AR, isReuse>(yLocal, xLocal, srcShape, true)；
• RA 模式（Pattern == 1）：沿第一个维度（外轴）归约。srcShape = {rLength, aLength}，即本样例采用的模式：ReduceMax<T_TYPE, AscendC::Pattern::Reduce::RA, isReuse>(yLocal, xLocal, srcShape, true)。

同时按 isReuse 区分临时空间的使用方式：

if constexpr (isReuse) {
    // 复用输入 buffer 的版本：不额外申请临时空间，接口内部直接改写输入
    AscendC::ReduceMax<T_TYPE, AscendC::Pattern::Reduce::RA, isReuse>(yLocal, xLocal, srcShape, true);
} else {
    // 显式传入共享临时 buffer 的版本
    AscendC::LocalTensor<uint8_t> sharedTmpBuffer = tmpQueue.AllocTensor<uint8_t>();
    AscendC::ReduceMax<T_TYPE, AscendC::Pattern::Reduce::RA, isReuse>(
        yLocal, xLocal, sharedTmpBuffer, srcShape, true);
    tmpQueue.FreeTensor(sharedTmpBuffer);
}

从底层实现看（reduce_max_v220_impl.h），ReduceMaxImpl 会先通过 static_assert(SupportType<T, half, float>(), ...) 与 static_assert(SupportType<pattern, Pattern::Reduce::AR, Pattern::Reduce::RA>(), ...) 在编译期校验类型与归约模式，再按模式分派：AR 走 BlockReduceByLastAxis<T, isReuseSource, ApiMode::API_MODE_MAX, Max<T, false>>，RA 走 BinaryReduceByFirstAxis<T, isReuseSource, Max<T, false>>。Ascend 950 系列的实现（reduce_max_3510_impl.h）则针对 32 字节对齐与非对齐、单/双寄存器位宽等场景细分了更多计算路径。

五、Tiling 侧实现：用 GetReduceMaxMaxMinTmpSize 精确计算临时空间

5.1 Tiling 数据结构

样例自定义了 ReduceCustomTilingData（reducemax.asc），通过 BEGIN_TILING_DATA_DEF/TILING_DATA_FIELD_DEF 宏定义 8 个 uint32_t 字段：

字段	含义
aLength	归约后输出的长度（非归约轴长度）
rLength	归约维度长度
xExtent	输入 Tensor 的元素个数
yExtent	输出 Tensor 的元素个数
srcInnerPad	输入是否需要 padding：0=不 padding，1=padding
pattern	归约模式：0=AR（沿第二个维度归约），1=RA（沿第一个维度归约）
isReuseSource	是否复用输入 buffer：0=不复用，1=复用
sharedTmpBufferSize	共享临时 buffer 大小（Kernel 侧据此 InitBuffer(tmpQueue, ...)）

Host 侧 ComputeTiling 将这些参数与临时空间大小一起写入 Tiling 结构，再由 GetTilingBuf 序列化到内存 Buffer，最终通过 aclrtMemcpy 拷贝到设备侧供核函数读取。

5.2 GetReduceMaxMaxMinTmpSize：临时空间的计算原理

GetReduceMaxMaxMinTmpSize 是 host 侧接口，用于获取 ReduceMax 完成计算所需的最大/最小临时空间（原型声明见 reduce_tiling.h）：

void GetReduceMaxMaxMinTmpSize(const AscendC::TensorShape& srcShape,
                               const AscendC::TensorDataType dataType,
                               ReducePattern pattern,
                               bool isSrcInnerPad,
                               bool isReuseSource,
                               uint32_t& maxValue,
                               uint32_t& minValue);

参数要点（详见 GetReduceMaxMaxMinTmpSize.md）：

• srcShape 与 dataType 必须与 Kernel 侧 ReduceMax 调用保持一致；
• pattern 只支持 ReducePattern::AR 与 ReducePattern::RA（ReducePattern 枚举同时定义 R/ARA/ARAR/.../RAR/RARA/... 等更多维度组合，但当前归约接口仅使用 AR/RA）；
• isSrcInnerPad 与 isReuseSource 同样需与 Kernel 侧参数一致；
• maxValue 是参考上限（超出该值的空间不会被接口使用），minValue 是功能正确性下限——预留空间不能小于 minValue；
• 当前实现中最大/最小临时空间相等，在范围内增大空间可获得一定性能收益，开发者需结合 UB 剩余空间选取。

从实现源码（reduce_tiling.cpp）可以还原其计算逻辑：

1. 平台与类型校验：通过 PlatformAscendCManager 获取当前 NPU 架构。Ascend 950 系列（DAV_3510/DAV_9201）允许 int8_t/uint8_t/int16_t/uint16_t/half/bfloat16_t/int32_t/uint32_t/float/int64_t/uint64_t；其余平台（如 dav-2201）仅允许 float 和 half。
2. 公共计算函数 GetReduceCommonMaxMinTmpSize（reduce_tiling.cpp）：
   • 先校验 shape 必须为二维、pattern 必须为 AR/RA；
   • 若 isReuseSource == true，直接返回 maxValue = minValue = 0——因为复用输入内存后接口不再需要额外临时空间，这也解释了样例中 IS_REUSE_SOURCE = 1 时 Tiling 传入的 sharedTmpBufferSize 为 0 且 Kernel 侧不再 InitBuffer(tmpQueue)；
   • AR 模式下，按内轴长度 last 与每 block/每 repeat 可承载元素数的关系，分三段计算临时空间（0、first * elePerBlk * typeSize、first * elePerRep * typeSize）；
   • RA 模式下，先通过 FindK(first) 求得归约分段数 k，再乘以对齐后的行字节数。

5.3 样例中的实际调用

样例在 ComputeTiling 中按如下方式调用（reducemax.asc）：

AscendC::TensorShape srcShape(shapeVec);   // {aLength, rLength}
AscendC::TensorDataType dataType = AscendC::TensorDataType::DT_FLOAT;
AscendC::ReducePattern reducePattern = (pattern == 0) ? AscendC::ReducePattern::AR : AscendC::ReducePattern::RA;
AscendC::GetReduceMaxMaxMinTmpSize(srcShape, dataType, reducePattern, isSrcInnerPad, isReuse, maxTmpSize, minTmpSize);
uint32_t localWorkspaceSize = minTmpSize;  // 使用最小临时空间
tiling.set_sharedTmpBufferSize(localWorkspaceSize);

样例选用 minTmpSize 作为 sharedTmpBufferSize，在保证功能正确的前提下最大程度节省片上内存；由于本样例开启了 isReuseSource = true，该值实际为 0，即完全复用输入 Buffer，不再额外占用临时空间。

六、Host 侧调用实现：数据准备、内核启动与精度校验

6.1 真值与输入数据生成

gen_data.py 使用 NumPy 生成输入与真值：

np.random.seed(0)
x_gm = np.random.uniform(-10, 10, [32, 136]).astype(np.float32)  # 输入 [32, 136]
y_gm = np.max(x_gm, axis=0).astype(np.float32)                    # 沿 axis=0 求列最大值，输出 [136]
x_gm.tofile("./input/input_x.bin")     # 写输入
y_gm.tofile("./output/golden.bin")     // 写真值

固定随机种子（seed=0）保证每次运行生成一致的数据，便于结果可复现；np.max(x_gm, axis=0) 与 Kernel 侧 RA 模式的归约语义严格对应。

6.2 内核启动流程

main（reducemax.asc）完整演示了基于 ACL（Ascend Computing Language）Runtime 的调用链：

1. GenerateTiling 生成 Tiling Buffer；
2. aclInit → aclrtSetDevice → aclrtCreateContext → aclrtCreateStream 完成 Runtime 初始化；
3. aclrtMallocHost/aclrtMalloc 分别申请 Host 与 Device 内存（workspace 大小 16 * 512 * 512 字节）；
4. 读取 ./input/input_x.bin，将输入与 Tiling 数据 aclrtMemcpy 到设备侧；
5. 使用内核调用符启动核函数：

   reduce_custom<<<numBlocks, 0, stream>>>(inputDevice, outputDevice, workspaceDevice, tilingDevice);
   

   其中 numBlocks = 1，表示单核执行；
6. aclrtSynchronizeStream 同步等待执行完成，将结果 aclrtMemcpy 回 Host，写入 ./output/output.bin；
7. 依次 aclrtFree/aclrtFreeHost 释放内存，销毁 stream/context/device 并 aclFinalize。

6.3 精度对比

CompareResult（reducemax.asc）读取 ./output/golden.bin 与计算结果逐元素比对，采用绝对误差与相对误差双重判据：EPS = 1e-4，当 |a - b| > EPS 且 |a - b| / |b| > EPS 时记为错误。全部通过则打印 test pass!。

该样例配套的自动化测试位于 tests/api/adv_api/reduce/reduce_max/test_operator_reduce_max.cpp（以及 950 系列对应的 test_operator_reduce_max_c310.cpp），可进一步验证接口在更多 shape 与参数组合下的正确性。

七、编译与运行：NPU 运行、CPU 调试与 NPU 仿真

7.1 环境准备

请先根据 CANN 开发套件包的安装方式（见 docs/zh/quick_start.md）配置环境变量：

source ${install_path}/cann/set_env.sh

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

7.2 NPU 模式编译与运行

在样例根目录 examples/01_simd_cpp_api/04_advanced_api/04_reduce/reducemax 下执行：

mkdir -p build && cd build;
cmake .. -DCMAKE_ASC_ARCHITECTURES=dav-2201;make -j; # 默认npu模式
python3 ../scripts/gen_data.py   # 生成测试输入数据
./demo

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

test pass!

7.3 CPU 调试与 NPU 仿真模式

使用 CPU 调试或 NPU 仿真模式时，在 cmake 阶段追加 -DCMAKE_ASC_RUN_MODE=cpu 或 -DCMAKE_ASC_RUN_MODE=sim：

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

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

7.4 编译选项说明

选项	可选值	说明
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 系列产品

7.5 工程文件要点

CMakeLists.txt 揭示了该样例工程的构建方式：

• find_package(ASC REQUIRED) 引入 Ascend C 编译工具链，project(kernel_samples LANGUAGES ASC CXX) 将 .asc 文件纳入 ASC 语言编译；
• CMAKE_ASC_ARCHITECTURES 以 CACHE STRING 形式暴露给用户，并通过 --npu-arch=${CMAKE_ASC_ARCHITECTURES} 传入编译选项；
• 可执行文件 demo 链接了 tiling_api、register、platform、unified_dlog、graph_base 等库，其中 tiling_api 提供 GetReduceMaxMaxMinTmpSize 等 Tiling 接口的实现。

八、约束与注意事项

结合 API 文档与源码实现，使用 ReduceMax 时需注意以下约束（详见 ReduceMax.md 的"约束说明"）：

• 地址对齐：当 srcInnerPad 为 true 时，输入向量内轴必须为 32 字节整数倍；若原始内轴长度 n 的字节数不是 32 的倍数，需自行向上补齐：inner = (n * sizeof(T) + 32 - 1) / 32 * 32 / sizeof(T)；
• 地址重叠：不支持源操作数与目的操作数地址重叠，也不支持 sharedTmpBuffer 与源/目的操作数地址重叠；
• 平台差异：Atlas A2/A3 系列 srcInnerPad 当前只支持 true，ReduceMax 仅支持 half/float；Ascend 950 系列数据类型支持更广，srcInnerPad 参数预留、由接口内部根据 srcShape 与 pattern 自行判断 32 字节对齐情况；
• 临时空间下限：显式传入 sharedTmpBuffer 时，空间大小不得小于 GetReduceMaxMaxMinTmpSize 返回的 minValue，否则无法保证功能正确；
• 编译期校验：ReduceMaxImpl 通过 static_assert 在编译期拒绝不支持的数据类型与非 AR/RA 的归约模式，出现编译错误时请优先核对这两项模板参数。

九、延伸阅读

若希望进一步深入，可继续阅读仓库内以下资料：

• 归约高阶接口声明与实现：include/adv_api/reduce/reduce.h、impl/adv_api/detail/reduce；
• Tiling 接口声明与实现：include/adv_api/reduce/reduce_tiling.h、impl/adv_api/tiling/reduce/reduce_tiling.cpp；
• API 参考文档：ReduceMax.md、GetReduceMaxMaxMinTmpSize.md；
• 自动化测试用例：tests/api/adv_api/reduce/reduce_max/test_operator_reduce_max.cpp；
• 相邻归约样例：同目录下的 reducemin、reducemean、reducesum 等样例，均可在本样例基础上通过接口替换快速复用。