CANN ops-transformer 算子 MoeGatingTopKSoftmaxV2：aclnnMoeGatingTopKSoftmaxV2 接口详解与源码剖析-北京尧图网络

CANN ops-transformer 算子 MoeGatingTopKSoftmaxV2aclnnMoeGatingTopKSoftmaxV2 接口详解与源码剖析【免费下载链接】ops-transformer本项目是CANN提供的transformer类大模型算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-transformer本技术指南聚焦于 CANN ops-transformer 算子库中 MoEMixture of Experts路由计算的核心算子MoeGatingTopKSoftmaxV2及其对外 aclnn 接口aclnnMoeGatingTopKSoftmaxV2完整讲解其在混合专家大模型场景下的功能语义、两段式调用范式、全部参数约束与错误码并结合仓库源码算子定义、Infershape、Tiling、Kernel 分发与 golden 校验剖析其底层实现原理。读完本文你将掌握该算子在 Ascend NPU 上的正确调用方式、renorm两种计算模式的差异以及如何读懂并调试其配套的测试用例。一、功能概述MoE 路由中的 Gating TopK Softmax在混合专家Mixture of Experts, MoE大模型结构中Router路由器为每个 token 计算其对所有专家的选择概率再通过 TopK 挑选出得分最高的若干专家并完成权重归一化。MoeGatingTopKSoftmaxV2正是将Softmax与TopK两个操作按可配置顺序融合为单算子实现避免中间张量落地 GMGlobal Memory从而减少访存开销、提升路由阶段在 NPU 上的计算效率。其核心功能可概括为若renorm0先对输入x沿最后一维做Softmax再对 Softmax 结果做TopK输出 TopK 的值与索引若renorm1先对输入x沿最后一维做TopK再对取出的 TopK 值做Softmax重新归一化即 renorm 的含义输出归一化后的 TopK 值与索引。同时算子支持可选的finishedOptional输入当某一行finished为True时该行不参与实际运算expertIdxOut直接填入num_expert值即x最后一个轴的大小用于在推理阶段标识已结束的序列/行。二、计算公式设输入x的形状为(..., N)N即专家数属性k为每个 token 选中的专家数Softmax 均沿最后一维axis-1执行则两种模式的计算过程如下。模式一renorm 0先 Softmax 再 TopKsoftmaxResultOutOptional softmax(x, axis-1) yOut, expertIdxOut topK(softmaxResultOutOptional, kk)模式二renorm 1先 TopK 再 SoftmaxtopkOut, expertIdxOut topK(x, kk) yOut softmax(topkOut, axis-1)两式中的yOut含义不同renorm0时是 Softmax 之后的 TopK 值renorm1时是先取 TopK 再做 Softmax 归一化得到的值。expertIdxOut则是 TopK 返回的索引即被选中的专家序号。对于finishedTrue的行expertIdxOut直接填充为num_expert即x最后一维大小。上述公式与 golden 参考实现一致见 tests/assets/golden.py其_softmax_numpy/_softmax_torch采用先减最大值再求指数、最后归一化的数值稳定 Softmax 写法且 float16 输入会先提升到 float32 计算这与算子在x为低精度输入时的内部精度处理策略相互印证。三、产品支持情况根据官方文档及算子配置该算子在以下产品上的支持情况如下产品是否支持Ascend 950PR / Ascend 950DT支持Atlas A3 训练系列产品 / Atlas A3 推理系列产品支持Atlas A2 训练系列产品 / Atlas A2 推理系列产品支持Atlas 200I/500 A2 推理产品不支持Atlas 推理系列产品不支持Atlas 训练系列产品不支持从源码看支持范围与 moe_gating_top_k_softmax_v2_def.cpp 中的 AICore 配置一一对应算子在ascend910b与ascend910_93对应 Atlas A2/A3 系列注册为动态编译 kernelmoe_gating_top_k_softmax_v2在ascend950注册为 RegBase 模式的 kernelmoe_gating_top_k_softmax_v2_apt。四、函数原型与两段式接口aclnnMoeGatingTopKSoftmaxV2属于 CANN 的两段式Two-Phaseaclnn 接口必须先调用第一段GetWorkspaceSize接口获取计算所需的 workspace 大小以及包含算子计算流程的执行器再调用第二段执行接口完成计算。两段式接口的通用约定可参考两段式接口说明。aclnnStatus aclnnMoeGatingTopKSoftmaxV2GetWorkspaceSize( const aclTensor *x, const aclTensor *finishedOptional, int64_t k, int64_t renorm, bool outputSoftmaxResultFlag, const aclTensor *yOut, const aclTensor *expertIdxOut, const aclTensor *softmaxResultOutOptional, uint64_t *workspaceSize, aclOpExecutor **executor)aclnnStatus aclnnMoeGatingTopKSoftmaxV2( void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)头文件为aclnnop/aclnn_moe_gating_top_k_softmax_v2.h调用方需要在代码中同时包含acl/acl.h与上述算子头文件。仓库中提供了可直接参考的完整可编译样例 examples/test_aclnn_moe_gating_top_k_softmax_v2.cpp。五、aclnnMoeGatingTopKSoftmaxV2GetWorkspaceSize 参数详解第一段接口负责入参校验与 workspace 大小计算各参数说明如下。参数名输入/输出描述使用说明数据类型数据格式维度(shape)非连续Tensorx输入公式中的 x待计算的输入要求是 2D/3D Tensor每一维大小应不大于 int32 最大值 2147483647FLOAT16、BFLOAT16、FLOAT32ND2-3√finishedOptional输入公式中的 finished表示该行是否参与运算shape 为 x_shape[:-1]BOOLND1-2√k输入TopK 的 k 值即专家数0 k x 的 -1 轴大小且 k 不大于 1024INT64---renorm输入Softmax 与 TopK 的执行顺序只支持取值 0 和 10 表示先 Softmax 再 TopK1 顺序相反INT64---outputSoftmaxResultFlag输入是否输出 Softmax 结果renorm0 时true 表示输出 Softmax 结果false 不输出renorm1 时该参数不生效BOOL---yOut输出对 x 做 Softmax 后取的 TopK 值非 -1 轴与 x 对应轴大小一致-1 轴大小等于 k与输入 x 一致ND2-3xexpertIdxOut输出TopK 的索引结果即对应的专家序号shape 与 yOut 一致INT32ND2-3xsoftmaxResultOutOptional输出计算过程中 Softmax 的结果shape 与 x 一致FLOAT32ND2-3xworkspaceSize输出需要在 Device 侧申请的 workspace 大小-----executor输出op 执行器包含算子计算流程-----关于几个关键参数的补充说明k与outputSoftmaxResultFlag在算子定义中的对应关系底层 IR 属性名为k必选Int 型、renorm可选默认 0、output_softmax_result_flag可选默认 false见 moe_gating_top_k_softmax_v2_def.cpp。在 binary.json 中同样声明了这三个属性及其默认值且k的默认值为 0、renorm默认值为 0、output_softmax_result_flag默认值为 false。softmaxResultOutOptional的数据类型固定为 FLOAT32即使输入x为 FLOAT16/BFLOAT16Softmax 中间结果也以 FLOAT32 输出这与 Infershape 中InferDataType强制将输出 2 设置为DT_FLOAT的逻辑一致见 moe_gating_top_k_softmax_v2_infershape.cpp也体现了算子内部以高精度累加保证 Softmax 数值稳定性的设计取向。finishedOptional为可选输入在 aclnn 示例中直接以nullptr传入表示所有行都参与计算。若传入其 shape 必须等于x_shape[:-1]即对 2D 输入为 1D 张量、对 3D 输入为 2D 张量逐行标记是否 finished。注意算子 IR 中该输入名称为finishedaclnn 层参数名为finishedOptional。输出 shape 推导规则yOut与expertIdxOut的 shape 为x_shape[:-1] [k]softmaxResultOutOptional的 shape 与x一致当renorm0且outputSoftmaxResultFlagfalse或不输出 Softmax 结果时softmaxResultOutOptional的 shape 被推导为[0]空张量。以上规则可对照 moe_gating_top_k_softmax_v2_infershape.cpp 中setOutShape的实现验证。非连续 Tensorx、finishedOptional均支持非连续输入表格中√yOut、expertIdxOut、softmaxResultOutOptional的非连续 Tensor列标记为x表示其连续性要求跟随输入x的连续性。非连续张量的处理机制可参考非连续Tensor说明。六、返回值与错误码两个接口均返回aclnnStatus状态码通用说明见 aclnn返回码。第一段接口在完成入参校验时可能出现以下报错场景返回值错误码描述ACLNN_ERR_PARAM_NULLPTR161001传入的必选输入、必选输出或必选属性为空指针ACLNN_ERR_PARAM_INVALID161002输入和输出的数据类型或数据格式不在支持范围内ACLNN_ERR_INNER_TILING_ERROR561002多个输入 tensor 之间的 shape 信息不匹配ACLNN_ERR_INNER_TILING_ERROR561002输入属性与输入 tensor 之间的 shape 信息不匹配ACLNN_ERR_INNER_TILING_ERROR561002x 的 shape 维度不为 2 或 3ACLNN_ERR_INNER_TILING_ERROR561002x 与 finishedOptional 的 shape 不匹配ACLNN_ERR_INNER_TILING_ERROR561002k 的值小于 0 或大于 x 的 -1 轴大小ACLNN_ERR_INNER_TILING_ERROR561002k 的值大于 1024ACLNN_ERR_INNER_TILING_ERROR561002renorm 的值不是 0 或 1ACLNN_ERR_INNER_TILING_ERROR561002softmaxResultOutOptional 的数据类型不在支持范围内这些校验约束在源码中均有对应实现k 0 || k MAX_K || k 最后一维的检查见 moe_gating_top_k_softmax_v2_infershape.cpp其中MAX_K 1024定义为文件级常量renorm取值范围检查见同一文件 L100-L106Tiling 阶段还会进一步校验输入输出数据类型、finished的 dtype 与 shape、y与expertIdx的 dtype 等见 moe_gating_top_k_softmax_v2_tiling_base.cpp。七、aclnnMoeGatingTopKSoftmaxV2 参数说明第二段接口负责在指定 Stream 上执行算子计算参数名输入/输出描述workspace输入在 Device 侧申请的 workspace 内存地址workspaceSize输入在 Device 侧申请的 workspace 大小由第一段接口 aclnnMoeGatingTopKSoftmaxV2GetWorkspaceSize 获取executor输入op 执行器包含算子计算流程stream输入指定执行任务的 Stream第二段接口返回aclnnStatus状态码含义同样参见 aclnn返回码。八、约束说明确定性计算aclnnMoeGatingTopKSoftmaxV2默认采用确定性实现即相同输入在多次运行中产生一致结果。确定性计算的通用说明可参考确定性计算。k上限k的值不大于 1024源码常量MAX_K。renorm取值只支持 0 和 1。维度上限x与finishedOptional每一维大小不大于 int32 最大值 2147483647。输入维度x仅支持 2D/3D ND 格式张量finishedOptional维度为x的 rank 减 1。类型约束x支持 FLOAT16/BFLOAT16/FLOAT32ND 格式yOut与x同类型expertIdxOut固定 INT32softmaxResultOutOptional固定 FLOAT32finishedOptional固定 BOOL。上述约束在算子定义中通过DataType/Format声明并在 Tiling 阶段逐项复核。九、调用示例以下示例代码来自 examples/test_aclnn_moe_gating_top_k_softmax_v2.cpp演示了x为{3, 4}3 行、4 个专家、k2、renorm0、输出 Softmax 中间结果的完整调用流程。具体编译与执行方法请参考编译与运行样例。#include acl/acl.h #include aclnnop/aclnn_moe_gating_top_k_softmax_v2.h #include iostream #include vector #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); 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初始化参考AscendCL对外接口列表 // 根据自己的实际device填写deviceId int32_t deviceId 0; aclrtStream stream; auto ret Init(deviceId, stream); // check根据自己的需要处理 CHECK_RET(ret 0, LOG_PRINT(Init acl failed. ERROR: %d\n, ret); return ret); // 2. 构造输入与输出需要根据API的接口自定义构造 std::vectorint64_t inputShape {3, 4}; std::vectorint64_t outShape {3, 2}; std::vectorint64_t expertIdOutShape {3, 2}; std::vectorint64_t softmaxResultOutOptionalShape {3, 4}; void* inputAddr nullptr; void* outAddr nullptr; void* expertIdOutAddr nullptr; void* softmaxResultOutOptionalAddr nullptr; aclTensor* input nullptr; aclTensor* out nullptr; aclTensor* expertIdOut nullptr; aclTensor* softmaxResultOutOptional nullptr; std::vectorfloat inputHostData {0.1, 1.1, 2.1, 3.1, 4.1, 5.1, 6.1, 7.1, 8.1, 9.1, 10.1, 11.1}; std::vectorfloat outHostData {0.1, 1.1, 2.1, 3.1, 4.1, 5.1}; std::vectorint32_t expertIdOutHostData {1, 1, 1, 1, 1, 1}; std::vectorfloat softmaxResultOutOptionalHostData {1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; // 创建输入x对应的aclTensor ret CreateAclTensor(inputHostData, inputShape, inputAddr, aclDataType::ACL_FLOAT, input); CHECK_RET(ret ACL_SUCCESS, return ret); // 创建输出yOut对应的aclTensor ret CreateAclTensor(outHostData, outShape, outAddr, aclDataType::ACL_FLOAT, out); CHECK_RET(ret ACL_SUCCESS, return ret); // 创建输出expertIdxOut对应的aclTensor ret CreateAclTensor(expertIdOutHostData, expertIdOutShape, expertIdOutAddr, aclDataType::ACL_INT32, expertIdOut); CHECK_RET(ret ACL_SUCCESS, return ret); // 创建输出softmaxResultOutOptional对应的aclTensor ret CreateAclTensor(softmaxResultOutOptionalHostData, softmaxResultOutOptionalShape, softmaxResultOutOptionalAddr, aclDataType::ACL_FLOAT, softmaxResultOutOptional); CHECK_RET(ret ACL_SUCCESS, return ret); // 3.调用CANN算子库API需要修改为具体的算子接口 uint64_t workspaceSize 0; aclOpExecutor* executor; // 调用aclnnMoeGatingTopKSoftmaxV2第一段接口 ret aclnnMoeGatingTopKSoftmaxV2GetWorkspaceSize(input, nullptr, 2, 0, true, out, expertIdOut, softmaxResultOutOptional, workspaceSize, executor); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnMoeGatingTopKSoftmaxV2GetWorkspaceSize 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); } // 调用aclnnMoeGatingTopKSoftmaxV2第二段接口 ret aclnnMoeGatingTopKSoftmaxV2(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret ACL_SUCCESS, LOG_PRINT(aclnnMoeGatingTopKSoftmaxV2 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.0f); ret aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outAddr, size * 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 size; i) { LOG_PRINT(result[%ld] is: %f\n, i, resultData[i]); } // 6. 释放aclTensor和aclScalar需要根据具体API的接口定义修改 aclDestroyTensor(input); aclDestroyTensor(out); aclDestroyTensor(expertIdOut); aclDestroyTensor(softmaxResultOutOptional); // 7. 释放device资源需要根据具体API的接口定义修改 aclrtFree(inputAddr); aclrtFree(outAddr); aclrtFree(expertIdOutAddr); aclrtFree(softmaxResultOutOptionalAddr); if (workspaceSize 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }示例中值得注意的几点第一段接口调用时finishedOptional传入nullptr表示不使用 finished 标记属性传参顺序为(k2, renorm0, outputSoftmaxResultFlagtrue)即先 Softmax 再 TopK并输出 Softmax 中间结果此时softmaxResultOutOptional的 shape 应为{3, 4}与x一致yOut{3, 2}与expertIdxOut{3, 2}的最后一维等于k与前述 shape 推导规则吻合。十、源码级原理剖析10.1 算子定义OpDef在 moe_gating_top_k_softmax_v2_def.cpp 中算子通过OpDef注册输入x必选FLOAT16/FLOAT/BF16ND 格式声明AutoContiguous()输入finished可选BOOLND 格式输出y必选对应 aclnn 层yOut、expert_idx必选对应expertIdxOutINT32、softmax_result可选FLOAT32属性k必选 Int、renorm可选默认 0、output_softmax_result_flag可选默认 false平台配置ascend910b与ascend910_93使用动态编译 kernelascend950使用 RegBase kernelmoe_gating_top_k_softmax_v2_apt且三个平台均开启动态 shape、动态 rank 支持。10.2 Infershape输出 shape 的推导Infershape 实现的核心逻辑setOutShape为yOut与expertIdxOut继承x除最后一维外的所有维度最后一维设为ksoftmaxResultOutOptional在需要输出时与x同 shape否则 shape 为[0]。同时 Infershape 完成三项关键校验x维度必须为 2 或 3、0 k 1024且k x最后一维、renorm只能取 0/1。值得注意的是softmaxFlag的生效条件是renorm 0 output_softmax_result_flag true——这与文档中renorm1 时 outputSoftmaxResultFlag 不生效的语义完全一致。10.3 Tiling按场景分发的分块策略Tiling 入口 moe_gating_top_k_softmax_v2_tiling.cpp 通过TilingRegistry分发到具体实现。仓库为不同数据规模准备了多套 Tiling 模板MoeGatingTopKSoftmaxV2EKFullLoadTilingData专家数与 k 均能整块加载E×K FullLoad的场景MoeGatingTopKSoftmaxV2KFullLoadTilingData仅 k 能整块加载K FullLoad的场景MoeGatingTopKSoftmaxV2PerfTilingData/PerfRegbaseTilingData面向列数 ≤1024 且满足 8 对齐等条件的性能优化路径。Tiling 数据字段如row、col、kAlignB16、blockNum、ubLoopOfFormerBlock等记录了行/列切分、UB 缓冲循环次数、Softmax 与 TopK 各自的内层 Tiling 参数见 moe_gating_top_k_softmax_v2_tiling.h。Tiling 阶段会先执行GetPlatformInfo获取 AIV 核数、UB 大小与GetShapeAttrsInfo校验 shape、dtype 与属性并规整出row/col其中 3D 输入会被展开为row dim0 * dim1、col dim2处理见 moe_gating_top_k_softmax_v2_tiling_base.cpp。Tiling Key 的编码规则也值得关注TOPK_SOFTMAX_TILING_KEY_BASE_ALL 100000、_SCENE 1000、_RENORM 100、_DTYPE 10再叠加列数区间8、8~64、64共同组成如103113之类的分发键1perf 场景0/1renorm0/1/3dtype float/half/bf161/2/3列数区间。这一编码在 moe_gating_top_k_softmax_v2_tiling.h 中有完整定义。10.4 Kernel多模板分发与 AIV 执行Kernel 入口 moe_gating_top_k_softmax_v2.cpp 声明为__global__ __aicore__并通过KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIV_1_0)指定由 AIV 执行g_coreType AIC时直接返回。随后根据 Tiling Key 分派到四个模板实现MoeGatingTopKSoftmaxV2EKFullLoadINPUT_TYPE, BUFFER_NUM, RENORM见 moe_gating_top_k_softmax_v2_ek_fullload.hMoeGatingTopKSoftmaxV2KFullLoadINPUT_TYPE, RENORM见 moe_gating_top_k_softmax_v2_k_fullload.hMoeGatingTopKSoftmaxV2KRenormINPUT_TYPE, RENORM见 moe_gating_top_k_softmax_v2_k_renorm.hMoeGatingTopKSoftmaxV2PerfINPUT_TYPE, RENORM, COL_RANGE见 moe_gating_top_k_softmax_v2_perf.h。BUFFER_NUM模板参数1 或 2控制双缓冲/单缓冲策略COL_RANGE对应列数区间。各模板在arch35目录下还有对应的架构优化版本如 moe_gating_top_k_softmax_v2_perf_arch35.h。10.5 golden 校验正确性的数学基准测试目录 tests/assets/golden.py 提供了 NumPy/Torch 双实现的 golden 参考先将x统一转 float32再按renorm分支执行先 Softmax 后 TopK或先 TopK 后 Softmax并对finished行填充num_expert索引。单测代码见 tests/ut/op_host/test_MoeGatingTopKSoftmaxV2_infershape.cpp 与 tests/ut/op_kernel/test_moe_gating_top_k_softmax_v2.cppST 用例配置文件位于 tests/st/aclnnMoeGatingTopKSoftmaxV2/atk_aclnnMoeGatingTopKSoftmaxV2.json可据此构造各种 shape、dtype、k、renorm与finished组合的回归用例。十一、总结MoeGatingTopKSoftmaxV2将 MoE 路由阶段高频出现的 Softmax 与 TopK 融合为单算子通过renorm属性灵活切换两种计算顺序配合finishedOptional与可选 Softmax 中间结果输出能够适配训练与推理含序列结束标记等不同路由场景。从仓库源码可以清晰看到从 aclnn 两段式接口、OpDef 定义、Infershape、Tiling 分块到 AIV Kernel 多模板分发的完整软件栈以及 golden 测试对数学正确性的兜底。开发者可直接复用调用示例快速集成并参考本仓库的算子目录深入研读实现细节。【免费下载链接】ops-transformer本项目是CANN提供的transformer类大模型算子库实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-transformer创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考

