So simple and Innocent

SOLECA_

209人浏览 · 2026-08-01 01:46:52

SOLECA_ · 2026-08-01 01:46:52 发布

今天是古法编程（）
写完AddCustom、SubCustom再补一点DivCustom

Hi Ascend

#include "acl/acl.h"
#include "kernel_operator.h"
using namespace AscendC;
__global__ __aicore__ void hi_ascend(){
    printf("Block[%lu/%lu]: Hi Ascend\n",GetBlockIdx(),GetBlockNum());//long unsigned
}
int32_t main(int argc ,char const *argv[]){
    aclInit(nullptr);
    int32_t device=0;
    aclrtSetDevice(deviceId);
    aclrtStream stream=nullptr;
    aclrtCreateStream(&stream);
    constexpr uint32_t blockDim=4;
    hi_ascend<<<blockDim,nullptr,stream>>>();
    aclrtSynchronizeStream(stream);
    aclrtDestoryStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();
    return 0;
}

教程 - 逐元素加法

数据类型：输入与输出均为 float 类型。
数据形状（Shape）：输入Shape为 (8, 2048)，输出Shape与输入相同。

首先是函数库

#include <cstdint>
#include <iostream>
#include <vector>
#include <algorithm>
#include <iterator>
#include "acl/acl.h" //host侧调用
#include "kernel_operator.h" //kernel侧调用
using namespace AscendC;
using namespace std;

BufferNum

constexpr uint32_t BUFFER_NUM=2;
constexpr uint32_t QUEUE_DEPTH=2;

tiling结构体

struct AddCustomTilingData{
    uint32_t totalLength;
    uint32_t tileNum;
};

KernelAdd类

class KernelAdd{
public:
    __aicore__ inline KernelAdd(){}
    __aicore__ inline void Init(GM_ADDR x,GM_ADDR y,GM_ADDR z,uint32_t totalLength,uint32_t tileNum);
    __aicore__ inline void Process();
private:
private:
    Tpipe pipe;//TPipe内存管理对象
    TQue<TPosition::VECIN,QUEUE_DEPTH>  inQueueX,inQueueY;
    TQue<TPosition::VECOUT,QUEUE_DEPTH> outQueueZ;
    AscendC::GlobalTensor<float> xGm;
    AscendC::GlobalTensor<float> yGm;
    AscendC::GlobalTensor<float> zGm;
    uint32_t blockLength;//单个AICore需要处理数据的总长度
    uint32_t tileNum; //核内分块的数量
    uint32_t tileLength;//核内分块的数据长度
}

初始化Init函数

__aicore__ inline void KernelAdd::Init(GM_ADDR x,GM_ADDR y,GM_ADDR z,uint32_t totalLength,uint32_t tileNum){
    this->blockLength=totalLength/AscendC::GetBlockNum();
    this->tileNum=tileNum;
    this->tileLength=this->blockLength/tileNum/BUFFER_NUM;
    xGm.SetGlobalBuffer((__gm__ float *)x + this->blockLength*GetBlockIdx(),this->blockLength);
    yGm.SetGlobalBuffer((__gm__ float *)y +this->blockLength*GetBlockIdx(),this->blockLength);
    zGm.SetGlobalBuffer((__gm__ float *)z+this->blockLength*GetBlockIdx(),this->blockLength);
    pipe.InitBuffer(inQueueX,BUFFER_NUM,this->tileLength*sizeof(float));
    pipe.InitBuffer(inQueueY,BUFFER_NUM,this->tileLength*sizeof(float));
    pipe.InitBuffer(outQueueZ,BUFFER_NUM,this->tileLength*sizeof(float));
}

核函数入口

__global__ __aicore__ void add_custom(GM_ADDR x,GM_ADDR y,GM_ADDR z,AddCustomTilingData tiling){
    KERNEL_TASK_TYPE_DEFAULT(KERNLE_TYPE_AIV_ONLY);
    KernelAdd op;
    op.Init(x,y,z,tiling.totalLength,total.tileNum);//初始GlobalMemory数据地址，数据总长度，分块数量
    op.Process();
}

Process逻辑

__aicore__ inline void KernelAdd::Process(){
    int32_t loopCount=this->tileNum*BUFFER_NUM;
    for(int32_t i=0;i<loopCount;i++){
        CopyIn(i);
        Compute(i);
        CopyOut(i);
    }
}

CopyIn任务

__aicore__ inline void KernelAdd::CopyIn(int32_t progress){
    LocalTensor<float> xLocal=inQueueX.AllocTensor<float>();
    LocalTensor<float> yLocal=inQueueY.AllocTensor<float>();
    DataCopy(xLocal,xGm[progress*this->tileLength],this->tileLength);
    DataCopy(yLocal,yGm[progress*this->tileLength],this->tileLength);
    inQueueX.EnQue(xLocal);
    inQueueY.EnQue(yLocal);
}

Compute任务

__aicore__ inline void KernelAdd::Compute(int32_t progress){
    LocalTensor<float> xLocal=inQueueX.DeQue<float>();
    LocalTensor<float> yLocal=inQueueY.DeQue<float>();
    LocalTensor<float> zLocal=outQueueZ.AllocTensor<float>();
    Add(zLocal,xLocal,yLocal,this->tileLength);
    outQueuZ.EnQue<float>(zLocal);
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}

