企业官网建设流程全解析

CANN ops-math 算子指南：aclnnGcd 两段式接口详解与 NPU 最大公约数计算实战

【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-math

导读

aclnnGcd是 CANN ops-math 数学算子库中用于在 NPU（Ascend 芯片）上按元素（element-wise）计算两个整数张量最大公约数（Greatest Common Divisor，GCD）的单算子 API。本文以 math/gcd/docs/aclnnGcd.md 为核心，结合仓库中该算子的 op_api、op_host、op_kernel 与测试代码，完整讲解其功能语义、产品支持范围、两段式接口（aclnnGcdGetWorkspaceSize+aclnnGcd）的参数与错误码、源码级实现原理以及可直接运行的调用示例。读完本文，你将掌握在 CANN 环境下正确调用aclnnGcd完成整数张量逐元素最大公约数计算的完整流程，并能理解其背后的数据类型推导、broadcast 校验与 AscendC kernel 调度机制。

功能说明：按元素计算整数最大公约数

aclnnGcd对给定的输入张量self和other计算逐元素（element-wise）维度的最大公约数。数学上，两个整数 a、b 的最大公约数是能同时整除 a 和 b 的最大正整数；该算子将这一运算推广到张量的每个对应元素上，且self与other都必须是整数类型张量。

out[i] = gcd(self[i], other[i])

从算子注册信息看，该算子在算子定义（OpDef）层面描述为：输入x1、x2，输出y，见 op_host/gcd_def.cpp。gcd_def.cpp中对x1、x2、y均声明了UINT8、INT8、INT16、INT32、INT64五种整数数据类型与 ND 格式支持，并开启动态 shape、动态 rank（DynamicRankSupportFlag(true)）、动态编译静态标志（DynamicCompileStaticFlag(true)）等配置。

在 kernel 实现层面（op_kernel/arch35/gcd_dag.h），算子对 int64 使用专门的GcdVecInt64SIMT 向量函数，对 int32/int16/int8/uint8 使用通用的GcdVec模板函数，二者均采用基于二进制移位与减法的 GCD 算法（处理了负数取绝对值、零值特判等边界），每个线程按threadIdx.x步进blockDim.x的方式遍历元素。

产品支持情况

根据 math/gcd/docs/aclnnGcd.md 中声明的产品支持矩阵：

说明：产品支持情况以文档声明为准，实际使用前请结合 CANN 版本配套关系确认目标芯片是否在支持范围内。

两段式接口与函数原型

aclnnGcd属于 CANN 单算子 API（aclnn 接口），遵循两段式接口调用规范（详见 docs/zh/context/two_phase_api.md）：必须先调用第一段接口aclnnGcdGetWorkspaceSize获取计算所需的 workspace 大小以及包含算子计算流程的执行器（executor），再按 workspaceSize 在 Device 侧申请内存后，调用第二段接口aclnnGcd真正执行计算。

第一段接口原型：

aclnnStatus aclnnGcdGetWorkspaceSize( const aclTensor* self, const aclTensor* other, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)

第二段接口原型：

aclnnStatus aclnnGcd( void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)

其中workspace指除输入/输出外算子在 NPU 上完成计算所需的临时内存。需要特别注意的是，第二段接口aclnnGcd(...)不能重复调用，同一 executor 只能执行一次，重复调用会出现异常。

aclnnGcdGetWorkspaceSize 参数详解

第一段接口完成入参校验、workspace 大小计算与执行器构建，各参数说明如下：

平台相关的数据类型差异

在上述输入 dtype 基础上，存在平台差异：

• Atlas A3 训练系列产品 / Atlas A3 推理系列产品：不支持 UINT8、INT8 数据类型。

这一差异在 op_api/aclnn_gcd.cpp 的源码中有更细粒度的体现。实现中按 SoC 版本维护了不同的 dtype 支持列表：

static const std::initializer_list<op::DataType> ASCEND910_DTYPE_SUPPORT_LIST = {op::DataType::DT_INT32}; static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_SUPPORT_LIST = { op::DataType::DT_INT32, op::DataType::DT_INT16, op::DataType::DT_INT64}; static const std::initializer_list<op::DataType> ASCEND950_DTYPE_SUPPORT_LIST = { op::DataType::DT_UINT8, op::DataType::DT_INT8, op::DataType::DT_INT16, op::DataType::DT_INT32, op::DataType::DT_INT64};

