CANN ops-math 算子指南aclnnGcd 两段式接口详解与 NPU 最大公约数计算实战【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-math导读aclnnGcd是 CANN ops-math 数学算子库中用于在 NPUAscend 芯片上按元素element-wise计算两个整数张量最大公约数Greatest Common DivisorGCD的单算子 API。本文以 math/gcd/docs/aclnnGcd.md 为核心结合仓库中该算子的 op_api、op_host、op_kernel 与测试代码完整讲解其功能语义、产品支持范围、两段式接口aclnnGcdGetWorkspaceSizeaclnnGcd的参数与错误码、源码级实现原理以及可直接运行的调用示例。读完本文你将掌握在 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、动态 rankDynamicRankSupportFlag(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 中声明的产品支持矩阵产品系列支持情况Ascend 950PR / Ascend 950DT支持Atlas A3 训练系列产品 / Atlas A3 推理系列产品支持Atlas A2 训练系列产品 / Atlas A2 推理系列产品支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品支持Atlas 训练系列产品支持说明产品支持情况以文档声明为准实际使用前请结合 CANN 版本配套关系确认目标芯片是否在支持范围内。两段式接口与函数原型aclnnGcd属于 CANN 单算子 APIaclnn 接口遵循两段式接口调用规范详见 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 大小计算与执行器构建各参数说明如下参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续TensorselfaclTensor*输入表示待转换的目标张量。数据类型与 other 的数据类型需满足数据类型推导规则。shape 与 other 的 shape 满足 broadcast 关系。UINT8、INT8、INT16、INT32、INT64ND1-8√otheraclTensor*输入表示待转换的目标张量。数据类型与 self 的数据类型需满足数据类型推导规则。shape 需要与 self 满足 broadcast 关系。UINT8、INT8、INT16、INT32、INT64ND1-8√outaclTensor*输出self 和 other 求最大公约数的结果。数据类型是 self 与 other 推导之后可转换的数据类型。shape 需要与 self 和 other 做 broadcast 后的 shape 一致。UINT8、INT8、UINT16、INT16、INT32、UINT32、INT64、UINT64ND1-8√workspaceSizeuint64_t*输出返回需要在 Device 侧申请的 workspace 大小。-----executoraclOpExecutor**输出返回 op 执行器包含了算子计算流程。-----平台相关的数据类型差异在上述输入 dtype 基础上存在平台差异Atlas A3 训练系列产品 / Atlas A3 推理系列产品不支持 UINT8、INT8 数据类型。这一差异在 op_api/aclnn_gcd.cpp 的源码中有更细粒度的体现。实现中按 SoC 版本维护了不同的 dtype 支持列表static const std::initializer_listop::DataType ASCEND910_DTYPE_SUPPORT_LIST {op::DataType::DT_INT32}; static const std::initializer_listop::DataType ASCEND910B_DTYPE_SUPPORT_LIST { op::DataType::DT_INT32, op::DataType::DT_INT16, op::DataType::DT_INT64}; static const std::initializer_listop::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仅支持 INT32Atlas A2/A3 系列910B 等支持 INT32、INT16、INT64Ascend 950RegBase 平台支持 UINT8、INT8、INT16、INT32、INT64 五种整数类型。头文件 op_api/aclnn_gcd.h 中亦注明推导后数据类型支持 INT32、INT16(Ascend910B)。因此在实际开发中应优先采用 INT32 以保证跨平台通用性。参数校验的实现逻辑从 op_api/aclnn_gcd.cpp 可以看到第一段接口的完整校验流程对应CheckParams中的三个步骤空指针检查CheckNotNull对self、other、out使用OP_CHECK_NULL检查任一为空指针则返回ACLNN_ERR_PARAM_NULLPTR错误码 161001。数据类型检查CheckDtypeValidself与other不能同时为 boolgcd not implemented for boolout必须是整型IsIntegralType否则报错通过op::PromoteType(self, other)推导公共数据类型推导失败DT_UNDEFINED报错推导结果必须落在当前平台 dtype 支持列表内见上文三个 support list推导后的类型必须能 cast 为out的类型OP_CHECK_RESULT_DTYPE_CAST_FAILED。shape 检查CheckShapeself、other维度不能超过 8MAX_DIM_LEN通过OP_CHECK_BROADCAST_AND_INFER_SHAPE校验二者可 broadcast 并得到广播后 shape广播后的 shape 必须与out的 shape 完全一致。执行器构建流程校验通过后第一段接口会构建算子计算图。由于self、other可能不是推导类型也可能不是连续张量aclnn_gcd.cpp 内部按如下固定流程组织计算计算promoteType PromoteType(self, other)对self调用l0op::Contiguous转连续若其 dtype 与 promoteType 不同则调用l0op::Cast转换l0 接口见 op_api/gcd.cpp 与 op_api/gcd.h 中的l0op::Gcd对other做同样的 Contiguous Cast调用l0op::Gcd(selfCasted, otherCasted, executor)完成核心计算若计算结果类型与out不一致调用l0op::Cast转换调用l0op::ViewCopy将结果写入outout可能是非连续张量通过uniqueExecutor-GetWorkspaceSize()返回所需 workspace 大小并将 executor 通过ReleaseTo转移给调用方。当self或other为空张量IsEmpty()时接口直接返回当前 executor 的 workspace 大小并成功退出跳过实际计算。aclnnGcd 参数详解第二段接口真正执行计算各参数说明如下参数名输入/输出描述workspace输入在 Device 侧申请的 workspace 内存地址。workspaceSize输入在 Device 侧申请的 workspace 大小由第一段接口 aclnnGcdGetWorkspaceSize 获取。executor输入op 执行器包含了算子计算流程。stream输入指定执行任务的 Stream。第二段接口的实现非常简洁见 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。其中常见的错误码包括返回值错误码描述ACLNN_SUCCESS0成功。ACLNN_ERR_PARAM_NULLPTR161001参数校验错误参数中存在非法的 nullptr。ACLNN_ERR_PARAM_INVALID161002参数校验错误如输入的两个数据类型不满足输入类型推导关系。ACLNN_ERR_RUNTIME_ERROR361001API 内部调用 npu runtime 的接口异常。ACLNN_ERR_INNER_XXX561xxxAPI 内部发生异常。对于aclnnGcdGetWorkspaceSize第一段接口完成入参校验以下场景会报错返回值错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的 self、other、out 是空指针。ACLNN_ERR_PARAM_INVALID161002self、other 推导后的数据类型不在支持范围之内或 out 的数据类型不在支持的范围之内。ACLNN_ERR_PARAM_INVALID161002self 和 other 的数据类型不满足数据类型推导规则。ACLNN_ERR_PARAM_INVALID161002self 和 other 进行数据类型转换后的数据类型无法转换为指定输出 out 的类型。ACLNN_ERR_PARAM_INVALID161002self 和 other 的 shape 不满足 broadcast 规则。ACLNN_ERR_PARAM_INVALID161002self 和 other 进行 broadcast 后的 shape 与 out 不一样。排查建议当返回 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的 dtypeuint8/int8/int16/int32/int64分别实例化BroadcastBaseTilingGcdOp::GcdComputeT::OpDag并设置tilingKey预留 32KB 的 DCACHEDCACHE_SIZE空间剩余 UB 空间用于计算PostTiling中通过SetLocalMemorySize(ubSize_ - DCACHE_SIZE)设定本地内存通过TilingPrepareForGcd获取 AIV 核数与 UB 大小等平台信息。kernel 侧op_kernel/gcd_apt.cpp 中定义了__global__ __aicore__ void gcd(...)入口函数借助BroadcastSchschMode, 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::vectorint64_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; } templatetypename T int CreateAclTensor(const std::vectorT hostData, const std::vectorint64_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::vectorint64_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::vectorint64_t selfShape {4, 2}; std::vectorint64_t otherShape {4, 2}; std::vectorint64_t outShape {4, 2}; void* selfDeviceAddr nullptr; void* otherDeviceAddr nullptr; void* outDeviceAddr nullptr; aclTensor* self nullptr; aclTensor* other nullptr; aclTensor* out nullptr; std::vectorint32_t selfHostData {0, 1, 2, 3, 4, 5, 6, 7}; std::vectorint32_t otherHostData {0, 1, 2, 3, 4, 5, 6, 7}; std::vectorint32_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::vectorint32_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 中对零值的特判输出 0gcd(3, 3)输出 3gcd(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.mdbroadcast 广播关系docs/zh/context/broadcast_relationship.mdaclnn 返回码说明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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考