OrangePi AIpro 香橙派昇腾 Ascend C 算子开发与调用

最新推荐文章于 2024-09-03 18:46:12 发布

西笑生

最新推荐文章于 2024-09-03 18:46:12 发布

阅读量478

点赞数 17

分类专栏：嵌入式深度学习文章标签：昇腾 AscendCL OrangePi AIpro 香橙派

本文链接：https://blog.csdn.net/flyfish1986/article/details/141815774

版权

嵌入式深度学习专栏收录该内容

80 篇文章 25 订阅

订阅专栏

OrangePi AIpro 香橙派昇腾 Ascend C 算子开发与调用

flyfish

开发Ascend C算子的基本流程图

算子类说明

数学表达式

Add算子的数学表达式为：

 z = x + y

Add算子有两个输入：x与y，输出为z。

输入和输出

算子的输入支持的数据类型为half（float16），算子输出的数据类型与输入数据类型相同。
算子输入支持shape（8，2048），输出shape与输入shape相同。
算子输入支持的format为：ND

计算逻辑

计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用计算接口完成两个输入参数相加，得到最终结果，再搬出到外部存储上。
即输入数据需要先搬运进AI Core的内部存储Local Memory，然后使用计算接口完成两个输入参数相加，得到最终结果，再搬出到外部存储Global Memory上。

核函数

核函数命名为add_custom
核函数有3个参数x，y，z；x，y为输入在Global Memory上的内存地址，z为输出在Global Memory上的内存地址

3个基本任务

Add算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。

CopyIn任务负责将Global Memory上的输入Tensor xGm和yGm搬运到Local Memory，分别存储在xLocal、yLocal
Compute任务负责对xLocal、yLocal执行加法操作，计算结果存储在zLocal中
CopyOut任务负责将输出数据从zLocal搬运至Global Memory上的输出Tensor zGm中

调用实现

1. CPU侧运行验证主要通过ICPU_RUN_KF CPU调测宏等CPU调测库提供的接口来完成；
1. NPU侧运行验证主要通过使用ACLRT_LAUNCH_KERNEL内核调用宏来完成。

应用程序通过ASCENDC_CPU_DEBUG 宏区分代码逻辑运行于CPU侧还是NPU侧。

运行样例算子

配置环境变量

默认路径，root用户安装CANN软件包

export ASCEND_INSTALL_PATH=/usr/local/Ascend/ascend-toolkit/latest

默认路径，非root用户安装CANN软件包

export ASCEND_INSTALL_PATH=$HOME/Ascend/ascend-toolkit/latest

指定路径install_path，安装CANN软件包

export ASCEND_INSTALL_PATH=${install_path}/ascend-toolkit/latest

运行命令

bash run.sh -r npu -v Ascend310B4
运行结果

算子代码kernel实现

CopyIn函数实现

使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor。
使用EnQue将LocalTensor放入VecIn的Queue中

Compute函数实现

使用DeQue从VecIn中取出LocalTensor。
使用Ascend C接口Add完成矢量计算。
使用EnQue将计算结果LocalTensor放入到VecOut的Queue中。
使用FreeTensor将释放不再使用的LocalTensor。

CopyOut函数实现

使用DeQue接口从VecOut的Queue中取出LocalTensor。
使用DataCopy接口将LocalTensor拷贝到GlobalTensor上。
使用FreeTensor将不再使用的LocalTensor进行回收。

算子类代码实现

add_custom.cpp

#include "kernel_operator.h"

constexpr int32_t TOTAL_LENGTH = 8 * 2048;                            // total length of data
constexpr int32_t USE_CORE_NUM = 8;                                   // num of core used
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;         // length computed of each core
constexpr int32_t TILE_NUM = 8;                                       // split data into 8 tiles for each core
constexpr int32_t BUFFER_NUM = 2;                                     // tensor num for each queue
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // seperate to 2 parts, due to double buffer

