【2024 鲲鹏&昇腾创新大赛集训营】Ascend C初体验

小铃铛第一次参加华为的线下培训，怀揣着非常激动的心情打开每一天！

今天已经是第二天的培训，正式开始了昇腾的一天学习！

插播一下昨天的一些知识，想直奔主题的小伙伴可以直接下滑哦~

        小铃铛以前对华为的昇腾全栈技术知之甚少，从昨天起开始了解一部分，以下是我的一些个人理解（也可能过于浅显，轻喷，适合小白入门了解）。

        CANN：底层与硬件交互的类似库的部分，Ascend C主要是将基于Python的机器学习模型从Python先编译为二进制文件，并通过C/C++调用，这样可以兼顾代码的复用和C/C++的高性能。是为了对标英伟达的CUDA而出现。

        AI框架：开发了MindSpore，同时支持TensorFlow/Pytorch等第三方框架，在技术专利方面优于英伟达。（但是昨天也听到同学说，在模型文件的读写方面仍有提高空间，并没有完全支持Pytorch中的模型文件类型）

        应用使能：主要说到了ModelArts，是类似于基于Python的Notebook的华为云提供的付费服务，支持MindSpore。用大白话来说，这一部分提供了这些模型与人们交互的平台。

        最后再打个广告：

        奖金丰厚，有意向的小伙伴们冲鸭！

先导知识

下面以一个例子讲解一下数据排布格式的概念。下面是同一张图片的三个不同通道。

R00	R01	R02
R10	R11	R12
R20	R21	R22

G00	G01	G02
G10	G11	G12
G20	G21	G22

B00	B01	B02
B10	B11	B12
B20	B21	B22

        说明：R、G、B分别代表红、绿、蓝三个通道，后面跟的数字可以视为二维数组里的坐标。

        我们分别按照NHWC与NCWH的顺序分别对上图进行遍历，注意遍历是按照从低维到高维进行的，即针对NHWC，遍历顺序是C轴->W轴->H轴->N轴；针对NCHW中，遍历顺序是W轴->H轴->C轴->N轴。

        上图中几个维度大小分别为N：1，H：3，W：3，C：3，不同的排布格式遍历顺序如下所示：

        NHWC：R00,G00,B00,R01,G01,B01,R02,G02,B02……

        NCHW：R00,R01,R02,R10,R11,R12,R20,R21,R22,G00……

书归正文

        本文使用的硬件环境：基于香橙派B130

        由于拍摄效果不太好，我尽量以文字说明。课后记录，且小铃铛不是很明确什么部分存在涉密，我仅展现部分内容，侵权删。

基于核函数的编程

Q：什么是核函数？

        Ascend C的核函数是算子在设备端AI Core上执行的入口，开发者可以在核函数中通过创建算子类对象和调用其成员函数来实现算子的所有功能。

算子：简单点来说就是数学公式，在这个特定语境中指“神经网络中的各种操作，例如矩阵乘法、卷积、池化、激活函数等。”

        核函数的定义需要使用`__global__`和`__aicore__`函数类型限定符，以标识它可以在设备上执行，并且是在AI Core上执行的。核函数的参数指针需要使用`__gm__`变量类型限定符来标识，表明指针指向的是Global Memory上的内存地址。核函数的调用使用特殊的内核调用符`<<<...>>>`，这与普通的C/C++函数调用不同，且核函数的调用是异步的。

小叮当有话说：一定注意__gm__只是一个限定符（前后分别两个下划线），在编程时千万不要和类型紧挨在一起，会导致两个标识符都无法被识别！！！

函数类型限定符

函数类型限定符	执行	调用	备注
__global__	在设备侧执行	用<<<...>>>来调用	必须为void返回值类型，如果有数据必须返回，则写在参数列表的指针位置里
__aicore__	在设备侧执行	仅从设备端调用

变量类型限定符

变量类型限定符	内存空间	意义
__gm__	驻留在Global Memory上	表明该指针变量指向Global Memory上某处内存地址

