AscendC从入门到精通系列(二)基于Kernel直调开发AscendC算子

简介: 本文介绍了AscendC算子的开发流程,包括核函数开发、算子类定义及其实现、核函数的CPU和NPU侧运行验证。通过具体示例`add_custom.cpp`,详细展示了如何使用Ascend C完成算子核函数的定义、初始化、数据搬运和计算过程,并提供了完整的CPU和NPU侧调用程序代码,帮助开发者理解和实践AscendC算子的开发。

本次主要讨论下AscendC算子的开发流程,基于Kernel直调工程的算子开发。

1 AscendC算子开发的基本流程

使用Ascend C完成Add算子核函数开发;
使用ICPU_RUN_KF CPU调测宏完成算子核函数CPU侧运行验证;
使用<<<>>>内核调用符完成算子核函数NPU侧运行验证。
在正式的开发之前,还需要先完成环境准备和算子分析工作,开发Ascend C算子的基本流程如下图所示:
image.png

2 核函数开发

本次以add_custom.cpp作为参考用例。Gitee也有对应工程和完整代码。
operator/AddCustomSample/KernelLaunch/AddKernelInvocationNeo · Ascend/samples - 码云 - 开源中国 (gitee.com)

2.1 核函数定义

首先要根据核函数定义 核函数-编程模型-Ascend C算子开发-算子开发-开发指南-CANN社区版8.0.RC3.alpha003开发文档-昇腾社区 (hiascend.com) 的规则进行核函数的定义,并在核函数中调用算子类的Init和Process函数。

// 给CPU调用
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();
}


// 给NPU调用
#ifndef ASCENDC_CPU_DEBUG
void add_custom_do(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, uint8_t *z)
{
   
    add_custom<<<blockDim, nullptr, stream>>>(x, y, z);
}
#endif

2.2 算子类定义

根据矢量编程范式实现算子类,本样例中定义KernelAdd算子类,其具体成员如下:

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;
};

核函数调用关系图
image.png

2.3 实现Init,CopyIn,Compute,CopyOut这个4个关键函数

Init函数初始化输入资源

__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));
    }
Process函数中通过如下方式调用这三个:
    __aicore__ inline void Process()
    {
   
        // loop count need to be doubled, due to double buffer
        constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;
        // tiling strategy, pipeline parallel
        for (int32_t i = 0; i < loopCount; i++) {
   
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

CopyIn函数中通过如下方式调用这三个:
1、使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor。
2、使用EnQue将LocalTensor放入VecIn的Queue中。

__aicore__ inline void CopyIn(int32_t progress)
    {
   
        // alloc tensor from queue memory
        AscendC::LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
        AscendC::LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();
        // copy progress_th tile from global tensor to local tensor
        AscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
        AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
        // enque input tensors to VECIN queue
        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
    }

Compute函数实现。
1、使用DeQue从VecIn中取出LocalTensor。
2、使用Ascend C接口Add完成矢量计算。
3、使用EnQue将计算结果LocalTensor放入到VecOut的Queue中。
4、使用FreeTensor将释放不再使用的LocalTensor。

__aicore__ inline void Compute(int32_t progress)
{
   
    // deque input tensors from VECIN queue
    AscendC::LocalTensor<half> xLocal = inQueueX.DeQue<half>();
    AscendC::LocalTensor<half> yLocal = inQueueY.DeQue<half>();
    AscendC::LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
    // call Add instr for computation
    AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH);
    // enque the output tensor to VECOUT queue
    outQueueZ.EnQue<half>(zLocal);
    // free input tensors for reuse
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}

CopyOut函数实现。
1、使用DeQue接口从VecOut的Queue中取出LocalTensor。
2、使用DataCopy接口将LocalTensor拷贝到GlobalTensor上。
3、使用FreeTensor将不再使用的LocalTensor进行回收。

 __aicore__ inline void CopyOut(int32_t progress)
{
   
    // deque output tensor from VECOUT queue
    AscendC::LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
    // copy progress_th tile from local tensor to global tensor
    AscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
    // free output tensor for reuse
    outQueueZ.FreeTensor(zLocal);
}

3 核函数的运行验证

异构计算架构中,NPU(kernel侧)与CPU(host侧)是协同工作的,完成了kernel侧核函数开发后,即可编写host侧的核函数调用程序,实现从host侧的APP程序调用算子,执行计算过程。

3.1 编写CPU侧调用程序

image.png

 // 使用GmAlloc分配共享内存,并进行数据初始化
    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);
    // 调用ICPU_RUN_KF调测宏,完成核函数CPU侧的调用
    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);
    // 调用GmFree释放申请的资源
    AscendC::GmFree((void *)x);
    AscendC::GmFree((void *)y);
    AscendC::GmFree((void *)z);

3.2 编写NPU侧运行算子的调用程序

