算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载导读本文围绕 CANN ops-math 数学算子库中的 Trunc 算子系统讲解其对外暴露的aclnnTrunc与aclnnInplaceTrunc两组单算子 API包括接口功能与数学定义、产品支持范围、两段式调用流程、四个接口的函数原型与全部参数约束、第一段接口的入参校验错误码以及可直接编译运行的单算子调用示例。同时结合 math/trunc 目录下的 op_api、op_host、op_kernel 源码与单元测试深入剖析该算子从参数校验、shape 推导、tiling 计算到 AI Core kernel 执行含截断取整的位运算实现的完整链路帮助开发者理解 Trunc 算子的调用方式与底层原理并能够据此迁移到其他 aclnn 数学算子的使用与调试。功能说明Trunc 算子对输入 Tensor 的每一个元素执行 trunc 运算即将数字的小数部分截去返回其整数部分。它属于逐元素elementwise一元算子输出 Tensor 与输入 Tensor 保持相同的 shape 与数据类型。计算公式为$$ out \text{trunc}(self) $$从语义上看trunc 与向零取整等价对于正数3.7得到3对于负数-3.7得到-3这与 floor向下取整-3.7 → -4和 ceil向上取整-3.7 → -3均不相同也是本算子在量化、坐标对齐等场景下常用的原因之一。仓库根 README 对该算子的描述与此一致math/trunc/README.md 中明确对输入张量的每一个元素取整舍弃小数部分。产品支持情况Trunc 算子在当前仓库版本中的产品支持情况如下产品是否支持Ascend 950PR / Ascend 950DT支持Atlas A3 训练系列产品 / Atlas A3 推理系列产品支持Atlas A2 训练系列产品 / Atlas A2 推理系列产品支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品支持Atlas 训练系列产品支持注意不同产品在数据类型支持上存在差异详见下文 aclnnTruncGetWorkspaceSize 一节中的产品差异说明。此外math/trunc/op_host/trunc_def.cpp 的算子定义中注册了ascend950与ascend350两个 AI Core 配置项并开启了动态 shape/动态 rank 支持与上述产品支持矩阵相对应。两段式接口与调用流程aclnnTrunc与aclnnInplaceTrunc实现相同的计算功能二者的区别仅在于输出结果的存放方式aclnnTrunc非原地版本需要调用方新建一个输出张量对象out来存储计算结果输入self在计算后保持不变。aclnnInplaceTrunc原地版本无需新建输出张量对象直接在输入张量selfRef的内存中覆盖存储计算结果可节省一份输出内存适合输入不再被复用的场景。两个接口均遵循 CANN 单算子 API 的两段式调用模式详见 两段式接口先调用第一段接口aclnnTruncGetWorkspaceSize或aclnnInplaceTruncGetWorkspaceSize完成入参校验并根据计算流程推导出所需 workspace 大小与 op 执行器按照第一段接口返回的workspaceSize在 Device 侧申请 workspace 内存再调用第二段接口aclnnTrunc或aclnnInplaceTrunc传入 workspace、workspaceSize、executor 与 stream 执行实际计算。其中 workspace 是指除输入/输出之外算子在 NPU 上完成计算所需的临时内存其大小由第一段接口计算得出。第二段接口不能重复调用同一个 executor 在一次GetWorkspaceSize → 执行配对中只能执行一次。函数原型四个接口的函数原型如下aclnnStatus aclnnTruncGetWorkspaceSize( const aclTensor* self, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)aclnnStatus aclnnTrunc( void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)aclnnStatus aclnnInplaceTruncGetWorkspaceSize( aclTensor* selfRef, uint64_t* workspaceSize, aclOpExecutor** executor)aclnnStatus aclnnInplaceTrunc( void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)接口声明可在 math/trunc/op_api/aclnn_trunc.h 中查阅其中标注了接口域为aclnn_math并给出了与文档一致的参数语义说明。aclnnTruncGetWorkspaceSize参数说明参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续TensorselfaclTensor*输入公式中的 self。shape 和数据类型需要与 out 一致。FLOAT、FLOAT16、BFLOAT16ND0-8√outaclTensor*输出公式中的 out。数据类型和 shape 需要与 self 一致。FLOAT、FLOAT16、BFLOAT16ND0-8√workspaceSizeuint64_t*输出返回需要在 Device 侧申请的 workspace 大小。-----executoraclOpExecutor**输出返回 op 执行器包含了算子计算流程。-----产品差异数据类型Atlas 训练系列产品不支持BFLOAT16 数据类型。Ascend 950PR / Ascend 950DT数据类型支持FLOAT、FLOAT16、BFLOAT16、INT32、INT8、UINT8。该产品差异与 math/trunc/op_api/aclnn_trunc.cpp 中按 NPU 架构区分的 dtype 支持列表完全对应源码中ASCEND910_DTYPE_DTYPE_SUPPORT_LISTFLOAT、FLOAT16、ASCEND910B_DTYPE_DTYPE_SUPPORT_LISTFLOAT、FLOAT16、BF16与ARCH3510_DTYPE_DTYPE_SUPPORT_LISTFLOAT、FLOAT16、BF16、INT8、UINT8、INT32会依据当前 NPU 架构GetCurrentPlatformInfo().GetCurNpuArch()动态选择。而 math/trunc/op_host/trunc_def.cpp 的 OpDef 中注册的数据类型全集为DT_BF16、DT_FLOAT16、DT_FLOAT、DT_INT32、DT_INT8、DT_UINT8格式均为FORMAT_ND——这解释了不同硬件上可用数据类型范围不同的原因。返回值与入参校验返回值为aclnnStatus状态码具体参见 aclnn返回码。第一段接口完成入参校验出现如下场景时报错返回值错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的 self 或 out 是空指针。ACLNN_ERR_PARAM_INVALID161002self 和 out 的数据类型不在支持的范围之内。ACLNN_ERR_PARAM_INVALID161002self 和 out 的数据类型不一致。ACLNN_ERR_PARAM_INVALID161002self 和 out 的维度大于 8。ACLNN_ERR_PARAM_INVALID161002self 和 out 的 shape 不一致。上述校验逻辑与源码实现一一对应math/trunc/op_api/aclnn_trunc.cpp 中的CheckParamsTrunc依次执行空指针检查CheckNotNull2Tensor、数据类型合法性及一致性检查CheckDtypeValid、shape 一致性检查CheckSameShape1In1Out任一不满足即返回对应的ACLNN_ERR_PARAM_NULLPTR或ACLNN_ERR_PARAM_INVALID。单元测试 math/trunc/tests/ut/op_api/test_aclnn_trunc.cpp 中对空指针 self/out、非法 dtype、dtype 不一致、shape 不一致等异常场景均有断言覆盖。aclnnTrunc参数说明参数名输入/输出描述workspace输入在 Device 侧申请的 workspace 内存地址。workspaceSize输入在 Device 侧申请的 workspace 大小由第一段接口 aclnnTruncGetWorkspaceSize 获取。executor输入op 执行器包含了算子计算流程。stream输入指定执行任务的 Stream。返回值返回aclnnStatus状态码具体参见 aclnn返回码。aclnnInplaceTruncGetWorkspaceSize参数说明参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续TensorselfRefaclTensor*输入|输出输入|输出张量。-FLOAT、FLOAT16、BFLOAT16ND0-8√workspaceSizeuint64_t*输出返回需要在 Device 侧申请的 workspace 大小。-----executoraclOpExecutor**输出返回 op 执行器包含了算子计算流程。-----原地版本只有selfRef一个张量参数它同时承担输入与输出角色因此无需也不允许单独构造 out。从源码看math/trunc/op_api/aclnn_trunc.cpp 中aclnnInplaceTruncGetWorkspaceSize直接执行auto out const_castaclTensor*(selfRef);将自身作为输出随后复用与普通版本完全相同的ExecTruncGetWorkspaceSize计算流程。返回值与入参校验返回aclnnStatus状态码具体参见 aclnn返回码。第一段接口完成入参校验出现如下场景时报错返回值错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的 selfRef 是空指针。ACLNN_ERR_PARAM_INVALID161002selfRef 数据类型和数据格式不在支持的范围之内。ACLNN_ERR_PARAM_INVALID161002selfRef 的维度大于 8。aclnnInplaceTrunc参数说明参数名输入/输出描述workspace输入在 Device 侧申请的 workspace 内存地址。workspaceSize输入在 Device 侧申请的 workspace 大小由第一段接口 aclnnInplaceTruncGetWorkspaceSize 获取。executor输入op 执行器包含了算子计算流程。stream输入指定执行任务的 Stream。返回值返回aclnnStatus状态码具体参见 aclnn返回码。约束说明确定性计算aclnnTrunc与aclnnInplaceTrunc默认即为确定性实现即相同输入在多次执行下会得到一致的计算结果。这一点对依赖可复现结果的训练与推理场景非常重要如需了解 CANN 对确定性计算的整体支持策略可参考 确定性计算。调用示例以下示例完整演示了在同一进程中依次以非原地与原地两种方式调用 Trunc 算子的全过程包含 ACL 初始化、输入输出 Tensor 构造、两段式接口调用、结果回拷与资源释放。示例代码仅作参考具体编译与运行步骤请参考 编译与运行样例仓库内对应的可运行示例位于 math/trunc/examples/test_aclnn_trunc.cpp。#include iostream #include vector #include acl/acl.h #include aclnnop/aclnn_trunc.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 shapeSize 1; for (auto i : shape) { shapeSize * i; } return shapeSize; } 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); return ret); ret aclrtCreateStream(stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtCreateStream failed. ERROR: %d\n, ret); return ret); return 0; } template typename 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() { // 1.固定写法device/stream初始化参考acl API手册 // 根据自己的实际device填写deviceId int32_t deviceId 0; aclrtStream stream; auto ret Init(deviceId, stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(Init acl failed. ERROR: %d\n, ret); return ret); // 2.构造输入与输出需要根据API的接口自定义构造 std::vectorint64_t selfShape {4, 2}; std::vectorint64_t outShape {4, 2}; void* selfDeviceAddr nullptr; void* outDeviceAddr nullptr; aclTensor* self nullptr; aclTensor* out nullptr; std::vectorfloat selfHostData {0, 1, 2, 3, 4, 5, 6, 7}; std::vectorfloat outHostData {0, 0, 0, 0, 0, 0, 0, 0}; // 创建self aclTensor ret CreateAclTensor(selfHostData, selfShape, selfDeviceAddr, aclDataType::ACL_FLOAT, self); CHECK_RET(ret ACL_SUCCESS, return ret); // 创建out aclTensor ret CreateAclTensor(outHostData, outShape, outDeviceAddr, aclDataType::ACL_FLOAT, out); CHECK_RET(ret ACL_SUCCESS, return ret); // 3.调用CANN算子库API需要修改为具体的API名称 uint64_t workspaceSize 0; aclOpExecutor* executor; // 调用aclnnTrunc第一段接口 ret aclnnTruncGetWorkspaceSize(self, out, workspaceSize, executor); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnTruncGetWorkspaceSize failed. ERROR: %d\n, ret); return ret); // 根据第一段接口计算出的workspaceSize申请device内存 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); } // 调用aclnnTrunc第二段接口 ret aclnnTrunc(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnTrunc failed. ERROR: %d\n, ret); return ret); uint64_t inplaceWorkspaceSize 0; aclOpExecutor* inplaceExecutor; // 调用aclnnInplaceTrunc第一段接口 ret aclnnInplaceTruncGetWorkspaceSize(self, inplaceWorkspaceSize, inplaceExecutor); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnInplaceTruncGetWorkspaceSize failed. ERROR: %d\n, ret); return ret); // 根据第一段接口计算出的workspaceSize申请device内存 void* inplaceWorkspaceAddr nullptr; if (inplaceWorkspaceSize 0) { ret aclrtMalloc(inplaceWorkspaceAddr, inplaceWorkspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(allocate workspace failed. ERROR: %d\n, ret); return ret); } // 调用aclnnInplaceTrunc第二段接口 ret aclnnInplaceTrunc(inplaceWorkspaceAddr, inplaceWorkspaceSize, inplaceExecutor, stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnInplaceTrunc failed. ERROR: %d\n, ret); return ret); // 4.固定写法同步等待任务执行结束 ret aclrtSynchronizeStream(stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclrtSynchronizeStream failed. ERROR: %d\n, ret); return ret); // 5.获取输出的值将device侧内存上的结果拷贝至host侧需要根据具体API的接口定义修改 auto size GetShapeSize(outShape); std::vectorfloat 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: %f\n, i, resultData[i]); } auto inplaceSize GetShapeSize(selfShape); std::vectorfloat inplaceResultData(inplaceSize, 0); ret aclrtMemcpy(inplaceResultData.data(), inplaceResultData.size() * sizeof(inplaceResultData[0]), selfDeviceAddr, inplaceSize * sizeof(float), 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 inplaceSize; i) { LOG_PRINT(inplaceResult[%ld] is: %f\n, i, inplaceResultData[i]); } // 6.释放aclTensor和aclScalar需要根据具体API的接口定义修改 aclDestroyTensor(self); aclDestroyTensor(out); // 7.释放device资源需要根据具体API的接口定义修改 aclrtFree(selfDeviceAddr); aclrtFree(outDeviceAddr); if (workspaceSize 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }对示例的几点补充说明示例中selfHostData {0, 1, 2, 3, 4, 5, 6, 7}均为整数仅用于演示流程实际使用时可将输入改为带小数的浮点数据如3.7、-3.7即可直观验证 trunc 的向零取整语义。原地版本的结果直接从selfDeviceAddr回拷inplaceResultData分支而非从单独的 out 内存回拷这是原地调用与普通调用的核心差异。workspace 的申请采用ACL_MEM_MALLOC_HUGE_FIRST策略当workspaceSize 0时无需申请第二段接口可直接传入空指针。示例中aclTensor的创建使用了连续 tensor 的 strides 计算方式strides[i] shape[i1] * strides[i1]而文档参数表显示 Trunc 支持非连续 Tensor第一段接口内部会先通过l0op::Contiguous将输入转为连续张量再执行 kernel输出侧则通过l0op::ViewCopy将结果写回可能非连续的 out详见下一节源码剖析。源码级剖析从 API 到 AI Core Kernel 的完整链路为了更深入地理解接口行为下面结合 math/trunc 目录源码梳理 Trunc 算子的完整实现链路。1. 第一段接口参数校验 → 连续性转换 → 计算图构建math/trunc/op_api/aclnn_trunc.cpp 中的ExecTruncGetWorkspaceSize是两段式接口的核心其执行序列为CREATE_EXECUTOR()创建 op 执行器aclOpExecutorCheckParamsTrunc完成空指针、dtype 支持列表与一致性、shape 一致性校验对空 tensor 输入self-IsEmpty()做提前短路处理直接返回 workspaceSize调用l0op::Contiguous(self, executor)将可能非连续的输入转为连续 tensor调用l0op::Trunc(selfContiguous, executor)构建截断计算节点调用l0op::ViewCopy(truncOpOut, out, executor)将结果写入用户提供的 out兼容非连续 outuniqueExecutor-GetWorkspaceSize()汇总全流程所需 workspace 大小并返回。其中l0op::Trunc与l0op::InplaceTrunc的算子级实现位于 math/trunc/op_api/trunc.cpp两者均通过TruncInferShape直接将输出 shape 设为输入 shape逐元素算子特性并通过ADD_TO_LAUNCHER_LIST_AICORE(Trunc, OP_INPUT(self), OP_OUTPUT(out))注册到 AICore 启动列表唯一区别在于InplaceTrunc使用const_castaclTensor*(self)使输出与输入共用同一内存地址从而实现在地in-place语义。2. 算子定义与 shape 推导动态 shape/rank 支持math/trunc/op_host/trunc_def.cpp 中Trunc算子定义了单个输入input_x与单个输出output_y数据类型支持DT_BF16 / DT_FLOAT16 / DT_FLOAT / DT_INT32 / DT_INT8 / DT_UINT8格式为NDAI Core 配置开启了DynamicCompileStaticFlag(true)、DynamicShapeSupportFlag(true)、DynamicRankSupportFlag(true)对应文档中维度 0-8、支持动态 shape的约束。math/trunc/op_host/trunc_infershape.cpp 通过IMPL_OP_INFERSHAPE(Trunc).InferShape(InferShape4Elewise)注册逐元素算子的 shape 推导规则即输出 shape 与输入完全一致。3. Tilingworkspace 分配与 Kernel 调度参数math/trunc/op_host/arch35/trunc_tiling_arch35.cpp 是 arch35 架构Ascend 950/350 系列的 tiling 实现从中可以看到固定申请ASCEND_WORKSPACE 16 * 1024 * 102416MB的 workspaceCalcInputDtype校验输入 dtype 必须在FLOAT16 / BF16 / FLOAT / INT8 / UINT8 / INT32范围内与文档及 OpDef 一致CheckShape强制要求输入与输出 shape 一致否则报错通过GET_TPL_TILING_KEY(tiling-baseTiling.scheMode, dType)生成 tilingKey并设置 block 维度供 kernel 侧按模板参数分派不同数据类型实现。4. Kernel截断取整的位运算实现math/trunc/op_kernel/trunc_apt.cpp 定义了__global__ __aicore__ void trunc(...)kernel依据编译期 dtype 模板参数TPL_FP16 / TPL_BF16 / TPL_FP32 / TPL_INT8 / TPL_UINT8 / TPL_INT32实例化不同的ElementwiseSch调度。值得关注的是浮点类型 trunc 的底层实现 math/trunc/op_kernel/arch35/trunc_dag.h它并不是简单地调用取整函数而是通过符号位保留 截断转换的位运算组合完成Reg::TruncateT, RoundMode::CAST_TRUNC, ...将浮点数据按截断向零舍入模式强制转换为整数得到小数部分被舍弃的整数值Reg::Duplicate(vregOutInt, UINT32_SIGN / UINT16_SIGN, mask)生成仅保留符号位的掩码对原始输入与掩码做And提取原符号位再与截断结果做Or将符号位写回对 FP32 使用 32 位符号掩码0x80000000对 FP16/BF16 使用 16 位掩码0x8000从而正确处理负数场景保证向零取整语义。对于 INT8 / UINT8 / INT32 等整数类型trunc_dag.h 中TruncDAGInt的 DAG 仅包含 CopyIn → CopyOut即整数数据本身没有小数部分直接透传即可。5. Kernel 二进制清单与单算子编译配置math/trunc/op_host/config/ascend950/trunc_binary.json 与 math/trunc/op_host/config/ascend350/trunc_binary.json 是单算子 Kernel 二进制包bin的清单文件其中op_type为Trunc针对bfloat16 / float16 / float32 / int32 / int8 / uint8六种 dtype 分别登记了对应的bin_filename如Trunc_1c4543fdfe7149b6b0717bd88f061562shape 字段为-2表示支持动态 shapeformat_match_mode为FormatAgnostic格式无关。这也从编译产物层面印证了上述六个 dtype 是当前仓库 Trunc 算子的完整支持集合。6. 测试验证op_api 单元测试math/trunc/tests/ut/op_api/test_aclnn_trunc.cpp 覆盖了 FLOAT16/FLOAT/BF16 三种 dtype、1D 至 8D 的多种 shape、9D非法维度返回ACLNN_ERR_PARAM_INVALID、空 tensor、空指针 self/out返回ACLNN_ERR_PARAM_NULLPTR、dtype 不合法与不一致、shape 不一致以及非连续 Tensor等场景math/trunc/tests/ut/op_api/test_aclnn_inplacetrunc.cpp 则针对原地版本做同等覆盖。op_host 单元测试math/trunc/tests/ut/op_host/test_trunc_infershape.cpp 与 math/trunc/tests/ut/op_host/arch35/test_trunc_tiling.cpp 分别验证 shape 推导与 tiling 逻辑。ST 场景用例math/trunc/tests/st/arch35/ttk_aclnn_trunc_st.csv 与 math/trunc/tests/st/arch35/ttk_kernel_trunc_st.csv 提供端到端与 kernel 级的用例清单测试数据生成脚本位于 math/trunc/tests/assets/golden.py。总结aclnnTrunc与aclnnInplaceTrunc是 CANN ops-math 中 Trunc 算子的两组两段式单算子 API功能等价、仅在是否复用输入内存上有所区别。使用时需严格遵循先GetWorkspaceSize获取 workspace 大小并完成入参校验再申请 Device 内存、调用执行接口的固定流程输入输出需满足 ND 格式、0-8 维、数据类型一致且位于产品支持范围等约束。通过源码可以看到Trunc 算子在 op_api 层通过 Contiguous/ViewCopy 完成非连续 Tensor 的适配在 kernel 层以截断转换 符号位回写的位运算实现了浮点向零取整并针对不同 NPU 架构提供差异化的 dtype 支持与动态 shape 能力。本文涉及的源码与测试文件均可直接在仓库 math/trunc 目录下进一步查阅相关通用概念可参考 两段式接口、aclnn返回码、编译与运行样例 与 确定性计算。赞分享算子库人工智能CANN【免费下载链接】ops-math本项目是CANN提供的数学类基础计算算子库实现网络在NPU上加速计算。项目地址https://gitcode.com/cann/ops-math点击查看免费下载相关推荐CANN ops-math 算子接口指南aclnnClampMin 两段式 API 调用与实现解析CANN ops math 算子接口指南aclnnClampMin 两段式 API 调用与实现解析 aclnnClampMin 是 CANN ops math算子库人工智能CANNCANN ops-math 算子 aclnnCastV3 接口使用指南CastV3 两段式调用与源码实现解析CANN ops math 算子 aclnnCastV3 接口使用指南CastV3 两段式调用与源码实现解析 aclnnCastV3 是 CANN ops m算子库人工智能CANNCANN ops-math 中 aclnnSin / aclnnInplaceSin 算子接口使用指南两段式 API 详解与源码级调用链分析CANN ops math 中 aclnnSin / aclnnInplaceSin 算子接口使用指南两段式 API 详解与源码级调用链分析 导读 aclnn算子库人工智能CANN创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考