CANN ops-transformer 算子 MoeGatingTopKSoftmaxV2：aclnnMoeGatingTopKSoftmaxV2 接口详解与源码剖析

相关推荐

Xftp超详细安装与配置教程：解决连接失败、中文乱码、服务异常

CLI驱动的Git Diff代码评审：LLM Agent如何重塑Code Review工作流

N_m3u8DL-RE 上手指南：MPD/M3U8/ISM 流媒体下载、解密与直播录制一次搞定

更多资讯

Atlas 300V 24G昇腾推理卡实战：从环境配置到YOLOv8部署调优

Ant Design Popconfirm 条件触发实战：用受控 open 与 onOpenChange 拦截实现“按需弹出“

Hugo Page.ReadingTime 方法详解：估算阅读时间的计算原理与多语言定制

N_m3u8DL-RE 流媒体下载实战：3 条命令把在线课程和直播存下来

信息化售后服务方案：从文档交付到自动运维闭环

微信小程序活动报名管理系统设计与实现

今日更新

校园网认证计费模式详解：AAA与代拨架构、原理及排障

Spring Boot+Vue垃圾分类毕设：分层、权限与排错实战

把 Cline 的 Base URL 改到 TaoToken 之后，DeepSeek 仍能修调试配置

本周精选

AI驱动的自动化专利分析平台技术解析

LeetCode-Go 题解：1252. Cells with Odd Values in a Matrix（奇数值单元格计数）双解法剖析

在线答疑系统Java毕设实战：Spring Boot前后端分离与状态流转

本月热门

基于YOLO的焊缝缺陷检测：从数据标注到工业部署全流程实践

分布式服务的决策记录

创业产品的首版范围

读完文章，想聊聊您的网站？