Add算子快速入门【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit本示例为入门级演示基于Ascend C SIMD实现Add算子引导您快速完成实践内容涵盖Device端核函数实现、Host端调用以及编译运行的完整流程助您建立整体认知。开始前请参考环境准备安装所需的CANN软件包。以下分别介绍基于C语言API和C Tensor的两种典型实现方式。Add算子功能介绍Add算子的数学表达式为计算逻辑为逐元素完成z x y。Device端代码实现后缀名为*.asc的代码文件可同时包含Host端与Device端代码Device端部分示例如下基于C语言API实现Memory矢量计算示例__vector__ __global__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { asc_init(); constexpr uint32_t block_length TOTAL_LENGTH / NUM_BLOCKS; __gm__ float* x_gm x block_idx * block_length; __gm__ float* y_gm y block_idx * block_length; __gm__ float* z_gm z block_idx * block_length; __ubuf__ float x_local[block_length]; __ubuf__ float y_local[block_length]; __ubuf__ float z_local[block_length]; asc_copy_gm2ub((__ubuf__ void*)x_local, (__gm__ void*)x_gm, block_length * sizeof(float)); asc_copy_gm2ub((__ubuf__ void*)y_local, (__gm__ void*)y_gm, block_length * sizeof(float)); asc_sync(); asc_add(z_local, x_local, y_local, block_length); asc_sync(); asc_copy_ub2gm((__gm__ void*)z_gm, (__ubuf__ void*)z_local, block_length * sizeof(float)); asc_sync(); }[!NOTE] 说明本Memory矢量计算示例支持以下型号Atlas A3训练系列产品/Atlas A3推理系列产品Atlas A2训练系列产品/Atlas A2推理系列产品SIMD算子的Kernel函数需要额外修饰符__vector__修饰符表明该算子仅在向量计算单元上执行。基于C Tensor实现Memory矢量计算示例template uint32_t blockLength __vector__ __global__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { AscendC::InitSocState(); AscendC::GlobalTensorfloat xGm, yGm, zGm; xGm.SetGlobalBuffer(x block_idx * blockLength, blockLength); yGm.SetGlobalBuffer(y block_idx * blockLength, blockLength); zGm.SetGlobalBuffer(z block_idx * blockLength, blockLength); AscendC::LocalMemAllocatorAscendC::Hardware::UB ubAllocator; AscendC::LocalTensorfloat xLocal ubAllocator.Allocfloat, blockLength(); AscendC::LocalTensorfloat yLocal ubAllocator.Allocfloat, blockLength(); AscendC::LocalTensorfloat zLocal ubAllocator.Allocfloat, blockLength(); AscendC::DataCopy(xLocal, xGm, blockLength); AscendC::DataCopy(yLocal, yGm, blockLength); AscendC::PipeBarrierPIPE_ALL(); AscendC::Add(zLocal, xLocal, yLocal, blockLength); AscendC::PipeBarrierPIPE_ALL(); AscendC::DataCopy(zGm, zLocal, blockLength); AscendC::PipeBarrierPIPE_ALL(); }[!NOTE] 说明该样例支持以下型号Ascend 950PR/Ascend 950DTAtlas A3训练系列产品/Atlas A3推理系列产品Atlas A2训练系列产品/Atlas A2推理系列产品SIMD算子的Kernel函数需要额外修饰符__vector__修饰符表明该算子仅在向量计算单元上执行。Host端代码实现Host端通过语法糖调用Device端核函数示例代码如下int32_t main(int argc, char const *argv[]) { ... constexpr uint32_t numBlocks 8; ... // Launch kernel NumBlocks, Block, Dynamic memory, Stream // numBlocks : Number of Blocks // 0 : No dynamic memory required in this sample // nullptr : Use default stream // Example OneCall add kernel which is implemented by SIMD C API add_customnumBlocks, 0(xDevice, yDevice, zDevice); // Example TwoCall add kernel which is implemented by SIMD C Basic API // add_customblockLengthnumBlocks, 0, stream(xDevice, yDevice, zDevice); aclrtSynchronizeDevice(); ... }算子编译与运行CMake配置文件示例cmake_minimum_required(VERSION 3.16) find_package(ASC REQUIRED) project(kernel_samples LANGUAGES ASC CXX) add_executable(c_api_add_example c_api_add.asc ) # # NPU 编译选项配置 # # 说明 # - 需根据实际部署的 NPU 硬件架构选择对应的 npu-arch 参数。 # target_compile_options(c_api_add_example PRIVATE $$COMPILE_LANGUAGE:ASC:--npu-archdav-2201 )编译与执行示例mkdir -p build cd build; # 创建并进入build目录 cmake ..;make -j; # 编译工程 ./c_api_add_example # 运行样例[!NOTE] 说明编译选项--npu-arch用于指定NPU架构版本dav-后面的数字为架构版本号请替换为您实际使用的版本。各AI处理器型号与架构版本的对应关系请查阅AI处理器型号和 __NPU_ARCH__ 的对应关系。此外基于C/C不同层级的编程接口和不同的矢量计算类型Add算子有多种实现方式具体可参考下表API层级矢量计算类型Add算子实例说明SIMD C API基于指针的Memory矢量计算Memory矢量计算Add算子示例同上述C API实现样例贴合C语言开发习惯易于上手SIMD C API基于指针和Reg计算的Reg矢量计算Reg矢量计算Add算子示例贴合C语言开发习惯性能上限更高基础API基于Tensor的Memory矢量计算Memory矢量计算Add算子示例同上述C Tensor实现样例匹配Tensor编程习惯易于上手基础API基于Tensor的Reg矢量计算Reg矢量计算Add算子示例匹配Tensor编程习惯性能上限更高[!NOTE] 说明 Ascend 950PR/Ascend 950DT新一代架构在传统UB缓存体系的基础上开放了寄存器Register可编程能力单个寄存器大小为256B。基于寄存器的矢量计算称为Reg矢量计算而基于传统UB的矢量计算称为Memory矢量计算。若要深入理解Ascend C的SIMD与SIMT编程模型请参阅Ascend C编程模型概述。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考