【2023 · CANN训练营第一季】昇腾AI高阶课（Ascend C算子开发入门-第一章）学习笔记

昇腾AI高阶课

Ascend C 算子开发入门

Ascend C 介绍

• 使用c/c++作为前端语言的算子开发工具四层接口抽象、并行编程范式、孪生调试等技术
• 使用c/c++愿语编程编程模型屏蔽硬件差异，编程范式提高开发效率多层级API封装，从简单到灵活，兼顾易用高效孪生调试，CPU侧模拟NPU侧行为

核函数

• Ascend C算子设备侧入口可以在核函数中实现算子逻辑的编写
• 是主机侧和设备侧连接的桥梁
• 直接在设备侧执行，需要规定数据访问与计算操作
• 限定符
• 需要使用函数类型限定符

变量类型限定符

• 限定符
• 需要使用函数类型限定符
• 指针入参变量统一为_gm_uint8_t*用户可统一使用uint8_t类型的指针并在使用时转化为实际的指针类型

• 调用核函数
• 使用<<<...>>>语法形式规定核函数的执行配置
• kernel_name<<<block, l2ctrl, stream>>>(argument list)
• blockDim 核数；l2ctrl 保留参数，暂时设置为固定值nullptr；stream 任务队列，用于管理并行
• 核函数调用时异步的，调用结束返回给主机侧
• 需要强制主机侧程序等待执行完毕的API
• aclrtSynchronizeStream

样例展示

• TIK c++模式：CPU模式与NPU模式
• CPU模式：算子功能调试，模拟NPU行为，不需要昇腾设备
•  NPU模式：算子功能/性能调试，可以使用NPU的强大算力进行运算加速
• 内置宏_CCE_KT_TEST_
• #ifdef CPU模式#ifndef NPU 模式

•    

API 接口介绍

• 常用数据结构
• GlobalTensor：用来存放GlobalMemory（外部存储）的全局数据
• template<typename T> class GlobalTensor{     void SetGlobalBuffer(__gm__ T* buffer, uint32_t bufferSize);}
• LocalTensor : 用于存放核上LocalMemory（内部数据）的数据
• template <typename Y> class LocalTensor
• 矢量计算指令接口：能启用AICore中的Vector进行计算
• 降低复杂指令使用难度跨代兼容性保障保留最大灵活度可能
• 三级接口 运算符重载进行连续矢量计算
• 二级接口 可以自定义运算量使用Local Tensor数据结构
• 0级接口 可以进行非连续计算Repeattimes表示迭代次数，一次最多计算256Bytes数据Repeatstride 相邻迭代相同block的地址步长（设置间隔）BlockStride同一迭代内不同block的地址步长Mask参数控制每次迭代内参与计算的元素，通过连续模式与逐比特模式