Global Memory：在昇腾处理器中，每一个核有一个自己的cache和memory，当需要与其他核共享数据时，将数据放在这个区域。它是在整个计算系统中可以被所有处理单元访问的内存区域。以下是Global Memory的一些重要特性和使用方式：

        1. 全局可访问性：Global Memory是所有AI Core都能够访问的内存区域，不同于每个AI Core内部的Local Memory（局部内存）。

        2. 数据存储：Global Memory通常用于存储输入数据、权重、中间结果和最终输出等，这些数据需要在不同的AI Core之间共享。

        3. 内存分配：在Ascend C中，可以使用特定的函数（如`GmAlloc`和`GmFree`）来分配和释放Global Memory。

        4. 数据搬运：由于Global Memory不是直接连接到AI Core的，因此需要通过数据搬运单元（如DMA）将数据从Global Memory移动到AI Core的Local Memory中，以便进行处理。

        5. 性能影响：频繁地从Global Memory读取或写入数据可能会导致性能瓶颈，因此在设计算子时，需要考虑如何有效地管理数据搬运，以减少对Global Memory的访问次数。

        6. 内存带宽：Global Memory的带宽是有限的，因此在设计算子时，应该尽量减少全局内存的访问频率，以避免带宽成为限制性能的因素。

        7. 内存对齐：为了优化内存访问速度，Global Memory中的数据通常需要按照特定的对齐要求进行存储。

        8. 内存管理：开发者需要注意Global Memory的内存管理，避免内存泄漏或访问非法内存地址。

        9. 编程模型：在Ascend C的编程模型中，Global Memory的使用是与SPMD（单程序多数据）模型相结合的，意味着相同的核函数代码会在多个AI Core上并行执行，每个AI Core处理数据的不同部分。

        10. 内存访问模式：Global Memory支持不同的内存访问模式，如连续访问、随机访问等，开发者可以根据具体的计算需求选择合适的访问模式。

         在Ascend C中，核函数通常与流水任务（Stage）结合使用，通过流水任务实现数据的并行处理，提升性能。Ascend C使用Queue队列来管理任务之间的数据通信和同步，而GlobalTensor和LocalTensor作为数据的基本操作单元。

临时变量的表达还需TBuf的参与，声明变量时较复杂。

        核函数开发的基本流程包括环境准备、算子分析、核函数开发、核函数运行验证等步骤。开发者需要分析算子的数学表达式、输入输出以及计算逻辑，并明确所需调用的Ascend C接口。核函数开发完成后，可以通过ICPU_RUN_KF宏在CPU侧进行运行验证，或使用内核调用符在NPU侧进行运行验证。

        Ascend C提供了多层级的API接口，从0级到3级，使得开发者可以根据自己的需求选择合适的API级别。3级接口提供运算符重载，2级接口针对一维tensor的连续计算问题，而0级接口则是最底层的开发接口，提供了最大的灵活性。

这里也就是说，如果希望优化性能，首选0级接口。

那么就会有好奇的小伙伴问：既然0级接口这么好，为什么不能都用0级接口？

A：因为0级接口中的参数即使隐去部分参数，仍有许多需要自己填写。包括但不限于：src_addr，des_addr，block_id等，小伙伴们可以自行查看。

编程简介

        每一个核函数都包含着三个部分：

CopyIn任务：将数据从GlobalTensor中拷贝到LocalTensor中

Compute任务：处理数据（具体需要执行的操作可参考api文档：接口概述-通用概念和约束（必读）-Ascend C API-Ascend C API-算子开发接口-CANN社区版8.0.RC2.alpha003开发文档-昇腾社区接口概述 Ascend C算子采用标准C++语法和一组类库API进行编程，您可以根据自己的需求选择合适的API。Ascend C编程类库API示意图如下所示，Ascend C API的操作数都是Tensor类型：GlobalTensor和LocalTensor；类库API分为基础API和高阶API。 基础API：实现对硬件能力的抽象，开放芯片的能力，保证完备性和兼容性。 高阶API：实现一些常用的https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/80RC2alpha003/apiref/opdevgapi/atlasascendc_api_07_0010.html）

CopyOut任务：将数据从LocalTensor中拷贝到GloablTensor中

Tiling

介又斯麻？

        Tiling，中文通常称为"瓦片化"或"分块"，是一种在并行计算和图形渲染中常用的技术，用于优化数据访问模式和提高计算效率。以下是Tiling的含义和设计原因：

        1. 数据切分：Tiling是将大规模的数据结构（如矩阵、图像或大型张量）切分成更小的、固定大小的块或"tiles"。

        2. 局部性优化：通过Tiling，可以增加数据的空间局部性，使得相关的数据元素在物理存储上彼此靠近，减少内存访问的跳转。

        3. 并行处理：Tiling允许同时对多个tiles进行并行处理，每个处理单元（如CPU核心、GPU线程或AI加速器的AI Core）可以独立地工作在分配给它的tile上。

        4. 减少访问延迟：小的tiles可以更快地加载到处理单元的本地缓存中，减少了从全局内存中读取数据的延迟。

        5. 适应硬件：Tiling可以根据不同硬件的内存层次结构和处理能力来设计，以最大化利用硬件资源。

        在Ascend C编程环境中，一个典型的算子实现包括Host端和Kernel端的代码。Host端通常负责准备数据、调度Kernel执行以及同步操作，而Kernel端则包含实际的计算逻辑。以下是Host端和Kernel端的示例代码：