CopyOut任务

__aicore__ inline void KernelAdd::CopyOut(int32_t progress){
    LocalTensor<float> zLocal=outQueueZ.DeQue<float>();
    DataCopy(zGm[progress*tileLength],zLocal,tileLength);
    outQueueZ.FreeTensor(zlocal);
}

Host侧调用

vector<float> kernel_add(vector<float> &x,vector<float> &y){
    constexpr uint32_t blockDim=8;//启动核数
    uint32_t totalLength=x.size();
    size_t totalByteSize=totalLength*sizeof(float);
    int32_t deviceId=0;
    aclrtStream stream=nullptr;
    AddCustomTilingData tiling={totalLength,8};
    uint8_t *xHost=reinterpret_cast<uint8_t *>(x.data());
    uint8_t *yHost=reinterpret_cast<uint8_t *>(y.data());
    uint8_t *zHost=nullptr;
    uint8_t xDevice=nullptr;
    uint8_t yDevice=nullptr;
    uint8_t zDevice=nullptr;
    aclInit(nullptr);
    aclrtSetDevice(deviceId);
    aclrtCreatStream(&stream);
    aclrtMallocHost((void**)(&zHost),totalByteSize);
    aclrtMalloc((void **)&xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void **)&yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void **)&zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMemcpy(xDevice, totalByteSize, xHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(yDevice, totalByteSize, yHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
    add_custom<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);//<<<blockDim, l2ctrl, stream>>> l2ctrl，保留参数，暂时设置为固定值nullptr。
    aclrtSynchronizeStream(stream);
    aclrtMemcpy(zHost, totalByteSize, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
    std::vector<float> z((float *)zHost, (float *)(zHost + totalByteSize));
    aclrtFree(xDevice);
    aclrtFree(yDevice);
    aclrtFree(zDevice);
    aclrtFreeHost(zHost);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();
    return z; 
}

验证函数VerrifyResult

uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
{
    auto printTensor = [](std::vector<float> &tensor, const char *name) {
        constexpr size_t maxPrintSize = 20;
        std::cout << name << ": ";
        std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
            std::ostream_iterator<float>(std::cout, " "));
        if (tensor.size() > maxPrintSize) {
            std::cout << "...";
        }
        std::cout << std::endl;
    };
    printTensor(output, "Output");
    printTensor(golden, "Golden");
    if (std::equal(golden.begin(), golden.end(), output.begin())) {
        std::cout << "[Success] Case accuracy is verification passed." << std::endl;
        return 0;
    } else {
        std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
        return 1;
    }
    return 0;
}

Host侧验证主程序main

int32_t main(int32_t argc, char *argv[])
{
    constexpr uint32_t totalLength = 8 * 2048;
    constexpr float valueX = 1.2f;
    constexpr float valueY = 2.3f;
    std::vector<float> x(totalLength, valueX);
    std::vector<float> y(totalLength, valueY);

    std::vector<float> output = kernel_add(x, y);

    std::vector<float> golden(totalLength, valueX + valueY);
    return VerifyResult(output, golden);
}

课后实践

请根据本节教程内容，在下面的代码框中补充实现sub_custom.asc，使用Ascend C编程实现矢量减法算子sub_custom。

数据类型：输入与输出均为 float 类型。
数据形状（Shape）：输入Shape为 (8, 2048)，输出Shape与输入相同。
数据布局（Format）：ND。

#include <cstdint>
#include <iostream>
#include <vector>
#include <algorithm>
#include <iterator>
#include "acl/acl.h"
#include "kernel_operator.h"
using namespace AscendC;
constexpr uint32_t BUFFER_NUM = 2; // tensor num for each queue
constexpr uint32_t QUEUE_DEPTH = 2;
constexpr uint32_t TILE_NUM=8;
constexpr uint32_t BLOCK_NUM=8;

struct SubCustomTilingData
{
    uint32_t totalLength;
    uint32_t tileNum;
};