class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)
    {
        xGm.SetGlobalBuffer((__gm__ half *)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
        yGm.SetGlobalBuffer((__gm__ half *)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
        zGm.SetGlobalBuffer((__gm__ half *)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
        pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half));
        pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half));
    }
    __aicore__ inline void Process()
    {
        int32_t loopCount = TILE_NUM * BUFFER_NUM;
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        AscendC::LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
        AscendC::LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();
        AscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
        AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
    }
    __aicore__ inline void Compute(int32_t progress)
    {
        AscendC::LocalTensor<half> xLocal = inQueueX.DeQue<half>();
        AscendC::LocalTensor<half> yLocal = inQueueY.DeQue<half>();
        AscendC::LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
        AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH);
        outQueueZ.EnQue<half>(zLocal);
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress)
    {
        AscendC::LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
        AscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
        outQueueZ.FreeTensor(zLocal);
    }

private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
    AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
    AscendC::GlobalTensor<half> xGm;
    AscendC::GlobalTensor<half> yGm;
    AscendC::GlobalTensor<half> zGm;
};

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

代码解释

constexpr int32_t TOTAL_LENGTH = 8 * 2048;                            // total length of data
constexpr int32_t USE_CORE_NUM = 8;                                   // num of core used
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;         // length computed of each core
constexpr int32_t TILE_NUM = 8;                                       // split data into 8 tiles for each core
constexpr int32_t BUFFER_NUM = 2;                                     // tensor num for each queue
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // seperate to 2 parts, due to double buffer

constexpr 用于声明一个表达式或变量在编译时是常量，而不是在运行时计算

 TOTAL_LENGTH: 数据总长度 8 * 2048。
 USE_CORE_NUM: 使用的核心数 8。
 BLOCK_LENGTH: 每个核心处理的数据长度 2048。
 TILE_NUM: 将数据分成的小块数量8，供每个核心处理。
 BUFFER_NUM: 每个队列中的张量数量2（由于使用双缓冲）。
 TILE_LENGTH: 由于双缓冲机制，进一步分成两部分。TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; 
2048/8/2

使用多核并行计算，即把数据进行分片，分配到多个核上进行处理。Ascend C核函数是在一个核上的处理函数，所以只处理部分数据。分配方案是：数据整体长度TOTAL_LENGTH为8* 2048，平均分配到8个核上运行，每个核上处理的数据大小BLOCK_LENGTH为2048。核函数只关注长度为BLOCK_LENGTH的数据应该如何处理。
即上面的代码

constexpr int32_t TOTAL_LENGTH = 8 * 2048;                            // total length of data
constexpr int32_t USE_CORE_NUM = 8;                                   // num of core used
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;         // length computed of each core

TILE_NUM（每个核上总计算数据分块个数）、BLOCK_LENGTH（每个核上总计算数据大小）、TILE_LENGTH（每个分块大小）等是固定的数值。

算子程序被调用时启动N个实例运行每个运行实例称为block，他们之间唯一区别是运行时实例ID不同:block idx的值不同
所有block被调用时，都执行相同的代码，有相同的参数
block概念类似线程(进程)，block idx类似ThreadID(进程ID)
算子程序需要使用block idx来进行数据并行计算切分。

每个核上处理的数据地址需要在起始地址上增加GetBlockIdx()*BLOCK_LENGTH（每个block处理的数据长度）的偏移来获取。这样也就实现了多核并行计算的数据切分。

以输入x为例，x + BLOCK_LENGTH * GetBlockIdx()即为单核处理程序中x在Global Memory上的内存偏移地址，获取偏移地址后，使用GlobalTensor类的SetGlobalBuffer接口设定该核上Global Memory的起始地址以及长度

