人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读本文围绕 examples/01_simd_cpp_api/04_advanced_api/05_sort 样例集展开讲解如何在 CANN asc-devkit 中使用 Sort 高阶 API 实现降序全排序、使用 TopK 高阶 API 实现沿最后一维取前 K 个最大/最小值及其索引的端到端算子工程。读完本文你将掌握 Kernel 直调算子工程的目录组织、tiling 参数计算、核函数编写、host 端调用与精度校验的完整流程并能理解GetSortMaxMinTmpSize、GetTopKMaxMinTmpSize、TopKTilingFunc等 tiling 接口在底层临时空间分配中的实际作用。该样例集包含两个子样例sort全排序与topkTopK 选择两者均基于 Kernel 直调算子工程支持 Ascend 950PR/Ascend 950DT、Atlas A3 训练/推理系列、Atlas A2 训练/推理系列产品。样例集总览排序操作的两类典型场景目录名称功能描述支持的产品sort基于 Kernel 直调算子工程调用 Sort 高阶 API 实现 sort 单算子按数值大小降序排序并保留原始索引信息Ascend 950PR/Ascend 950DTAtlas A3 训练系列产品/Atlas A3 推理系列产品Atlas A2 训练系列产品/Atlas A2 推理系列产品topk基于 Kernel 直调算子工程调用 TopK 高阶 API 实现 topk 单算子获取最后一维的前 K 个最大值或最小值及其对应索引Ascend 950PR/Ascend 950DTAtlas A3 训练系列产品/Atlas A3 推理系列产品Atlas A2 训练系列产品/Atlas A2 推理系列产品两个样例的计算流程在结构上完全一致输入数据先搬入片上存储GM → Local调用高阶 API 完成排序/选择计算再将结果搬出Local → GM。差别仅在于所调用的高阶 API、输入输出规格与 tiling 参数的来源。下文分别对两个样例做源码级剖析。Sort 样例详解全量降序排序并保留索引功能与输入输出规格Sort 样例对输入张量按数值大小进行降序排序同时保留原始索引信息支持float数值与uint32_t索引的联合排序。样例规格如下对应 sort/README_en.md项目名称shape数据类型format输入srcLocalValue[1, 128]floatND输入srcLocalIndex[1, 128]uint32_tND输出dstLocalValue[1, 128]floatND输出dstLocalIndex[1, 128]uint32_tND核函数名sort_custom---该样例使用固定 shape输入srcLocalValue[1, 128]、srcLocalIndex[1, 128]输出dstLocalValue[1, 128]、dstLocalIndex[1, 128]。目录结构sort ├── scripts │ └── gen_data.py // 生成输入数据与期望golden数据 ├── CMakeLists.txt // 构建工程文件 ├── data_utils.h // 数据读写函数 ├── sort.asc // Ascend C 样例实现与调用示例 └── README.md // 样例文档其中 sort.asc 是核心实现文件同时包含 tiling 计算、核函数类、host 端 main 调用三部分data_utils.h 提供ReadFile/WriteFile二进制文件读写工具CMakeLists.txt 以ASC语言将sort.asc编译为可执行文件demo。Tiling 实现GetSortMaxMinTmpSize 计算临时空间Sort 的 tiling 在ComputeTiling中完成sort.ascvoid ComputeTiling(uint32_t elementCount, optiling::SortCustomTilingData tiling) { std::vectorint64_t shapeVec {elementCount}; AscendC::TensorShape srcShape(shapeVec); AscendC::SortConfig config; config.type AscendC::SortType::RADIX_SORT; config.isDescend false; uint32_t maxTmpSize; uint32_t minTmpSize; // 获取 Sort 操作所需的最大和最小临时空间大小 AscendC::GetSortMaxMinTmpSize( srcShape, AscendC::TensorDataType::DT_FLOAT, AscendC::TensorDataType::DT_UINT32, false, config, maxTmpSize, minTmpSize); uint32_t localWorkspaceSize minTmpSize; tiling.set_elementCount(elementCount); tiling.set_sharedTmpBufferSize(localWorkspaceSize); }关键点tiling 数据结构通过宏BEGIN_TILING_DATA_DEF/TILING_DATA_FIELD_DEF/END_TILING_DATA_DEF定义包含elementCount待排序元素个数与sharedTmpBufferSize共享临时 buffer 大小两个字段并通过REGISTER_TILING_DATA_CLASS(SortCustom, SortCustomTilingData)注册。GetSortMaxMinTmpSize的接口声明位于 include/adv_api/sort/sort_tiling_intf.hsrcShape为输入 shapevalueType/indexType为数值与索引数据类型isReuseSource表示是否可复用源内存config为排序配置输出maxValue与minValue分别为最大/最小临时空间。SortConfig结构体见 sort_tiling_intf.h 与 sort_utils.h包含四个成员type排序算法类型默认SortType::RADIX_SORT即基数排序、isDescend是否降序、hasSrcIndex、hasDstIndex。本样例取minTmpSize作为共享临时空间大小以最小化片上 buffer 占用。生成后的 tiling 数据通过SaveToBuffer序列化到内存并由GenerateTiling返回给 host 端最终经aclrtMemcpy拷入设备端 tiling 地址。核函数实现CopyIn → Sort/Extract → CopyOut核函数类KernelSortT, isFullSortsort.asc遵循经典的 Ascend C 三段式流水CopyIn通过AscendC::DataCopy将mValueGlobal、mIndexGlobal两个 GlobalTensor 搬入 VECIN 队列入队等待计算。Compute出队后先经AscendC::Concat拼接数据重复次数mConcatRepeatTimes elementCount / 16随后调用核心计算sort.ascAscendC::SortT, isFullSort(sortedLocal, concatLocal, indexLocal, sortTmp, mSortRepeatTimes);Sort模板参数T指定数据类型本样例为floatisFullSort指定是否全排序本样例为true。排序完成后调用AscendC::Extract(dstValueLocal, dstIndexLocal, sortedLocal, mExtractRepeatTimes)将排序结果与索引分别抽取到输出张量。临时工作空间有两路来源当 tiling 给出的sharedTmpBufferSize 0时使用共享临时 bufferdstQueSharedTmp否则使用核内自建 buffersortTmpLocal体现以 tiling 结果驱动资源分配的设计。队列均通过TPipe::InitBuffer初始化输入队列深度为 2双缓冲输出队列与临时队列深度为 1。CopyOut将 VECOUT 队列中的结果DataCopy回 Global 内存dstGmValue/dstGmIndex。核函数入口sort_customsort.asc声明为extern C __global__ __vector__接收源值、源索引、目标值、目标索引、workspace、tiling 六路__gm__地址参数host 侧通过CopyTiling将 GM 上的 tiling 数据解析到结构体后调用op.Process()。host 端调用与数据校验main函数sort.asc的流程是典型的 Kernel 直调模板初始化aclInit→aclrtSetDevice(0)→aclrtCreateContext→aclrtCreateStream。内存准备aclrtMallocHost分配 host 侧内存aclrtMalloc(..., ACL_MEM_MALLOC_HUGE_FIRST)分配设备侧内存。固定ELEMENTCOUNT 128workspace 固定16 * 1024 * 1024字节tiling 数据为 2 个uint32_t。数据搬运ReadFile读取./input/input_srcGmValue.bin、./input/input_srcGmIndex.bin经aclrtMemcpy(ACL_MEMCPY_HOST_TO_DEVICE)上板GenerateTiling(ELEMENTCOUNT)生成的 tiling 同样拷入设备。核函数启动sort_customnumBlocks, 0, stream( srcGmValueDevice, srcGmIndexDevice, dstGmValueDevice, dstGmIndexDevice, workspaceDevice, tilingDevice);内核调用运算符传入 block 数固定 1、l2 模式0与 stream。结果回读与校验aclrtSynchronizeStream同步后将输出ACL_MEMCPY_DEVICE_TO_HOST回拷WriteFile落盘为./output/output_dstGmValue.bin与./output/output_dstGmIndex.binCompareResult以绝对误差/相对误差均不超1e-4为阈值与./output/golden_*.bin逐元素比对。测试数据与 golden 生成gen_data.py 使用 numpy 生成测试数据x1为[-60000, 60000]均匀分布的 128 个 float32 随机数x2为 128 个 0 组成的 int32 索引golden 通过-np.sort(-x1)技巧实现降序排序并落盘。索引张量初始全 0 且 Sort 输出会重排索引因此该样例侧重验证数值排序结果与索引随动关系。TopK 样例详解沿最后一维取前 K 个极值功能与输入输出规格TopK 样例在 Normal 模式下对输入张量进行 TopK 计算沿最后一维获取前 K 个最大值或最小值及其对应索引支持float数值与int32_t索引样例规格见 topk/README_en.md项目名称shape数据类型format输入srcLocalValue[2, 32]floatND输入srcLocalIndex[1, 32]int32_tND输入srcLocalFinish[1, 32]int32_tND输出dstLocalValue[2, 8]floatND输出dstLocalIndex[2, 8]int32_tND核函数名topk_custom---与 Sort 不同TopK 多了一个srcLocalFinishfinish 标志输入该标志用于标识某行是否已完成若置位则输出索引被改写为总长度 N。样例固定 shapesrcLocalValue[2, 32]、srcLocalIndex[1, 32]、srcLocalFinish[1, 32]输出dstLocalValue[2, 8]、dstLocalIndex[2, 8]K1 时按 8 对齐的 padding 输出。Tiling 实现TopKTilingFunc GetTopKMaxMinTmpSizeTopK 的 tiling 计算在 topk.asc 的GenerateTiling中完成AscendC::TopKTilingFunc( ascendCPlatform, inner, outter, k, dtypesize, false, AscendC::TopKMode::TOPK_NORMAL, isLargest, tiling.topKTilingData); AscendC::GetTopKMaxMinTmpSize( ascendCPlatform, inner, outter, false, false, AscendC::TopKMode::TOPK_NORMAL, isLargest, dtypesize, maxsize, minsize); tiling.set_minsize(minsize);关键点tiling 数据类TopKCustomTilingData内嵌TopkTiling topKTilingData结构体外加k、outter、inner、n、minsize、isLargest六个字段其中outter/inner分别对应输入张量第一维行数与第二维每行长度n为数据总长度isLargest为 true 表示取最大值、false 表示取最小值。TopKTilingFunc声明见 include/adv_api/sort/topk_tiling.h根据平台信息、inner/outter、K 值、数据类型与模式生成 TopK 接口所需的 tiling 信息。GetTopKMaxMinTmpSizetopk_tiling.h计算 TopK 计算所需的最大/最小临时片上空间样例取minsize作为局部临时 buffer 大小并做 32 字节对齐修正LOCAL_BYTES 32。平台对象PlatformAscendC由PlatformAscendCManager::GetInstance获取CPU 调试模式下按__NPU_ARCH__2201 对应Ascend910B13510 对应Ascend950PR_9589显式指定NPU 模式则直接获取当前运行平台。核函数实现模板参数驱动的 TopK 调用核函数类KernelTopKT, U, isInitIndex, isHasfinish, isReuseSrc, topkMode, tmpLocaltopk.asc通过模板参数表达多种组合是否初始化索引、是否携带 finish 标志、是否复用源数据、TopK 模式、是否使用局部临时空间。本样例实例化为KernelTopKfloat, int32_tisHasfinish true、tmpLocal true。Compute 阶段的核心调用topk.ascAscendC::TopKT, isInitIndex, isHasfinish, isReuseSrc, AscendC::TopKMode::TOPK_NORMAL( dstLocalValue, dstLocalIndex, srcLocalValue, srcLocalIndex, srcLocalFinish, k, topKTilingData, topKInfo, isLargest);要点TopKInfo定义见 include/adv_api/sort/topk_utils.h携带outter、inner32 字节对齐后的 inner、n张量实际长度三个维度信息。TopKMode枚举topk_utils_constants.h包含TOPK_NORMALNormal 大数据量模式与TOPK_NSMALL小数据量模式本样例固定TOPK_NORMAL但核函数模板在topkMode true时切换到TOPK_NSMALL分支展示两种模式在统一模板下的适配能力。输出前先用AscendC::Duplicate将dstLocalValue/dstLocalIndex清零K 值按数据类型对齐为kPadfloat 按 8 对齐、half 按 16 对齐对应K_FLOAT 8、K_HALF 16输出 shape 随之扩展为kPad × outter。临时空间路径tmpLocal true时在 VECCALC 位置初始化TBuf临时 buffer 并传入 TopKtmpLocal false时直接传入topKTilingData完成计算展示两种资源供给方式。核函数入口topk_customtopk.asc接收七路__gm__参数源值、源索引、finish、目标值、目标索引、workspace、tilinghost 侧固定OUTTER 2、INNER 32、N 32、K 1、ISLARGE false即取每行最小值。host 端调用与 finish 语义验证TopK 的 host 流程与 Sort 一致aclInit→ 内存分配 → 读输入 → 搬 tiling →topk_customnumBlocks, 0, stream→ 同步 → 回拷 → 比对差异在于多读一路input_finishGm.bin且 tiling 大小按TILINGDATA_SIZE 33 6个uint32_t计算33 为内嵌TopkTiling结构体字段数。golden 生成topk/scripts/gen_data.py通过np.argsort(x1, axis-1, kindstable)得到稳定排序索引取前 K 列作为期望索引再用np.take_along_axis取回对应数值同时按 K 的 8 对齐补零 padding并对 finish 置位的行将索引改写为 N——这正好呼应了核函数中 finish 标志的语义。构建与运行三种运行模式在两个样例各自目录sort/topk下按以下步骤构建运行。配置环境变量根据当前环境的 CANN 开发套件安装方式配置环境变量source ${install_path}/cann/set_env.sh注意${install_path}为 CANN 包安装目录未指定安装目录时默认安装路径为/usr/local/Ascend。运行样例NPU 模式默认mkdir -p build cd build; cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # NPU 模式默认 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo切换 CPU 调试 / NPU 仿真模式cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # CPU 调试模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # NPU 仿真模式注意切换构建模式前需在 build 目录中删除CMakeCache.txt清理 cmake 缓存后重新执行 cmake。构建选项说明选项取值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2 训练/推理系列与 Atlas A3 训练/推理系列产品dav-3510 对应 Ascend 950PR/Ascend 950DT构建层面两个样例的 CMakeLists.txt 结构相同find_package(ASC REQUIRED)引入 ASC 语言工具链project(kernel_samples LANGUAGES ASC CXX)声明工程语言add_executable(demo sort.asc)或topk.asc生成可执行文件并链接tiling_api、register、platform、unified_dlog、dl、m、graph_base等库通过--npu-arch${CMAKE_ASC_ARCHITECTURES}编译选项把目标架构传给 ASC 编译器。执行结果以下输出表示精度比对成功test pass!若比对失败程序输出test failed!且CompareResult会打印首个不一致元素的输出值与 golden 值便于定位问题。从样例到高阶 API 接口关键能力清单将样例中的调用点与仓库头文件对应可以得到排序高阶 API 的完整能力视图接口声明位置样例中的用途SortT, isFullSort核函数高阶 API全量降序排序并携带索引sort.ascGetSortMaxMinTmpSizeinclude/adv_api/sort/sort_tiling_intf.h计算 Sort 最大/最小临时空间取最小空间建 bufferSortConfig/SortType::RADIX_SORTinclude/adv_api/sort/sort_utils.h指定基数排序算法与升降序、源/目标索引开关TopKT, ...include/adv_api/sort/topk.hNormal 模式取前 K 个极值及索引topk.ascTopKTilingFuncinclude/adv_api/sort/topk_tiling.h生成 TopK 接口所需的 tiling 信息GetTopKMaxMinTmpSizeinclude/adv_api/sort/topk_tiling.h计算 TopK 最大/最小临时空间TopKInfoinclude/adv_api/sort/topk_utils.h携带 outter/inner/n 维度信息TopKMode/TopKAlgo/TopKConfiginclude/adv_api/sort/topk_utils_constants.hNormal/NSMALL 模式、RADIX_SELECT/MERGE_SORT 算法、排序配置值得注意的实现细节临时空间的两级决策Sort 样例以GetSortMaxMinTmpSize的minTmpSize为共享 buffer 大小TopK 样例将minsize对齐到 32 字节后作为 VECCALC 局部 buffer当最小空间为零或不可用时核内自建 buffer 作为兜底。这套tiling 定尺寸、核内按需取用的模式是可复用的资源规划范式。模板参数驱动行为差异KernelTopK的七个模板参数将是否初始化索引、是否有 finish、是否复用源、大小模式、临时空间位置全部编译期展开避免运行期分支开销这也是高阶 API 样例常见的组织方式。数据生成与校验闭环gen_data.py 生成的 golden 与核内逻辑严格对应如 finish 改写索引、K 的 8 对齐 padding保证样例可以自证正确性。总结与扩展建议本样例集以最小可运行的 Kernel 直调工程完整示范了昇腾排序类算子的两种典型写法Sort面向全量排序float uint32_t数值/索引联合输出适合需要排完还要知道原来位置的场景TopK面向极值选择支持 Normal/NSMALL 双模式、最大值/最小值切换与 finish 语义适合只关心前 K 个结果的高效检索场景。实际开发中你可以在此基础上做三类扩展一是将固定 shape 改为动态 shape让elementCount/inner/outter由运行时输入推导二是通过修改SortConfig.isDescend与isLargest参数切换升/降序、最大/最小 K三是参照TopKAlgo的RADIX_SELECT与MERGE_SORT两种算法枚举topk_utils_constants.h对比不同数据规模下的取舍。构建与调试时建议先用cpu模式在宿主环境快速验证逻辑再切sim仿真最后上 NPU 实测配合样例自带的 golden 比对机制即可完成全链路精度确认。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐Wand-Enhancer 完整指南本地补丁拆除 2 小时限制Wand Enhancer 完整指南本地补丁拆除 2 小时限制 Wand Enhancer 是面向 WandWeMod客户端的开源本地增强工具。它把客户端人工智能深度学习算子库CANNAscendCANN asc-devkit Math API 算子样例全解析从 Acosh 到 Xor 的昇腾高阶数学算子实现与端到端验证CANN asc devkit Math API 算子样例全解析从 Acosh 到 Xor 的昇腾高阶数学算子实现与端到端验证 导读 本文以 CANN asc人工智能深度学习算子库CANNAscend昇腾 Ascend C Arange 索引计算算子实战基于 asc-devkit 高阶 API 的等差数列生成与端到端实现解析昇腾 Ascend C Arange 索引计算算子实战基于 asc devkit 高阶 API 的等差数列生成与端到端实现解析 导读 本文围绕 asc dev人工智能深度学习算子库CANNAscend上一篇终极宽屏体验PvZWidescreen让植物大战僵尸焕然新生 下一篇从GROMACS轨迹到结合自由能为什么你需要gmx_MMPBSA创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考