概述

面向算子开发场景的编程语言Ascend C，原生支持C和C++标准规范，最大化匹配用户开发习惯；通过多层接口抽象、自动并行计算、孪生调试等关键技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署。

使用Ascend C进行自定义算子开发的突出优势有：

• C/C++原语编程
• 编程模型屏蔽硬件差异，编程范式提高开发效率
• 类库API封装，从简单到灵活，兼顾易用与高效
• 孪生调试，CPU侧模拟NPU侧的行为，可优先在CPU侧调试

  环境准备

  基础概念

• 开发环境：指编译开发代码的环境。
• 运行环境：指运行算子、推理程序、训练程序等的环境。运行环境必须带昇腾AI处理器。
• 开发环境与运行环境合设场景：指带昇腾AI处理器的机器既作为开发环境又作为运行环境。此种场景下，代码开发与代码运行在同一台机器上。
• 开发环境与运行环境分设场景：开发环境和运行环境不在同一台机器上，开发者使用带有昇腾AI处理器的机器作为运行环境；使用其他独立机器进行代码开发与编译，作为开发环境。

CANN、达芬奇架构与算子

AI Core的逻辑架构

1. AI Core是昇腾AI处理器的计算核心，采用华为自研的达芬奇架构，通常也叫做DaVinci Core
2. 根据不同的处理器版本，AI Core的计算、存储喝带宽资源有不同的规格
3. 达芬奇架构的主要部分
   1. 计算单元——包含三种基础计算资源（矩阵、向量、标量计算单元）
      1. 矩阵运算单元：L0a（左矩阵）+L0b（右矩阵）——>L0c
   2. 存储单元——AI Core的片上存储单元和相应的数据通路构成了存储系统
   3. 控制单元——整个计算过程提供了指令控制，相当于AI Core的司令部，负责整个AI Core的运行

      AI Core的内部并行计算架构抽象

4. 异步指令流——Scalar运算单元读取指令序列，并把指令发射给对应单元的指令队列，各个计算单元异步的并行执行接收到的指令
5. 同步信号流——指令见可能有依赖关系
6. 计算数据流

   算子是什么

7. 算子对应网络中层或节点的计算逻辑，有输入、输出、计算逻辑。e.g.：Caffe网络、Tensorflow网络都用到了基础单元——算子
8. 数学定义：一个函数空间到函数空间上的映射O：X->X；广义的讲，对任何函数进行某一操作都可称作一个算子
9. 基本概念：
   1. 算子名称（Name）：用于标志网络中的某一个算子，同一网络中的算子名称需要保持唯一e.g.：卷积算子Conv1，Conv2
   2. 算子类型（Type）：同类型的算子的实现逻辑相同
   3. 数据容器（Tensor）：张量（Tensor）是存储算子输入数据、输出数据的容器，而张量描述符（TensoeDesc）是对输入数据、输出数据的描述。数据结构如下
      1. 名称（Name）：用于索引，保持唯一
      2. 形状（Shape）：比如（10，）、（1024，1024）、（2，3，4）、（i1，i2，i3...）
      3. 数据类型（dtype）：FP16...
      4. 数据排布格式（format）：
         1. 深度学习领域，多维数据通过多维数组存储，比如卷积神经网络用四维数组
            1. N：Batch数量，比如图像的数目
            2. H：特征图高度
            3. W：特征图宽度
            4. C：channel
         2. Caffe->NCHW
   4. 算子属性
      1. 轴——表示数据维度
10. 什么是Ascend C算子——面向算子开发场景的编程语言Ascend C，原生支持C和C++标准规范，最大化匹配用户开发习惯；通过多层接口抽象、自动并行计算、孪生调试等关键技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署。
11. 使用Ascend C进行自定义算子开发的突出优势有：
    • C/C++原语编程
    • 编程模型屏蔽硬件差异，编程范式提高开发效率
    • 类库API封装，从简单到灵活，兼顾易用与高效
    • 孪生调试，CPU侧模拟NPU侧的行为，可优先在CPU侧调试
12. 开发自定义算子的需求场景：
    1. 推理场景
    2. 训练场景
    3. 网络调优
13. Device模块

昇腾AscendC快速入门

核函数

1. 核函数（Kernel Function）是Ascend C算子设备侧实现的入口。在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，当核函数被调用时，多个核都执行相同的核函数代码，具有相同的参数，并行执行。
2. Ascend C允许用户使用核函数这种C/C++函数的语法扩展来管理设备端的运行代码，用户在核函数中进行算子类对象的创建和其成员函数的调用，由此实现该算子的所有功能。
3. 使用变量限定符

核函数定义

extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z);

以上是一个核函数声明的例子，编写核函数时需要遵循以下规则。

• 使用函数类型限定符 除了需要按照C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__global__和__aicore__。 使用__global__函数类型限定符来标识它是一个核函数，可以被<<<...>>>调用；使用__aicore__函数类型限定符来标识该核函数在设备端AI Core上执行：

  __global__ __aicore__ void kernel_name(argument list);

  编程中使用到的函数可以分为三类：核函数（device侧执行）、host侧执行函数、device侧执行函数（除核函数之外的）。三者的调用关系如下图所示：

  • host侧执行函数可以调用同类的host执行函数，也就是通用C/C++编程中的函数调用；也可以通过<<<>>>调用核函数。
  • device侧执行函数（除核函数之外的）可以调用同类的device侧执行函数。
  • 核函数可以调用device侧执行函数（除核函数之外的）。

• 使用变量类型限定符

  指针入参变量需要增加变量类型限定符__gm__。表明该指针变量指向Global Memory上某处内存地址。

• 其他规则或建议

  1. 规则：核函数必须具有void返回类型。
  2. 规则：仅支持入参为指针或C/C++内置数据类型（Primitive data types），如：half* s0、float* s1、int32_t c。
  3. 建议：为了统一表达，建议使用GM_ADDR宏来修饰入参，GM_ADDR宏定义如下：

     #define GM_ADDR __gm__ uint8_t*

     使用GM_ADDR修饰入参的样例如下：

     extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)

     这里统一使用uint8_t类型的指针，在后续的使用中需要将其转化为实际的指针类型。

核函数调用

核函数的调用语句是C/C++函数调用语句的一种扩展。实际进行算子开发时，可以通过该方式调用，可以通过接入CANN框架的方式调用，更多细节请参考算子调用。

常见的函数调用方式是如下的形式：

function_name(argument list);

核函数使用内核调用符<<<...>>>这种语法形式，来规定核函数的执行配置：

kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);

内核调用符仅可在NPU侧编译时调用，CPU侧编译无法识别该符号。 执行配置由3个参数决定：

• blockDim，规定了核函数将会在几个核上执行。每个执行该核函数的核会被分配一个逻辑ID，即block_idx，可以在核函数的实现中调用GetBlockIdx来获取block_idx；
• l2ctrl，保留参数，暂时设置为固定值nullptr，开发者无需关注；
• stream，类型为aclrtStream，stream用于维护一些异步操作的执行顺序，确保按照应用程序中的代码调用顺序在device上执行。