xGm.SetGlobalBuffer((__gm__ half *)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
yGm.SetGlobalBuffer((__gm__ half *)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
zGm.SetGlobalBuffer((__gm__ half *)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);

在这里插入图片描述单核上的处理数据，可以进行数据切块（Tiling），将数据切分成8块（并不意味着8块就是性能最优）。切分后的每个数据块再次切分成2块，即可开启double buffer，实现流水线之间的并行。

这样单核上的数据（2048个数）被切分成16块，每块TILE_LENGTH（128）个数据。Pipe为inQueueX分配了两块大小为TILE_LENGTH * sizeof(half)个字节的内存块，每个内存块能容纳TILE_LENGTH（128）个half类型数据。在这里插入图片描述即上面的代码

 TILE_NUM: 将数据分成的小块数量8，供每个核心处理。
 BUFFER_NUM: 每个队列中的张量数量2（由于使用双缓冲）。
 TILE_LENGTH: 由于双缓冲机制，进一步分成两部分。TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; 
2048/8/2 = 2048/16 = 128

语法

const int x = 10; // x 是一个不可修改的常量，但它的值可能在运行时才能确定
constexpr int y = 20; // y 是一个编译时常量，编译器可以在编译时直接使用它的值

在以上代码中，y 的值可以在编译时确定，所以编译器可以将其直接插入到程序的指令中，减少运行时的开销。而 x 只保证了其不可修改，但不强制要求编译时确定其值。

CopyIn任务:将Global Memory上的输入Tensor xGm和yGm搬运至Local Memory，分别存储在xLocal, yLocal
Compute任务:对xLocal,yLocal执行加法操作，计算结果存储在zLocal中
CopyOut任务:将输出数据从zLocal搬运至GlobalMemory上的输出Tensor zGm中
CopyIn，Compute任务间通过VECIN队列inQueueX，inQueueY进行通信和同步
Compute，CopyOut任务间通过VECOUT队列outQueueZ进行通信和同步
pipe内存管理对象对任务间交互使用到的内存、临时变量使用到的内存统一进行管理
在这里插入图片描述

简化算子类

class KernelAdd {
public:
    __aicore__ inline KernelAdd(){}
    // 初始化函数，完成内存初始化相关操作
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z){}
    // 核心处理函数，实现算子逻辑，调用私有成员函数CopyIn、Compute、CopyOut完成矢量算子的三级流水操作
    __aicore__ inline void Process(){}

private:
    // 搬入函数，完成CopyIn阶段的处理，被核心Process函数调用
    __aicore__ inline void CopyIn(int32_t progress){}
    // 计算函数，完成Compute阶段的处理，被核心Process函数调用
    __aicore__ inline void Compute(int32_t progress){}
    // 搬出函数，完成CopyOut阶段的处理，被核心Process函数调用
    __aicore__ inline void CopyOut(int32_t progress){}

private:
    AscendC::TPipe pipe;  //Pipe内存管理对象
    AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;  //输入数据Queue队列管理对象，QuePosition为VECIN
    AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueZ;  //输出数据Queue队列管理对象，QuePosition为VECOUT
    AscendC::GlobalTensor<half> xGm;  //管理输入输出Global Memory内存地址的对象，其中xGm, yGm为输入，zGm为输出
    AscendC::GlobalTensor<half> yGm;
    AscendC::GlobalTensor<half> zGm;
};

在这里插入图片描述

最后一段

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

使用__global__函数类型限定符来标识它是一个核函数，可以被<<<>>>调用；
使用__aicore__函数类型限定符来标识该核函数在设备端AI Core上执行。指针入参变量需要增加变量类型限定符__gm__，表明该指针变量指向Global Memory上某处内存地址。为了统一表达，使用GM_ADDR宏来修饰入参，GM_ADDR宏定义如下

#define GM_ADDR __gm__ uint8_t*

在这里插入图片描述算子类的Init函数，完成内存初始化相关工作
Process函数完成算子实现的核心逻辑

测试代码

main.cpp

#include "data_utils.h"
#ifndef ASCENDC_CPU_DEBUG
#include "acl/acl.h"
#include "aclrtlaunch_add_custom.h"
#else
#include "tikicpulib.h"
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z);
#endif

int32_t main(int32_t argc, char *argv[])
{
    uint32_t blockDim = 8;
    size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);
    size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);