image.png

  // AscendCL初始化
    CHECK_ACL(aclInit(nullptr));
    // 运行管理资源申请
    int32_t deviceId = 0;
    CHECK_ACL(aclrtSetDevice(deviceId));
    aclrtStream stream = nullptr;
    CHECK_ACL(aclrtCreateStream(&stream));
    // 分配Host内存
    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));
    // 分配Device内存
    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));
    // Host内存初始化
    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));
    // 用内核调用符<<<>>>调用核函数完成指定的运算,add_custom_do中封装了<<<>>>调用
    add_custom_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice);
    CHECK_ACL(aclrtSynchronizeStream(stream));
    // 将Device上的运算结果拷贝回Host
    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));
    // AscendCL去初始化
    CHECK_ACL(aclrtDestroyStream(stream));
    CHECK_ACL(aclrtResetDevice(deviceId));
    CHECK_ACL(aclFinalize());

3.3 完整main.cpp

/**
 * @file main.cpp
 *
 * Copyright (C) 2024. Huawei Technologies Co., Ltd. All rights reserved.
 *
 * This program is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
 */
#include "data_utils.h"
#ifndef ASCENDC_CPU_DEBUG
#include "acl/acl.h"
extern void add_custom_do(uint32_t blockDim, 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[])
{
   
    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));

    add_custom_do(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;
}

整体运行起来,请参考operator/AddCustomSample/KernelLaunch/AddKernelInvocationNeo · Ascend/samples - 码云 - 开源中国 (gitee.com)

相关文章
|
10月前
|
人工智能 算法 编译器
AscendC从入门到精通系列(一)初步感知AscendC
Ascend C是CANN推出的一种编程语言,专为算子开发设计,支持C/C++标准,旨在提高开发效率与运行性能。通过Ascend C,开发者能在昇腾AI处理器上高效实现自定义算法。本文档介绍了如何使用Ascend C编写和运行一个简单的“Hello World”程序,包括核函数的编写、主程序调用及CMake配置,展示了Ascend C的基本使用流程。
|
资源调度
一天掌握latex论文编辑,从标题作者,段落,数学公式,图片,图表,到参考文献全流程
一天掌握latex论文编辑,从标题作者,段落,数学公式,图片,图表,到参考文献全流程
1246 0
|
6月前
|
容器
vllm+vllm-ascend本地部署QwQ-32B
本指南介绍如何下载、安装和启动基于Ascend的vLLM模型。首先,可通过华为镜像或Hugging Face下载预训练模型;其次,安装vllm-ascend,支持通过基础镜像(如`quay.io/ascend/vllm-ascend:v0.7.3-dev`)或源码编译方式完成;最后,使用OpenAI兼容接口启动模型,例如运行`vllm serve`命令,设置模型路径、并行规模等参数。适用于大模型推理场景,需注意显存需求(如QwQ-32B需70G以上)。
2377 17
|
10月前
|
存储 人工智能 JSON
AscendC从入门到精通系列(三)基于自定义算子工程开发AscendC算子
本文介绍了基于Ascend C的自定义算子开发流程,涵盖从工程创建、代码编写、编译部署到运行验证的全过程。以动态shape的AddCustom算子为例,详细描述了如何利用CANN提供的工具msOpGen生成开发工程,实现算子核函数与host侧代码,以及如何编译、部署和测试自定义算子。
|
11月前
|
人工智能 Apache 流计算
参与Flink社区活动,免费赢取FFA大会两日通票~
Flink Forward Asia 2024 将于 11 月 29-30 日在上海举行,庆祝 Apache Flink 诞生十周年。大会将回顾 Flink 的技术成就,展望未来十年的发展,并介绍 Flink 2.0 版本。通过三种参与方式,您有机会免费赢取大会两日通票和 Flink 专属周边。
517 13
参与Flink社区活动,免费赢取FFA大会两日通票~
|
10月前
|
机器学习/深度学习 存储 缓存
ATB概念之:算子tiling
算子 tiling 是一种优化技术,用于提高大规模张量运算的计算效率。它通过将大任务分解为小块,优化内存使用、支持并行计算,并防止内存溢出。在ATB中,tiling data指kernel的分片参数,用于指导计算。ATB提供了三种 tiling data 搬移策略:整体搬移、多stream搬移及随kernel下发搬移,旨在优化内存拷贝任务,提高计算效率。
|
10月前
|
人工智能 算法 PyTorch
ATB是什么?
ATB加速库专为华为Ascend AI处理器设计,针对Transformer模型的训练和推理进行了深度优化。它通过算法、硬件和软件层面的优化,大幅提升模型性能,降低能耗与成本。ATB支持PyTorch、MindSpore等多种框架,提供高效的基础算子及图算子技术,适用于各种应用场景。其软件架构主要包括基础Operation、Plugin机制和Graph Frame三部分,通过优化算子计算和数据传输,实现性能的显著提升。
|
10月前
|
域名解析 运维 网络协议
网络诊断指南:网络故障排查步骤与技巧
网络诊断指南:网络故障排查步骤与技巧
3233 7
|
11月前
|
传感器 前端开发 Android开发
在 Flutter 开发中,插件开发与集成至关重要,它能扩展应用功能,满足复杂业务需求
在 Flutter 开发中,插件开发与集成至关重要,它能扩展应用功能,满足复杂业务需求。本文深入探讨了插件开发的基本概念、流程、集成方法、常见类型及开发实例,如相机插件的开发步骤,同时强调了版本兼容性、性能优化等注意事项,并展望了插件开发的未来趋势。
296 2
|
Python
python实现:旋转矩阵转换为四元数
python实现:旋转矩阵转换为四元数
370 0