即：Atlas 训练系列产品（Ascend 910）仅支持 INT32；Atlas A2/A3 系列（910B 等）支持 INT32、INT16、INT64；Ascend 950（RegBase 平台）支持 UINT8、INT8、INT16、INT32、INT64 五种整数类型。头文件 op_api/aclnn_gcd.h 中亦注明"推导后数据类型支持 INT32、INT16(Ascend910B)"。因此在实际开发中，应优先采用 INT32 以保证跨平台通用性。

参数校验的实现逻辑

从 op_api/aclnn_gcd.cpp 可以看到第一段接口的完整校验流程，对应CheckParams中的三个步骤：

1. 空指针检查（CheckNotNull）：对self、other、out使用OP_CHECK_NULL检查，任一为空指针则返回ACLNN_ERR_PARAM_NULLPTR（错误码 161001）。
2. 数据类型检查（CheckDtypeValid）：
   • self与other不能同时为 bool（gcd not implemented for bool）；
   • out必须是整型（IsIntegralType），否则报错；
   • 通过op::PromoteType(self, other)推导公共数据类型，推导失败（DT_UNDEFINED）报错；
   • 推导结果必须落在当前平台 dtype 支持列表内（见上文三个 support list）；
   • 推导后的类型必须能 cast 为out的类型（OP_CHECK_RESULT_DTYPE_CAST_FAILED）。
3. shape 检查（CheckShape）：
   • self、other维度不能超过 8（MAX_DIM_LEN）；
   • 通过OP_CHECK_BROADCAST_AND_INFER_SHAPE校验二者可 broadcast 并得到广播后 shape；
   • 广播后的 shape 必须与out的 shape 完全一致。

执行器构建流程

校验通过后，第一段接口会构建算子计算图。由于self、other可能不是推导类型，也可能不是连续张量，aclnn_gcd.cpp 内部按如下固定流程组织计算：

1. 计算promoteType = PromoteType(self, other)；
2. 对self调用l0op::Contiguous转连续，若其 dtype 与 promoteType 不同则调用l0op::Cast转换（l0 接口见 op_api/gcd.cpp 与 op_api/gcd.h 中的l0op::Gcd）；
3. 对other做同样的 Contiguous + Cast；
4. 调用l0op::Gcd(selfCasted, otherCasted, executor)完成核心计算；
5. 若计算结果类型与out不一致，调用l0op::Cast转换；
6. 调用l0op::ViewCopy将结果写入out（out可能是非连续张量）；
7. 通过uniqueExecutor->GetWorkspaceSize()返回所需 workspace 大小，并将 executor 通过ReleaseTo转移给调用方。

当self或other为空张量（IsEmpty()）时，接口直接返回当前 executor 的 workspace 大小并成功退出，跳过实际计算。

aclnnGcd 参数详解

第二段接口真正执行计算，各参数说明如下：

第二段接口的实现非常简洁（见 op_api/aclnn_gcd.cpp），核心是调用框架统一入口CommonOpExecutorRun(workspace, workspaceSize, executor, stream)完成计算，并打点L2_DFX_PHASE_2(aclnnGcd)用于 DFX 追踪。所有算子相关逻辑（参数校验、tiling、kernel 调度）都在第一段接口与 executor 构建阶段完成。

返回值与错误码

两个接口均返回aclnnStatus状态码，通用状态码说明可参考 docs/zh/context/aclnn_return_code.md。其中常见的错误码包括：

对于aclnnGcdGetWorkspaceSize，第一段接口完成入参校验，以下场景会报错：

排查建议：当返回 161002 时，可结合报错日志中的具体描述定位是 dtype 推导、cast 还是 shape broadcast 环节失败；当返回 561xxx 系列内部错误时（如ACLNN_ERR_INNER_OPP_PATH_NOT_FOUND、ACLNN_ERR_INNER_OPP_KERNEL_PKG_NOT_FOUND），通常与算子二进制 kernel 库未正确安装或ASCEND_OPP_PATH环境变量未配置有关。