Host端示例代码（C++）：

#include "acl/acl.h"
#include "ascend_c_datatypes.h"

// 假设Add算子的输入输出数据类型为float16，即half
using half = uint16_t;

int main() {
    // 初始化ACL环境
    aclError ret = acl_init(nullptr);
    if (ret != ACL_ERROR_NONE) {
        printf("Acl init failed, errorCode: %d.\n", ret);
        return -1;
    }

    // 计算资源，这里假设使用0号设备
    aclrtContext context;
    aclrtSetDevice(0);
    ret = aclrtCreateContext(&context, 0);
    if (ret != ACL_ERROR_NONE) {
        printf("Create context failed, errorCode: %d.\n", ret);
        return -1;
    }

    // 分配内存，假设已知输入输出大小为data_size
    size_t data_size = 1024; // 示例大小
    half *input1 = reinterpret_cast<half*>(aclrtMalloc(data_size * sizeof(half), ACL_MEM_MALLOC_NORMAL_ONLY));
    half *input2 = reinterpret_cast<half*>(aclrtMalloc(data_size * sizeof(half), ACL_MEM_MALLOC_NORMAL_ONLY));
    half *output = reinterpret_cast<half*>(aclrtMalloc(data_size * sizeof(half), ACL_MEM_MALLOC_NORMAL_ONLY));

    // 假设数据初始化代码填充input1和input2...

    // 定义Kernel执行的参数
    void *inputs[] = {input1, input2};
    void *outputs[] = {output};
    size_t dims[] = {data_size};

    // 调用Kernel，这里add_custom是Kernel函数名
    add_custom(inputs, outputs, dims);

    // 同步流，确保Kernel执行完成
    aclrtStream stream;
    aclrtCreateStream(&stream);
    aclrtSynchronizeStream(stream);

    // 释放资源
    aclrtFree(input1);
    aclrtFree(input2);
    aclrtFree(output);
    aclrtDestroyStream(stream);
    aclrtDestroyContext(context);
    aclFinalize();

    return 0;
}

Kernel端示例代码： 

// 核函数的声明，使用__global__和__aicore__限定符
extern "C" __global__ __aicore__ void add_custom(
    __attribute__((annotate("input"))) const half* input1,
    __attribute__((annotate("input"))) const half* input2,
    __attribute__((annotate("output"))) half* output,
    size_t size) {
    
    // 每个AI Core处理的数据部分
    size_t block_size = size / USE_CORE_NUM; // 假设USE_CORE_NUM定义了使用的核心数量
    size_t start_idx = GET_BLOCK_IDX * block_size;
    size_t end_idx = start_idx + block_size;

    // 遍历每个AI Core负责的数据范围
    for (size_t i = start_idx; i < end_idx; ++i) {
        output[i] = input1[i] + input2[i];
    }
}

        以上是我由Kimi生成的一份代码，基本含义达到，但是在修改时一般至少包含三个文件：Host端Tiling头文件、Host端Tiling的cpp文件、Kernel端的具体的类似核函数的形式的cpp文件。

        由于实战代码涉及机密，无法给出具体代码，我在昇腾社区搜索了如下文章。

Ascend C教程文档全新来袭！囊括最佳实践、硬核原理、开发技巧...快来get新技能~-技术干货-昇腾社区本文介绍Ascend C教程文档的近期改进亮点。https://www.hiascend.com/zh/developer/techArticles/20240531-1

算子性能

总结

对比

感悟

        小铃铛以前使用C/C++就是简单地用于笔试的算法以及写地图和五子棋的课设，今天才算是真正接触了底层的编程，与我想象的很不同。

        我认为C/C++与Java的区别本质不在于封装等特性，如果说Java中有类的概念，C/C++中也可以方便地定义结构体；如果说Java中有继承的概念，C/C++中结构体中也有类似概念。

        区别在于对于底层编程，我们会更多地对指针、对地址来进行操作，而这些是很有可能导致系统崩溃的隐式的错误。小铃铛在今天下午的编程中，对一个队列（此时分配空间是必要的，不是queue）操作，而没有为其分配空间，从而出现了编译通过而测试不通过的场面。助教老师对我的代码逻辑进行检查，主讲老师对我的代码逐行检查，才发现了错因。而这仅仅是一个非常简单的示例，可想而知底层开发对系统的影响之大。

更多学习资源

[1]Ascend C一站式学习资源：昇腾Ascend C-入门课程-学习资源-算子文档-昇腾社区

[2]AscendCL应用开发指南：快速入门-AscendCL应用开发（C&C++）-应用开发-CANN社区版8.0.RC2.alpha003开发文档-昇腾社区