#ifdef ASCENDC_CPU_DEBUG
    uint8_t *x = (uint8_t *)AscendC::GmAlloc(inputByteSize);
    uint8_t *y = (uint8_t *)AscendC::GmAlloc(inputByteSize);
    uint8_t *z = (uint8_t *)AscendC::GmAlloc(outputByteSize);

    ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);
    ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize);

    AscendC::SetKernelMode(KernelMode::AIV_MODE);
    ICPU_RUN_KF(add_custom, blockDim, x, y, z); // use this macro for cpu debug

    WriteFile("./output/output_z.bin", z, outputByteSize);

    AscendC::GmFree((void *)x);
    AscendC::GmFree((void *)y);
    AscendC::GmFree((void *)z);
#else
    CHECK_ACL(aclInit(nullptr));
    int32_t deviceId = 0;
    CHECK_ACL(aclrtSetDevice(deviceId));
    aclrtStream stream = nullptr;
    CHECK_ACL(aclrtCreateStream(&stream));

    uint8_t *xHost, *yHost, *zHost;
    uint8_t *xDevice, *yDevice, *zDevice;

    CHECK_ACL(aclrtMallocHost((void **)(&xHost), inputByteSize));
    CHECK_ACL(aclrtMallocHost((void **)(&yHost), inputByteSize));
    CHECK_ACL(aclrtMallocHost((void **)(&zHost), outputByteSize));
    CHECK_ACL(aclrtMalloc((void **)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc((void **)&yDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc((void **)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

    ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);
    ReadFile("./input/input_y.bin", inputByteSize, yHost, inputByteSize);

    CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
    CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));

    ACLRT_LAUNCH_KERNEL(add_custom)(blockDim, stream, xDevice, yDevice, zDevice);
    CHECK_ACL(aclrtSynchronizeStream(stream));

    CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));
    WriteFile("./output/output_z.bin", zHost, outputByteSize);

    CHECK_ACL(aclrtFree(xDevice));
    CHECK_ACL(aclrtFree(yDevice));
    CHECK_ACL(aclrtFree(zDevice));
    CHECK_ACL(aclrtFreeHost(xHost));
    CHECK_ACL(aclrtFreeHost(yHost));
    CHECK_ACL(aclrtFreeHost(zHost));

    CHECK_ACL(aclrtDestroyStream(stream));
    CHECK_ACL(aclrtResetDevice(deviceId));
    CHECK_ACL(aclFinalize());
#endif
    return 0;
}

代码解释

main函数中通过对ASCENDC_CPU_DEBUG宏定义的判断来区分CPU和NPU侧的运行程序

#include "data_utils.h"
#ifndef ASCENDC_CPU_DEBUG
#include "acl/acl.h"
extern void add_custom_do(uint32_t coreDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z);
#else
#include "tikicpulib.h"
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z);
#endif

int32_t main(int32_t argc, char* argv[])
{
    size_t inputByteSize = 8 * 2048 * sizeof(uint16_t);  // uint16_t represent half
    size_t outputByteSize = 8 * 2048 * sizeof(uint16_t);  // uint16_t represent half
    uint32_t blockDim = 8;

#ifdef ASCENDC_CPU_DEBUG
    // 用于CPU调试的调用程序
    
#else
    // NPU侧运行算子的调用程序

#endif
    return 0;
}

语法解释

#define CHECK_ACL(x)                                                                        \
    do {                                                                                    \
        aclError __ret = x;                                                                 \
        if (__ret != ACL_ERROR_NONE) {                                                      \
            std::cerr << __FILE__ << ":" << __LINE__ << " aclError:" << __ret << std::endl; \
        }                                                                                   \
    } while (0);

定义了一个名为 CHECK_ACL 的宏，用于检查 aclError 类型的错误并在出现错误时输出错误信息。

do { ... } while (0) 结构 ：
这个结构是为了确保宏在使用时的语法行为正确。例如，在使用 if 语句包裹宏时，可以避免由于缺少大括号导致的语法错误。这是常见的宏定义技巧，do { ... } while (0) 使得宏的定义相当于一个单一语句。

尽管这是一个循环结构，但由于条件为 0，循环体仅执行一次，相当于一个代码块被包裹在一个单一的语句中。
举个例子 ：假设没有使用 do { ... } while (0)，直接定义宏：