约束说明

• 确定性计算：aclnnGcd默认采用确定性实现。即相同输入、相同运行环境下，多次执行结果可复现、不受调度抖动影响。

源码级实现原理：从图推导到 kernel 调度

形状推导与数据类型推导

在 graph 侧，op_graph/gcd_graph_infer.cpp 通过IMPL_OP(Gcd).InferDataType(InferDataType4Gcd)声明输出y的数据类型与输入x1保持一致；而 shape 推导复用框架的广播推导能力：op_host/gcd_infershape.cpp 中IMPL_OP_INFERSHAPE(Gcd).InferShape(Ops::Base::InferShape4Broadcast)，即输出 shape 为两个输入 broadcast 后的 shape——这与 aclnn 层CheckShape中"广播后 shape 必须与 out 一致"的校验逻辑一一对应。

Tiling 与调度

op_host/arch35/gcd_tiling_arch35.cpp 实现了 arch35 平台的 tiling：

• 校验x1、x2、y三种 dtype 必须一致，否则返回GRAPH_FAILED；
• 按x1的 dtype（uint8/int8/int16/int32/int64）分别实例化BroadcastBaseTiling<GcdOp::GcdCompute<T>::OpDag>，并设置tilingKey；
• 预留 32KB 的 DCACHE（DCACHE_SIZE）空间，剩余 UB 空间用于计算，PostTiling中通过SetLocalMemorySize(ubSize_ - DCACHE_SIZE)设定本地内存；
• 通过TilingPrepareForGcd获取 AIV 核数与 UB 大小等平台信息。

kernel 侧，op_kernel/gcd_apt.cpp 中定义了__global__ __aicore__ void gcd(...)入口函数，借助BroadcastSch<schMode, OpDag> sch(tiling); sch.Process(x1, x2, y);完成广播场景下的张量搬运与计算调度；op_kernel/arch35/gcd_struct.h 则声明了算子的 tiling key 模板参数（BRC_TEMP_SCH_MODE_KEY_DECL(schMode)）。

二进制与测试

算子二进制配置见 op_host/config/ascend950/gcd_binary.json，为 uint8/int8/int16/int32/int64 五种 dtype 分别声明了Gcd_UINT8、Gcd_INT8、Gcd_INT16、Gcd_INT32、Gcd_INT64五个 bin 文件，输入输出均为 ND 格式、动态 shape（-2）且format_match_mode为FormatAgnostic。仓库同时提供了 UT 测试（tests/ut/op_api/test_aclnn_gcd.cpp、tests/ut/op_host/op_api/test_aclnn_gcd.cpp）与可直接运行的示例 examples/test_aclnn_gcd.cpp。

调用示例

以下示例代码基于 math/gcd/docs/aclnnGcd.md 文档中的调用示例，演示完整的"初始化 ACL → 创建 tensor → 第一段接口获取 workspace → 申请 workspace → 第二段接口执行 → 同步并取回结果 → 释放资源"流程，具体编译与执行过程请参考 docs/zh/context/compile_and_run_sample.md：