class KernelSub {
public:
    __aicore__ inline KernelSub(){}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
    {
        // 请补充……
        blockLength=totalLegnth/GetBlockNum();
        tileNum=tileNum;
        tileLength=blockLength/tileNum/BUFFER_NUM;
        //从哪传来的 要分配什么
        //分配内存，每个核的GlobalMemory访问地址和队列的初始buffer
        xGm.SetGlobalBuffer((__gm__ float *)x+this->blockLength*GetBlockId(),tileLength);
        yGm.SetGlobalBuffer((__gm__ float *)y+GetBlockNum()*tileLength,tileLength);//这个__gm__ float *是什么写法
        zGm.SetGlobalBuffer((__gm__ float *)z+tileLength*GetBloackNum(),tileLength);
        //然后下面是用pipe分配内存的
        pipe.InitBuffer(inQueueX,BUFFER_NUM,this->tileLength*sizeof(float));
        pipe.InitBuffer(inQueueY,BUFFER_NUM,tileLength*sizeof(float));//后面这一项填充的是字节数
        pipe.InitBuffer(inQueueY,BUFFER_NUM,tileLength*sizeof(float));
        pipe.InitBuffer(outQueue,BUFFER_NUM,tileLength*sizeof(float));
        
        
    }
    __aicore__ inline void Process()
    {
        // 请补充……
        int32_t loopNum=tileNum*BUFFER_NUM;
        for(int32_t i=0;i<loopNum;i++){
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        // 请补充……
        LocalTensor<float>xLocal=inQueueX.AllocTensor<float>();
        LocalTensor<float>yLocal=inQueueY.AllocTensor<float>();
        DataCopy(xLocal,xGm[process*tileLength],tileLength);
        DataCopy(xLocal,yGm[process*tileLength].tileLength);
        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
        
    }
    __aicore__ inline void Compute(int32_t progress)
    {
        // 请补充……
        LocalTensor<float>xLocal=inQueueX.Deque<float>();
        LocalTensor<float>yLocal=inQueueY.Deque<float>();
        LocalTensor<float>zLocal=outQueueZ.AllocTensor<float>();
        Sub(zLocal,xLocal,yLocal,tileLength);
        outQueueZ.EnQue<float>(zLocal);
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress)
    {
        // 请补充……
        LocalTensor<float> zLocal=outQueuZ.DeQue<float>();
        DataCopy(zGm[progress*tileLength],zLocal,tileLength);
        outQueueZ.FreeTensor(zLocal);
    }

private:
    // 请补充……
    Tpipe pipe;
    TQue<AscendC::TPosition::VECIN,QUEUE_DEPTH> inQueueX,inQueueY;
    TQue<TPosition::VECOUT,QUEUE_DEPTH> outQueueZ;
    GlobalTensor<float> xGm,yGm,zGm;
    uint32_t blockLength,tileNum,tileLength;//tiling结构体 totalLength，tileNum
};

__global__ __aicore__ void sub_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, SubCustomTilingData tiling)
{
    // 请补充……
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
    KernelSub op;
    op.Init(x,y,z,tiling.totalLength,tiling.tileNum);
    op.process();
}

std::vector<float> kernel_sub(std::vector<float> &x, std::vector<float> &y)
{
    // 请补充……
    constexpr uint32_t blockDim=BLOCK_NUM;
    uint32_t totalLength=x.size();
    size_t totalByteSize=totalLength*sizeof(float);
    int32_t deviceId=0;
    aclrtStream stream=nullptr;
    AddCustomTilingData tiling={totalLength,TILE_NUM}
    uint8_t *xHost=reinterpret_cast<uint8_t*>(x.data());//按字节粒度操作内存,二进制层面直接重解释指针类型,强制重解释为1字节无符号字节指针uint8_t *,用这个指针可以逐字节访问 Host 内存
    uint8_t *yHost=reinterpret_cast<uint8_t*>(y.data());
    uint8_t *zHost =nullptr;
    uint8_t *xDevice=nullptr;
    uint8_t *yDevice=nullptr;
    uint8_t *zDevice=nullptr;
    aclInit(nullptr);
    aclrtSetDevice(deviceId);
    aclrtCreatStream(&stream);
    aclrtMallocHost((void**)(&zHost),totalByteSize);
 
    aclrtMalloc((void**)&xDevice,totalByteSize,ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMemcpy(xDevice,totalByteSize,xHost,totalByteSize,ACL_MEMCPY_HOST_TO_DEVICE);
    
    sub_custom<<<BlockDim,nullptr,stream>>>(xDevice,yDevice,zDevice,tiling);//device作资源管理指针
    aclrtSynchronizeStream(stream);
    aclrtMemcpy(zHost, totalByteSize, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
    std::vector<float> z((float *)zHost, (float *)(zHost + totalByteSize));
    aclrtFree(xDevice);
    aclrtFree(yDevice);
    aclrtFree(zDevice);
    aclrtFreeHost(zHost);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();
    return z; 
}

uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
{
    auto printTensor = [](std::vector<float> &tensor, const char *name) {
        constexpr size_t maxPrintSize = 20;
        std::cout << name << ": ";
        std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
            std::ostream_iterator<float>(std::cout, " "));
        if (tensor.size() > maxPrintSize) {
            std::cout << "...";
        }
        std::cout << std::endl;
    };
    printTensor(output, "Output");
    printTensor(golden, "Golden");
    if (std::equal(golden.begin(), golden.end(), output.begin())) {
        std::cout << "[Success] Case accuracy is verification passed." << std::endl;
        return 0;
    } else {
        std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
        return 1;
    }
    return 0;
}

int32_t main(int32_t argc, char *argv[])
{
    constexpr uint32_t totalLength = 8 * 2048;
    constexpr float valueX = 1.2f;
    constexpr float valueY = 2.3f;
    std::vector<float> x(totalLength, valueX);
    std::vector<float> y(totalLength, valueY);

    // 请补充……
    std::vector<float> output=kernel_sub(x,y);

    std::vector<float> golden(totalLength, valueX - valueY);
    return VerifyResult(output, golden);
}