#define CHECK_ACL(x) \
    aclError __ret = x; \
    if (__ret != ACL_ERROR_NONE) { \
        std::cerr << __FILE__ << ":" << __LINE__ << " aclError:" << __ret << std::endl; \
    }

使用宏的代码 ：

if (condition)
    CHECK_ACL(aclInit());
else
    handleError();

宏展开后的代码 ：

if (condition)
    aclError __ret = aclInit();
    if (__ret != ACL_ERROR_NONE) {
        std::cerr << __FILE__ << ":" << __LINE__ << " aclError:" << __ret << std::endl;
    };
else
    handleError();

产生的问题 ：if 语句只控制第一行（aclError __ret = aclInit();），而第二个 if 语句会被视为独立的，这导致 else 与第一个 if 不匹配，引发语法错误或逻辑错误。

算子核函数NPU侧运行验证

资源初始化

  CHECK_ACL(aclInit(nullptr));

调用aclInit接口初始化AscendCL：

使用AscendCL接口开发应用时，必须先初始化AscendCL ，否则可能会导致后续系统内部资源初始化出错，进而导致其它业务异常。

Device，用于指定计算设备。

Device的生命周期源于首次调用aclrtSetDevice接口。
每次调用aclrtSetDevice接口，系统会进行引用计数加1；调用aclrtResetdevice接口，系统会进行引用计数减1。
当引用计数减为零时，在本进程中Device上的资源不可用。
Context，在Device下，一个Context一定属于一个唯一的Device。

int32_t deviceId = 0;
CHECK_ACL(aclrtSetDevice(deviceId));
CHECK_ACL(aclrtResetDevice(deviceId));

Context 隐式创建

隐式创建的Context（即默认Context），生命周期始于调用aclrtSetDevice接口，终结于调用aclrtResetdevice接口使引用计数为零时。隐式Context只会被创建一次，调用aclrtSetDevice接口重复指定同一个Device，只增加隐式创建的Context的引用计数。

Stream 显式创建

Stream，是Device上的执行流，在同一个stream中的任务执行严格保序。

Stream分隐式创建和显式创建。
每个Context都会包含一个默认Stream，这个属于隐式创建，隐式创建的stream生命周期同归属的Context。
用户可以显式创建stream，显式创建的stream生命周期始于调用aclrtCreateStream，终结于调用aclrtDestroyStream接口。显式创建的stream归属的Context被销毁或生命周期结束后，会影响该stream的使用，虽然此时stream没有被销毁，但不可再用。

aclrtStream stream = nullptr;
CHECK_ACL(aclrtCreateStream(&stream));
CHECK_ACL(aclrtDestroyStream(stream));

数据传输

Host上的内存

可以使用AscendCL提供的aclrtMallocHost接口申请内存
aclrtMallocHost会尝试申请物理地址连续的内存，后续在Host与Device数据交互时性能更优。调用aclrtMallocHost接口后、使用内存前，建议先调用aclrtMemset接口初始化内存，清除内存中的随机数。

Device上的内存

使用AscendCL提供的aclrtMalloc接口申请内存

CHECK_ACL(aclrtMallocHost((void **)(&xHost), inputByteSize));
CHECK_ACL(aclrtMallocHost((void **)(&yHost), inputByteSize));
CHECK_ACL(aclrtMallocHost((void **)(&zHost), outputByteSize));
CHECK_ACL(aclrtMalloc((void **)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMalloc((void **)&yDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMalloc((void **)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

通过内存复制实现数据传输。
数据传输可以通过内存复制的方式实现，分为同步内存复制、异步内存复制：

同步内存复制：调用aclrtMemcpy接口。
异步内存复制：调用aclrtMemcpyAsync接口，再调用aclrtSynchronizeStream接口实现Stream内任务的同步等待。
对于Host内的数据传输、Device内的数据传输、Host与Device之间的数据传输，可以调用内存复制的接口实现，也可以直接通过指针传递数据

CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));

ACLRT_LAUNCH_KERNEL(add_custom)(blockDim, stream, xDevice, yDevice, zDevice);
CHECK_ACL(aclrtSynchronizeStream(stream));

CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));