#include <iostream> #include <vector> #include "acl/acl.h" #include "aclnnop/aclnn_gcd.h" #define CHECK_RET(cond, return_expr) \ do { \ if (!(cond)) { \ return_expr; \ } \ } while (0) #define LOG_PRINT(message, ...) \ do { \ printf(message, ##__VA_ARGS__); \ } while (0) int64_t GetShapeSize(const std::vector<int64_t>& shape) { int64_t shape_size = 1; for (auto i : shape) { shape_size *= i; } return shape_size; } int Init(int32_t deviceId, aclrtStream* stream) { // 固定写法,资源初始化 auto ret = aclInit(nullptr); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); ret = aclrtSetDevice(deviceId); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); aclFinalize(); return ret); ret = aclrtCreateStream(stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); aclrtResetDevice(deviceId); aclFinalize(); return ret); return 0; } template<typename T> int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType, aclTensor** tensor) { auto size = GetShapeSize(shape) * sizeof(T); // 调用aclrtMalloc申请device侧内存 auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上 ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret); // 计算连续tensor的strides std::vector<int64_t> strides(shape.size(), 1); for (int64_t i = shape.size() - 2; i >= 0; i--) { strides[i] = shape[i + 1] * strides[i + 1]; } // 调用aclCreateTensor接口创建aclTensor *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), *deviceAddr); return 0; } int main() { int32_t deviceId = 0; aclrtStream stream; auto ret = Init(deviceId, &stream); CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); std::vector<int64_t> selfShape = {4, 2}; std::vector<int64_t> otherShape = {4, 2}; std::vector<int64_t> outShape = {4, 2}; void* selfDeviceAddr = nullptr; void* otherDeviceAddr = nullptr; void* outDeviceAddr = nullptr; aclTensor* self = nullptr; aclTensor* other = nullptr; aclTensor* out = nullptr; std::vector<int32_t> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7}; std::vector<int32_t> otherHostData = {0, 1, 2, 3, 4, 5, 6, 7}; std::vector<int32_t> outHostData = {1, 1, 1, 1, 0, 0, 0, 0}; ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_INT32, &self); CHECK_RET(ret == ACL_SUCCESS, return ret); ret = CreateAclTensor(otherHostData, otherShape, &otherDeviceAddr, aclDataType::ACL_INT32, &other); CHECK_RET(ret == ACL_SUCCESS, return ret); ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_INT32, &out); CHECK_RET(ret == ACL_SUCCESS, return ret); uint64_t workspaceSize = 0; aclOpExecutor* executor; ret = aclnnGcdGetWorkspaceSize(self, other, out, &workspaceSize, &executor); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGcdGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); void* workspaceAddr = nullptr; if (workspaceSize > 0) { ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret); } ret = aclnnGcd(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGcd failed. ERROR: %d\n", ret); return ret); ret = aclrtSynchronizeStream(stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); auto size = GetShapeSize(outShape); std::vector<int32_t> resultData(size, 0); ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret); for (int64_t i = 0; i < size; i++) { LOG_PRINT("result[%ld] is: %d\n", i, resultData[i]); } // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改 aclDestroyTensor(self); aclDestroyTensor(other); aclDestroyTensor(out); // 7. 释放device资源 aclrtFree(selfDeviceAddr); aclrtFree(otherDeviceAddr); aclrtFree(outDeviceAddr); if (workspaceSize > 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }

示例中selfHostData与otherHostData均为{0, 1, 2, 3, 4, 5, 6, 7}，out初始值无关紧要（会被覆盖）。计算结果满足result[i] = gcd(self[i], other[i])，例如gcd(0, 0)按 kernel 语义（op_kernel/arch35/gcd_dag.h 中对零值的特判）输出 0，gcd(3, 3)输出 3，gcd(5, 5)输出 5。

使用要点提醒

• self、other、out均通过aclCreateTensor创建，strides需按连续张量规则计算；本文示例均为 ND 连续张量，实际业务中三者也可以是非连续张量（参数表标注为"√"），此时 aclnn 层会自动通过Contiguous/ViewCopy处理。
• workspaceSize为 0 时无需申请 workspace，直接传nullptr即可。
• executor由第一段接口生成，第二段接口执行后不可复用；整个流程建议与示例一致，在aclrtSynchronizeStream同步完成后再读取out数据。
• 若需要self与othershape 不一致（满足 broadcast 关系）的场景，可参考 docs/zh/context/broadcast_relationship.md 中关于广播规则的说明，输出 shape 取广播后结果。

扩展阅读

• 两段式接口规范：docs/zh/context/two_phase_api.md
• 数据类型推导规则：docs/zh/context/deduction_relationship.md
• broadcast 广播关系：docs/zh/context/broadcast_relationship.md
• aclnn 返回码说明：docs/zh/context/aclnn_return_code.md
• 编译与运行样例：docs/zh/context/compile_and_run_sample.md
• 算子实现源码：math/gcd/op_api/aclnn_gcd.cpp、math/gcd/op_host/gcd_def.cpp、math/gcd/op_kernel/arch35/gcd_dag.h
• 可运行示例：math/gcd/examples/test_aclnn_gcd.cpp

【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-math