人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读本文以 asc-devkit 仓库中examples/01_simd_cpp_api/04_advanced_api/06_index/arange样例为核心系统讲解如何基于昇腾 AI 处理器的 Ascend C 高阶 API——Arange等差序列生成接口实现给定起始值、等差值与长度在片上UB存储中生成指定长度的等差数列并搬运回外部存储GM的完整算子开发流程。读完本文你将掌握 Arange 接口的函数原型、参数约束、底层实现原理与性能特征能够独立完成样例的编译、三种运行模式NPU 运行、CPU 调试、NPU 仿真切换、输入数据与真值生成以及精度对比验证同时理解该样例在仓库中的源码组织结构与配套测试证据。样例概述与功能定位Arange 样例实现的是等差数列生成功能给定起始值firstValue、等差值diffValue和长度count返回一个满足dst[i1] dst[i] diffValue的等差数列。在算子开发实践中等差数列序列常被用作索引生成、序列构造、查表前处理等场景的基础数据属于典型的数据构造类算子。样例将输入起始值与等差值作为标量参数shape 为[1,1]的 float 数据通过 Arange 高阶 API 在片上存储中生成指定长度的等差数列再通过 DataCopy 将结果搬运至外部存储。核函数名为arange_custom。样例输入输出规格名称shape数据类型格式firstGm输入[1,1]floatNDdiffGm输入[1,1]floatNDoutputGm输出[104,1]floatND需要注意输出 shape 为[104,1]而非[100,1]样例固定生成 100 个 float 元素而片上搬运与存储按 32 字节对齐100个 float400 字节向上对齐后实际写入 104 个元素(1008-1)/8*8104每 8 个 float 为 32 字节对齐块。该对齐逻辑在 arange.asc 中通过count_aligned计算并与 gen_data.py 保持严格一致。支持的产品及 CANN 软件版本产品CANN 软件版本Ascend 950PR950DT 系列产品 CANN 9.1.0Atlas A3 系列产品 CANN 9.0.0Atlas A2 系列产品 CANN 9.0.0从仓库中 Arange 接口的产品支持矩阵docs/zh/api/SIMD-API/adv_api/index_compute/Arange.md可以看到更完整的边界Atlas 推理系列产品 AI Core 支持、Vector Core 不支持Atlas 训练系列产品与 Atlas 200I/500 A2 推理产品不支持。因此本样例覆盖的三个产品系列是 Arange 接口在当前仓库中的主要应用对象。目录结构与工程骨架样例根目录结构如下├── arange │ ├── scripts │ │ └── gen_data.py // 输入数据和真值数据生成脚本 │ ├── CMakeLists.txt // 编译工程文件 │ ├── data_utils.h // 数据读入写出函数 │ ├── arange.asc // Ascend C 样例实现 调用样例 │ └── README.md // 样例说明文档编译工程CMakeLists.txtCMakeLists.txt 展示了基于 ASC 语言的最小可运行编译工程cmake_minimum_required(VERSION 3.16) find_package(ASC REQUIRED) project(kernel_samples LANGUAGES ASC CXX) add_executable(demo arange.asc) set(CMAKE_ASC_ARCHITECTURES dav-2201 CACHE STRING CMAKE_ASC_ARCHITECTURES, e.g. dav-2201, dav-3510) target_link_libraries(demo PRIVATE tiling_api register platform unified_dlog dl graph_base) target_compile_options(demo PRIVATE $$COMPILE_LANGUAGE:ASC:--npu-arch${CMAKE_ASC_ARCHITECTURES})关键点说明find_package(ASC REQUIRED)引入 asc-devkit 提供的 ASC 编译器与构建支撑相关配置见 cmake/asc/asc-config.cmake工程默认将CMAKE_ASC_ARCHITECTURES缓存为dav-2201并以--npu-arch编译选项传给 ASC 编译器链接库tiling_api提供 tiling 相关支撑本样例使用 Arange 相关的临时空间查询接口register用于核函数注册platform提供运行时平台能力。数据读写工具data_utils.hdata_utils.h 提供两个二进制文件读写函数ReadFile(filePath, fileSize, buffer, bufferSize)以二进制方式读取文件内部校验文件存在性、是否为普通文件、文件大小不超过缓冲区WriteFile(filePath, buffer, size)以O_RDWR|O_CREAT|O_TRUNC方式创建/覆盖写入二进制文件。host 端代码通过这两个函数读入input/first.bin、input/diff.bin输入并写出output/output.bin结果文件。Arange 高阶 API 原理与参数约束函数原型Arange 高阶 API 的公开接口定义位于 include/adv_api/index/arithprogression.htemplate typename T __aicore__ inline __in_pipe__(S) __out_pipe__(V, S) void Arange( const LocalTensorT dst, const T firstValue, const T diffValue, const int32_t count) { ArithProgressionImpl(dst, firstValue, diffValue, count); }参数说明参数输入/输出描述dst输出目的操作数 LocalTensor大小应大于等于count * sizeof(T)支持的 TPosition 为 VECIN/VECCALC/VECOUTfirstValue输入等差数列的首个元素值diffValue输入等差数列相邻元素之间的差值应大于等于 0count输入等差数列的长度count 0支持的数据类型Ascend 950PR950DT 系列产品int16_t、half、int32_t、float、int64_tAtlas A3 / Atlas A2 系列产品int16_t、half、int32_t、float。该约束在 3510 实现中通过static_assert(SupportTypeT, int16_t, int32_t, half, float, int64_t(), ...)进行编译期检查见 arithprogression_3510_impl.h。格式约束为仅支持 ND 格式。参数校验机制在 CPU 调试ASCENDC_CPU_DEBUG与内核侧检查场景下arithprogression_check_common.h 通过CheckArithProgressionParamsClass完成运行时校验count必须大于 0否则报错 The count parameter cannot be %d, should be greater than 0diffValue必须大于等于 0否则报错 The diffValue parameter cannot be %f, should be greater than or equal to 0数据类型校验仅允许half/float/int16_t/int32_t输出 Tensor 必须位于 UBVECIN/VECOUT/VECCALC并校验count与 dst 长度的匹配关系。配套的校验测试 kernel_arithprogression_check.cpp 验证了合法调用不触发异常ArithProgressionAPICheckTestSuccess、非法调用uint8_t 类型 非正 count/diffValue恰好触发 3 次异常ArithProgressionAPICheckTestFailure两条路径。底层实现原理Arange 内部算法在 Vector 上进行核心思路是先生成基础等差数列再通过向量加法指数级扩充。API 文档docs/zh/api/SIMD-API/adv_api/index_compute/Arange.md将其概括为三步长度 8 以内按 firstValue 与 diffValue 的值用 SetValue 直接生成基础序列最大 8 个元素若目标长度小于 8 则算法结束长度 8 至 64对基础序列使用 Adds 逐块扩充最大循环 7 次扩充至 64 个元素若目标长度小于 64 则算法结束长度 64 以上对 64 元素序列使用 Adds 继续按大步长循环扩充直至达到目标长度。仓库中的公共实现 arithprogression_common_impl.h 与文档描述完全对应GetBaseArithProgressionT()用标量循环dstLocal.SetValue(i, firstValue diffValue * i)生成基础序列并对half做了特化先转 float 计算再转回 half避免精度截断当count BLOCK_NUM一个 32 字节块可容纳的元素数float 下为 8时先生成一块基础序列再通过Adds以GetArithProgressionStepT(diffValue, BLOCK_NUM)为增量逐块复制扩充当count REPEAT_NUM一个 Vector repeat 可处理的元素数float 下为 64时先扩充至 64 元素再以GetArithProgressionStepT(diffValue, REPEAT_NUM)为步长按 repeat 粒度扩充尾部不足一个 repeat 时用SetVectorMask掩码处理当count BLOCK_NUM时直接走标量 SetValue 路径并插入 S-V 事件同步SetFlag/WaitFlagHardEvent::S_V保证标量与向量流水之间的数据依赖。在 Ascend 950PR950DT3510 架构上实现则切换到寄存器级路径 arithprogression_3510_impl.h先用Reg::Arange生成寄存器内的 0..N 索引序列再经Reg::Muls(dstReg, dstReg, diffValue, fullMask)乘上等差值、Reg::Adds(dstReg, dstReg, firstValue, fullMask)加上起始值最后按 repeat 循环Reg::StoreAlign写入 UB并在每个 repeat 后对寄存器序列做一次Reg::Add平移。该路径还处理了int64_t8 字节元素场景此时使用双寄存器 traitRegTraitNumTwo且 repeat 次数减半。由此可见无论哪种实现路径Arange 都是纯计算、无外部数据依赖的接口——它不读任何输入 Tensor仅依据标量参数生成序列因此不需要申请额外 workspace 临时空间。Tiling 配套接口GetArangeMaxMinTmpSizeinclude/adv_api/index/arithprogression_tiling.h 提供了 Arange 的 tiling 辅助接口void GetArangeMaxMinTmpSize(uint32_t maxValue, uint32_t minValue);由于 Arange 内部不需要临时空间其实现arithprogression_tiling_impl.cpp中maxValue与minValue均返回 0。开发者可在 tiling 阶段调用该接口确认无需为 Arange 预留 workspace从而最大化 UB 利用率。注意旧接口GetArithProgressionMaxMinTmpSize已废弃并将在后续版本移除见 docs/zh/api/SIMD-API/adv_api/index_compute/GetArangeMaxMinTmpSize.md请使用新接口。样例实现拆解arange.ascarange.asc 是样例的唯一实现源文件包含 Tiling 定义、核函数与 host 端调用三大部分。第一步Tiling 数据定义与生成namespace optiling { BEGIN_TILING_DATA_DEF(ArithprogressionCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, outSize); END_TILING_DATA_DEF; REGISTER_TILING_DATA_CLASS(ArithprogressionCustom, ArithprogressionCustomTilingData) } // namespace optilingtiling 数据仅包含一个outSize字段用于把输出序列长度从 host 侧传递到核函数。ComputeTiling负责赋值GetTilingBuf将 tiling 数据序列化到缓冲区GenerateTiling封装完整生成流程。host 端通过aclrtMemcpy将该缓冲区拷贝至设备端 tiling 内存后随核函数下发。第二步核函数实现核函数内部封装了KernelArangeT类遵循 Ascend C 典型的Init → ProcessCopyIn → Compute → CopyOut三段式结构Init将起始值、等差值、长度保存为成员变量用SetGlobalBuffer绑定全局输出 Tensor并通过pipe-InitBuffer(outDst, 1, ...)在 UB 上按 32 字节对齐申请输出缓冲Compute核心调用即 Arange 高阶 API// dstLocal: 输出Tensor存储生成的等差数列 // static_castT(firstValue_): 起始值等差数列的第一个元素 // static_castT(diffValue_): 等差值相邻元素之间的差值 // count_: 元素个数生成的等差数列的长度 AscendC::ArangeT(dstLocal, static_castT(firstValue_), static_castT(diffValue_), count_);计算结果经outDst.EnQue入队CopyOutoutDst.DeQue取出结果用AscendC::DataCopy按 32 字节对齐后的块数BLOCK_NUM 32 / sizeof(T)即 float 下每 8 个元素一块将结果从 UB 拷贝至全局内存最后FreeTensor释放。核函数本体arange_custom的签名如下__global__ __vector__ void arange_custom( __gm__ uint8_t* firstGm, __gm__ uint8_t* diffGm, __gm__ uint8_t* dstGm, __gm__ uint8_t* workspace, __gm__ uint8_t* tiling)它先从 tiling 全局内存中解析出outSizeCopyTiling按 4 字节粒度逐字段拷贝再读取firstGm、diffGm指向的标量值初始化KernelArangefloat并执行。核函数声明为__vector__即运行在 Vector 计算单元上与 Arange 的向量实现路径一致。第三步host 端调用与验证host 端main函数完整演示了基于 ACLAscend CANN LiteRuntime 的算子调用流程运行环境初始化aclInit→aclrtSetDevice(0)→aclrtCreateContext→aclrtCreateStream内存分配与数据准备host 侧用aclrtMallocHost设备侧用aclrtMalloc(..., ACL_MEM_MALLOC_HUGE_FIRST)分别申请标量first/diff、输出、workspace 与 tiling 四类内存从./input/first.bin、./input/diff.bin读入输入aclrtMemcpy拷入设备核函数下发使用内核调用符以单 blockNUM_BLOCKS 1方式调用arange_custom随后aclrtSynchronizeStream同步等待完成结果回读与精度对比aclrtMemcpy将输出拷回 hostWriteFile写出./output/output.bin再调用CompareResult与./output/golden.bin逐元素对比详见下文精度对比小节资源回收依次释放设备/宿主内存、销毁 stream、context、reset 设备并aclFinalize。整个调用链展示了从Tiling 生成 → 数据搬入 → 核函数执行 → 结果搬出 → 精度校验的完整闭环可作为编写其他高阶 API 算子的通用模板。数据生成与精度对比机制输入与真值生成gen_data.pygen_data.py 使用 NumPy 生成样例所需的全部二进制文件start np.float32(1.234) difference np.float32(5.678) length np.uint32((100 8 - 1) // 8 * 8) # 100个float32元素按32字节向上对齐实际生成104个元素 stop start difference * (length - 1) golden np.linspace(start, stop, length, dtypenp.float32)first.bin/diff.bin输入起始值与等差值此处取 1.234 与 5.678便于观察序列行为length.bin对齐后的序列长度104与核函数侧count_aligned一致golden.bin使用np.linspace生成的等差序列真值供精度对比使用。精度对比逻辑CompareResult函数读取golden.bin对输出中每个 float 元素计算绝对误差ae与相对误差re ae / abs(b)当ae EPS re EPSEPS 1e-5时判定为错误并打印失配索引与数值全部元素通过后打印CompareResult golden.bin success!host 端最终打印test pass!表示精度对比成功。编译运行指南1. 配置环境变量请根据当前环境上 CANN 开发套件包的安装方式配置环境变量source ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。2. NPU 模式执行在样例根目录下执行如下命令mkdir -p build cd build; # 创建并进入 build 目录 cmake .. -DCMAKE_ASC_ARCHITECTURESdav-2201; make -j; # 编译工程 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行编译生成的可执行程序执行样例3. CPU 调试与 NPU 仿真模式使用 CPU 调试或 NPU 仿真模式时添加-DCMAKE_ASC_RUN_MODEcpu或-DCMAKE_ASC_RUN_MODEsim参数即可cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # CPU调试模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # NPU仿真模式注意切换编译模式前需清理 cmake 缓存可在 build 目录下执行rm CMakeCache.txt后重新 cmake。4. 编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2 系列产品和 Atlas A3 系列产品dav-3510 对应 Ascend 950PR950DT 系列产品其中dav-2201作为默认值同时固化在样例 CMakeLists.txt 的set(CMAKE_ASC_ARCHITECTURES dav-2201 CACHE STRING ...)中。三个架构号与 Arange 接口实现的分支对应dav-2201走公共向量实现arithprogression_common_impl.hdav-3510走寄存器实现arithprogression_3510_impl.h。CPU 调试模式下还会额外启用ASCENDC_CPU_DEBUG宏从而激活 Arange 的参数校验路径见 arithprogression_common_impl.h。5. 执行结果test pass!输出test pass!即表示核函数计算结果与 golden 真值在1e-5误差阈值内完全一致精度对比成功。仓库配套测试与更多参考资料Arange 功能在当前仓库中有成体系的测试支撑功能测试test_operator_arange.cpp以 half 类型对 count15/64/260 三组参数化用例调用Arange并执行DataCopy输出验证不同长度跨度的生成正确性覆盖 8 以内、864、64 以上三条实现路径API 校验测试kernel_arithprogression_check.cpp验证合法调用零异常、非法参数精确触发 3 次异常Tiling 测试在 tests/api/adv_api/tiling/test_tiling.cpp 中通过platfrom_stub_set_chip_version切换芯片版本后调用临时空间查询接口验证返回值为 0。如需进一步深入接口语义可查阅 API 文档 Arange含算法框图与约束说明与 GetArangeMaxMinTmpSize样例的英文版本见 README_en.md。本样例位于examples/01_simd_cpp_api/04_advanced_api/06_index/目录下与同目录的 index 系列样例共同构成 Ascend C 高阶 API 索引类算子的学习素材。小结通过本文可以确认Arange 高阶 API 在 asc-devkit 中的落地形态是一个零临时空间、纯向量计算的序列生成接口接口层提供Arange(dst, firstValue, diffValue, count)的统一签名公共实现通过SetValue 生成基础序列 Adds 指数级扩充最大化向量利用率3510 架构则进一步下沉到寄存器级Reg::Arange/Muls/Adds路径样例arange.asc完整演示了从 tiling 定义、核函数封装到 ACL 运行时调用的全过程配合gen_data.py的数据生成与1e-5精度的 golden 对比构成了可复制、可验证的算子开发范式。无论用于学习高阶 API 用法还是作为自研序列生成算子的起点本样例都值得直接对照阅读与运行。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐昇腾 Ascend C Arange 索引计算算子实战基于 asc-devkit 高阶 API 的等差数列生成与端到端实现解析昇腾 Ascend C Arange 索引计算算子实战基于 asc devkit 高阶 API 的等差数列生成与端到端实现解析 导读 本文围绕 asc dev人工智能深度学习算子库CANNAscend基于 Arange 高阶 API 实现等差数列生成CANN asc-devkit 索引算子Index示例深度解析基于 Arange 高阶 API 实现等差数列生成CANN asc devkit 索引算子Index示例深度解析 本篇文章围绕 CANN asc devk人工智能深度学习算子库CANNAscendasc-devkit PhiloxRandom 随机数生成样例详解基于 Ascend C 高阶 API 的端到端实现asc devkit PhiloxRandom 随机数生成样例详解基于 Ascend C 高阶 API 的端到端实现 导读 本文以 asc devkit 开源人工智能深度学习算子库CANNAscend上一篇Python Machine Learning 项目技术栈全景从 NumPy、scikit-learn 到 Theano 与 Keras下一篇终极指南如何快速搭建个人照片管理系统创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读本文以 asc-devkit 仓库中examples/01_simd_cpp_api/04_advanced_api/06_index/arange样例为核心系统讲解如何基于昇腾 AI 处理器的 Ascend C 高阶 API——Arange等差序列生成接口实现给定起始值、等差值与长度在片上UB存储中生成指定长度的等差数列并搬运回外部存储GM的完整算子开发流程。读完本文你将掌握 Arange 接口的函数原型、参数约束、底层实现原理与性能特征能够独立完成样例的编译、三种运行模式NPU 运行、CPU 调试、NPU 仿真切换、输入数据与真值生成以及精度对比验证同时理解该样例在仓库中的源码组织结构与配套测试证据。样例概述与功能定位Arange 样例实现的是等差数列生成功能给定起始值firstValue、等差值diffValue和长度count返回一个满足dst[i1] dst[i] diffValue的等差数列。在算子开发实践中等差数列序列常被用作索引生成、序列构造、查表前处理等场景的基础数据属于典型的数据构造类算子。样例将输入起始值与等差值作为标量参数shape 为[1,1]的 float 数据通过 Arange 高阶 API 在片上存储中生成指定长度的等差数列再通过 DataCopy 将结果搬运至外部存储。核函数名为arange_custom。样例输入输出规格名称shape数据类型格式firstGm输入[1,1]floatNDdiffGm输入[1,1]floatNDoutputGm输出[104,1]floatND需要注意输出 shape 为[104,1]而非[100,1]样例固定生成 100 个 float 元素而片上搬运与存储按 32 字节对齐100个 float400 字节向上对齐后实际写入 104 个元素(1008-1)/8*8104每 8 个 float 为 32 字节对齐块。该对齐逻辑在 arange.asc 中通过count_aligned计算并与 gen_data.py 保持严格一致。支持的产品及 CANN 软件版本产品CANN 软件版本Ascend 950PR950DT 系列产品 CANN 9.1.0Atlas A3 系列产品 CANN 9.0.0Atlas A2 系列产品 CANN 9.0.0从仓库中 Arange 接口的产品支持矩阵docs/zh/api/SIMD-API/adv_api/index_compute/Arange.md可以看到更完整的边界Atlas 推理系列产品 AI Core 支持、Vector Core 不支持Atlas 训练系列产品与 Atlas 200I/500 A2 推理产品不支持。因此本样例覆盖的三个产品系列是 Arange 接口在当前仓库中的主要应用对象。目录结构与工程骨架样例根目录结构如下├── arange │ ├── scripts │ │ └── gen_data.py // 输入数据和真值数据生成脚本 │ ├── CMakeLists.txt // 编译工程文件 │ ├── data_utils.h // 数据读入写出函数 │ ├── arange.asc // Ascend C 样例实现 调用样例 │ └── README.md // 样例说明文档编译工程CMakeLists.txtCMakeLists.txt 展示了基于 ASC 语言的最小可运行编译工程cmake_minimum_required(VERSION 3.16) find_package(ASC REQUIRED) project(kernel_samples LANGUAGES ASC CXX) add_executable(demo arange.asc) set(CMAKE_ASC_ARCHITECTURES dav-2201 CACHE STRING CMAKE_ASC_ARCHITECTURES, e.g. dav-2201, dav-3510) target_link_libraries(demo PRIVATE tiling_api register platform unified_dlog dl graph_base) target_compile_options(demo PRIVATE $$COMPILE_LANGUAGE:ASC:--npu-arch${CMAKE_ASC_ARCHITECTURES})关键点说明find_package(ASC REQUIRED)引入 asc-devkit 提供的 ASC 编译器与构建支撑相关配置见 cmake/asc/asc-config.cmake工程默认将CMAKE_ASC_ARCHITECTURES缓存为dav-2201并以--npu-arch编译选项传给 ASC 编译器链接库tiling_api提供 tiling 相关支撑本样例使用 Arange 相关的临时空间查询接口register用于核函数注册platform提供运行时平台能力。数据读写工具data_utils.hdata_utils.h 提供两个二进制文件读写函数ReadFile(filePath, fileSize, buffer, bufferSize)以二进制方式读取文件内部校验文件存在性、是否为普通文件、文件大小不超过缓冲区WriteFile(filePath, buffer, size)以O_RDWR|O_CREAT|O_TRUNC方式创建/覆盖写入二进制文件。host 端代码通过这两个函数读入input/first.bin、input/diff.bin输入并写出output/output.bin结果文件。Arange 高阶 API 原理与参数约束函数原型Arange 高阶 API 的公开接口定义位于 include/adv_api/index/arithprogression.htemplate typename T __aicore__ inline __in_pipe__(S) __out_pipe__(V, S) void Arange( const LocalTensorT dst, const T firstValue, const T diffValue, const int32_t count) { ArithProgressionImpl(dst, firstValue, diffValue, count); }参数说明参数输入/输出描述dst输出目的操作数 LocalTensor大小应大于等于count * sizeof(T)支持的 TPosition 为 VECIN/VECCALC/VECOUTfirstValue输入等差数列的首个元素值diffValue输入等差数列相邻元素之间的差值应大于等于 0count输入等差数列的长度count 0支持的数据类型Ascend 950PR950DT 系列产品int16_t、half、int32_t、float、int64_tAtlas A3 / Atlas A2 系列产品int16_t、half、int32_t、float。该约束在 3510 实现中通过static_assert(SupportTypeT, int16_t, int32_t, half, float, int64_t(), ...)进行编译期检查见 arithprogression_3510_impl.h。格式约束为仅支持 ND 格式。参数校验机制在 CPU 调试ASCENDC_CPU_DEBUG与内核侧检查场景下arithprogression_check_common.h 通过CheckArithProgressionParamsClass完成运行时校验count必须大于 0否则报错 The count parameter cannot be %d, should be greater than 0diffValue必须大于等于 0否则报错 The diffValue parameter cannot be %f, should be greater than or equal to 0数据类型校验仅允许half/float/int16_t/int32_t输出 Tensor 必须位于 UBVECIN/VECOUT/VECCALC并校验count与 dst 长度的匹配关系。配套的校验测试 kernel_arithprogression_check.cpp 验证了合法调用不触发异常ArithProgressionAPICheckTestSuccess、非法调用uint8_t 类型 非正 count/diffValue恰好触发 3 次异常ArithProgressionAPICheckTestFailure两条路径。底层实现原理Arange 内部算法在 Vector 上进行核心思路是先生成基础等差数列再通过向量加法指数级扩充。API 文档docs/zh/api/SIMD-API/adv_api/index_compute/Arange.md将其概括为三步长度 8 以内按 firstValue 与 diffValue 的值用 SetValue 直接生成基础序列最大 8 个元素若目标长度小于 8 则算法结束长度 8 至 64对基础序列使用 Adds 逐块扩充最大循环 7 次扩充至 64 个元素若目标长度小于 64 则算法结束长度 64 以上对 64 元素序列使用 Adds 继续按大步长循环扩充直至达到目标长度。仓库中的公共实现 arithprogression_common_impl.h 与文档描述完全对应GetBaseArithProgressionT()用标量循环dstLocal.SetValue(i, firstValue diffValue * i)生成基础序列并对half做了特化先转 float 计算再转回 half避免精度截断当count BLOCK_NUM一个 32 字节块可容纳的元素数float 下为 8时先生成一块基础序列再通过Adds以GetArithProgressionStepT(diffValue, BLOCK_NUM)为增量逐块复制扩充当count REPEAT_NUM一个 Vector repeat 可处理的元素数float 下为 64时先扩充至 64 元素再以GetArithProgressionStepT(diffValue, REPEAT_NUM)为步长按 repeat 粒度扩充尾部不足一个 repeat 时用SetVectorMask掩码处理当count BLOCK_NUM时直接走标量 SetValue 路径并插入 S-V 事件同步SetFlag/WaitFlagHardEvent::S_V保证标量与向量流水之间的数据依赖。在 Ascend 950PR950DT3510 架构上实现则切换到寄存器级路径 arithprogression_3510_impl.h先用Reg::Arange生成寄存器内的 0..N 索引序列再经Reg::Muls(dstReg, dstReg, diffValue, fullMask)乘上等差值、Reg::Adds(dstReg, dstReg, firstValue, fullMask)加上起始值最后按 repeat 循环Reg::StoreAlign写入 UB并在每个 repeat 后对寄存器序列做一次Reg::Add平移。该路径还处理了int64_t8 字节元素场景此时使用双寄存器 traitRegTraitNumTwo且 repeat 次数减半。由此可见无论哪种实现路径Arange 都是纯计算、无外部数据依赖的接口——它不读任何输入 Tensor仅依据标量参数生成序列因此不需要申请额外 workspace 临时空间。Tiling 配套接口GetArangeMaxMinTmpSizeinclude/adv_api/index/arithprogression_tiling.h 提供了 Arange 的 tiling 辅助接口void GetArangeMaxMinTmpSize(uint32_t maxValue, uint32_t minValue);由于 Arange 内部不需要临时空间其实现arithprogression_tiling_impl.cpp中maxValue与minValue均返回 0。开发者可在 tiling 阶段调用该接口确认无需为 Arange 预留 workspace从而最大化 UB 利用率。注意旧接口GetArithProgressionMaxMinTmpSize已废弃并将在后续版本移除见 docs/zh/api/SIMD-API/adv_api/index_compute/GetArangeMaxMinTmpSize.md请使用新接口。样例实现拆解arange.ascarange.asc 是样例的唯一实现源文件包含 Tiling 定义、核函数与 host 端调用三大部分。第一步Tiling 数据定义与生成namespace optiling { BEGIN_TILING_DATA_DEF(ArithprogressionCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, outSize); END_TILING_DATA_DEF; REGISTER_TILING_DATA_CLASS(ArithprogressionCustom, ArithprogressionCustomTilingData) } // namespace optilingtiling 数据仅包含一个outSize字段用于把输出序列长度从 host 侧传递到核函数。ComputeTiling负责赋值GetTilingBuf将 tiling 数据序列化到缓冲区GenerateTiling封装完整生成流程。host 端通过aclrtMemcpy将该缓冲区拷贝至设备端 tiling 内存后随核函数下发。第二步核函数实现核函数内部封装了KernelArangeT类遵循 Ascend C 典型的Init → ProcessCopyIn → Compute → CopyOut三段式结构Init将起始值、等差值、长度保存为成员变量用SetGlobalBuffer绑定全局输出 Tensor并通过pipe-InitBuffer(outDst, 1, ...)在 UB 上按 32 字节对齐申请输出缓冲Compute核心调用即 Arange 高阶 API// dstLocal: 输出Tensor存储生成的等差数列 // static_castT(firstValue_): 起始值等差数列的第一个元素 // static_castT(diffValue_): 等差值相邻元素之间的差值 // count_: 元素个数生成的等差数列的长度 AscendC::ArangeT(dstLocal, static_castT(firstValue_), static_castT(diffValue_), count_);计算结果经outDst.EnQue入队CopyOutoutDst.DeQue取出结果用AscendC::DataCopy按 32 字节对齐后的块数BLOCK_NUM 32 / sizeof(T)即 float 下每 8 个元素一块将结果从 UB 拷贝至全局内存最后FreeTensor释放。核函数本体arange_custom的签名如下__global__ __vector__ void arange_custom( __gm__ uint8_t* firstGm, __gm__ uint8_t* diffGm, __gm__ uint8_t* dstGm, __gm__ uint8_t* workspace, __gm__ uint8_t* tiling)它先从 tiling 全局内存中解析出outSizeCopyTiling按 4 字节粒度逐字段拷贝再读取firstGm、diffGm指向的标量值初始化KernelArangefloat并执行。核函数声明为__vector__即运行在 Vector 计算单元上与 Arange 的向量实现路径一致。第三步host 端调用与验证host 端main函数完整演示了基于 ACLAscend CANN LiteRuntime 的算子调用流程运行环境初始化aclInit→aclrtSetDevice(0)→aclrtCreateContext→aclrtCreateStream内存分配与数据准备host 侧用aclrtMallocHost设备侧用aclrtMalloc(..., ACL_MEM_MALLOC_HUGE_FIRST)分别申请标量first/diff、输出、workspace 与 tiling 四类内存从./input/first.bin、./input/diff.bin读入输入aclrtMemcpy拷入设备核函数下发使用内核调用符以单 blockNUM_BLOCKS 1方式调用arange_custom随后aclrtSynchronizeStream同步等待完成结果回读与精度对比aclrtMemcpy将输出拷回 hostWriteFile写出./output/output.bin再调用CompareResult与./output/golden.bin逐元素对比详见下文精度对比小节资源回收依次释放设备/宿主内存、销毁 stream、context、reset 设备并aclFinalize。整个调用链展示了从Tiling 生成 → 数据搬入 → 核函数执行 → 结果搬出 → 精度校验的完整闭环可作为编写其他高阶 API 算子的通用模板。数据生成与精度对比机制输入与真值生成gen_data.pygen_data.py 使用 NumPy 生成样例所需的全部二进制文件start np.float32(1.234) difference np.float32(5.678) length np.uint32((100 8 - 1) // 8 * 8) # 100个float32元素按32字节向上对齐实际生成104个元素 stop start difference * (length - 1) golden np.linspace(start, stop, length, dtypenp.float32)first.bin/diff.bin输入起始值与等差值此处取 1.234 与 5.678便于观察序列行为length.bin对齐后的序列长度104与核函数侧count_aligned一致golden.bin使用np.linspace生成的等差序列真值供精度对比使用。精度对比逻辑CompareResult函数读取golden.bin对输出中每个 float 元素计算绝对误差ae与相对误差re ae / abs(b)当ae EPS re EPSEPS 1e-5时判定为错误并打印失配索引与数值全部元素通过后打印CompareResult golden.bin success!host 端最终打印test pass!表示精度对比成功。编译运行指南1. 配置环境变量请根据当前环境上 CANN 开发套件包的安装方式配置环境变量source ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。2. NPU 模式执行在样例根目录下执行如下命令mkdir -p build cd build; # 创建并进入 build 目录 cmake .. -DCMAKE_ASC_ARCHITECTURESdav-2201; make -j; # 编译工程 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行编译生成的可执行程序执行样例3. CPU 调试与 NPU 仿真模式使用 CPU 调试或 NPU 仿真模式时添加-DCMAKE_ASC_RUN_MODEcpu或-DCMAKE_ASC_RUN_MODEsim参数即可cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # CPU调试模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # NPU仿真模式注意切换编译模式前需清理 cmake 缓存可在 build 目录下执行rm CMakeCache.txt后重新 cmake。4. 编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2 系列产品和 Atlas A3 系列产品dav-3510 对应 Ascend 950PR950DT 系列产品其中dav-2201作为默认值同时固化在样例 CMakeLists.txt 的set(CMAKE_ASC_ARCHITECTURES dav-2201 CACHE STRING ...)中。三个架构号与 Arange 接口实现的分支对应dav-2201走公共向量实现arithprogression_common_impl.hdav-3510走寄存器实现arithprogression_3510_impl.h。CPU 调试模式下还会额外启用ASCENDC_CPU_DEBUG宏从而激活 Arange 的参数校验路径见 arithprogression_common_impl.h。5. 执行结果test pass!输出test pass!即表示核函数计算结果与 golden 真值在1e-5误差阈值内完全一致精度对比成功。仓库配套测试与更多参考资料Arange 功能在当前仓库中有成体系的测试支撑功能测试test_operator_arange.cpp以 half 类型对 count15/64/260 三组参数化用例调用Arange并执行DataCopy输出验证不同长度跨度的生成正确性覆盖 8 以内、864、64 以上三条实现路径API 校验测试kernel_arithprogression_check.cpp验证合法调用零异常、非法参数精确触发 3 次异常Tiling 测试在 tests/api/adv_api/tiling/test_tiling.cpp 中通过platfrom_stub_set_chip_version切换芯片版本后调用临时空间查询接口验证返回值为 0。如需进一步深入接口语义可查阅 API 文档 Arange含算法框图与约束说明与 GetArangeMaxMinTmpSize样例的英文版本见 README_en.md。本样例位于examples/01_simd_cpp_api/04_advanced_api/06_index/目录下与同目录的 index 系列样例共同构成 Ascend C 高阶 API 索引类算子的学习素材。小结通过本文可以确认Arange 高阶 API 在 asc-devkit 中的落地形态是一个零临时空间、纯向量计算的序列生成接口接口层提供Arange(dst, firstValue, diffValue, count)的统一签名公共实现通过SetValue 生成基础序列 Adds 指数级扩充最大化向量利用率3510 架构则进一步下沉到寄存器级Reg::Arange/Muls/Adds路径样例arange.asc完整演示了从 tiling 定义、核函数封装到 ACL 运行时调用的全过程配合gen_data.py的数据生成与1e-5精度的 golden 对比构成了可复制、可验证的算子开发范式。无论用于学习高阶 API 用法还是作为自研序列生成算子的起点本样例都值得直接对照阅读与运行。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐昇腾 Ascend C Arange 索引计算算子实战基于 asc-devkit 高阶 API 的等差数列生成与端到端实现解析昇腾 Ascend C Arange 索引计算算子实战基于 asc devkit 高阶 API 的等差数列生成与端到端实现解析 导读 本文围绕 asc dev人工智能深度学习算子库CANNAscend基于 Arange 高阶 API 实现等差数列生成CANN asc-devkit 索引算子Index示例深度解析基于 Arange 高阶 API 实现等差数列生成CANN asc devkit 索引算子Index示例深度解析 本篇文章围绕 CANN asc devk人工智能深度学习算子库CANNAscendasc-devkit PhiloxRandom 随机数生成样例详解基于 Ascend C 高阶 API 的端到端实现asc devkit PhiloxRandom 随机数生成样例详解基于 Ascend C 高阶 API 的端到端实现 导读 本文以 asc devkit 开源人工智能深度学习算子库CANNAscend上一篇Python Machine Learning 项目技术栈全景从 NumPy、scikit-learn 到 Theano 与 Keras下一篇终极指南如何快速搭建个人照片管理系统创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
到O(n)的思维转换)