NPU侧调用方式 - Kernel Launch方式

ACLRT_LAUNCH_KERNEL调用接口的使用方法如下：

ACLRT_LAUNCH_KERNEL(kernel_name)(blockDim, stream, argument list);

kernel_name：算子核函数的名称。
blockDim：规定了核函数将会在几个核上执行。每个执行该核函数的核会被分配一个逻辑ID，即block_idx，可以在核函数的实现中调用GetBlockIdx来获取block_idx。
stream，类型为aclrtStream，stream用于维护一些异步操作的执行顺序，确保按照应用程序中的代码调用顺序在Device上执行。
argument list：参数列表，与核函数的参数列表保持一致。

ACLRT_LAUNCH_KERNEL(add_custom)(blockDim, stream, xDevice, yDevice, zDevice);

解释

// blockDim设置为8表示在8个核上调用了add_custom核函数，每个核都会独立且并行地执行该核函数，该核函数的参数列表为x，y，z。
ACLRT_LAUNCH_KERNEL(add_custom)(8, stream, x, y, z)

blockDim，规定了核函数将会在几个核上执行。每个执行该核函数的核会被分配一个逻辑ID，即block_idx，可以在核函数的实现中调用GetBlockIdx来获取block_idx；
BlockDim是逻辑核的概念，取值范围为[1,65535]。为了充分利用硬件资源，一般设置为物理核的核数或其倍数。对于耦合架构和分离架构，OrangePi AIpro 香橙派是耦合架构。

耦合架构：由于其Vector、Cube单元是集成在一起的，BlockDim用于设置启动多个AICore核实例执行，不区分Vector、Cube。AI Core的核数可以通过GetCoreNumAiv或者GetCoreNumAic获取。

内核调用符方式，核函数可以使用内核调用符<<<…>>>这种语法形式

// blockDim设置为8表示在8个核上调用了add_custom核函数，每个核都会独立且并行地执行该核函数，该核函数的参数列表为x，y，z。
add_custom<<<8, nullptr, stream>>>(x, y, z);

内存管理

使用aclrtMalloc接口申请的内存，需要通过aclrtFree接口释放内存。
频繁调用aclrtMalloc接口申请内存、调用aclrtFree接口释放内存，会损耗性能，建议用户提前做内存预先分配或二次管理，避免频繁申请/释放内存。

使用aclrtMallocHost接口申请的内存，需要通过aclrtFreeHost接口释放内存。
使用aclrtMallocHost接口分配过多的锁页内存，将导致操作系统用于分页的物理内存变少，从而降低系统整体的性能

AscendCL去初始化

CHECK_ACL(aclFinalize());

调用aclFinalize接口实现AscendCL去初始化。

西笑生

关注

17
点赞
踩
9

收藏

觉得还不错? 一键收藏
打赏
0
评论
OrangePi AIpro 香橙派昇腾 Ascend C 算子开发与调用

flyfish
复制链接

扫一扫

专栏目录

OrangePi AIpro 香橙派 昇腾 Ascend C 算子开发 与 调用

OrangePi AIpro 香橙派 昇腾 Ascend C 算子开发 与 调用

算子类说明

数学表达式

输入和输出

计算逻辑

核函数

3个基本任务

调用实现

运行样例算子

运行命令

算子代码kernel实现

CopyIn函数实现

Compute函数实现

CopyOut函数实现

算子类代码实现

代码解释

语法

简化算子类

测试代码

代码解释

语法解释

算子核函数NPU侧运行验证

资源初始化

Device，用于指定计算设备。

Context 隐式创建

Stream 显式创建

数据传输

Host上的内存

Device上的内存

NPU侧调用方式 - Kernel Launch方式

内存管理

AscendCL去初始化

OrangePi AIpro 香橙派昇腾 Ascend C 算子开发与调用

OrangePi AIpro 香橙派昇腾 Ascend C 算子开发与调用