算子开发-昇腾CANN训练营-Ascend C-Mmad 算子迁移

1. 项目介绍

		C = A * B + Bias

A、B为源操作数，A为左矩阵，形状为[M, K]；B为右矩阵，形状为[K, N]。
C为目的操作数，存放矩阵乘结果的矩阵，形状为[M, N]。
Bias为矩阵乘偏置，形状为[N]。对A*B结果矩阵的每一行都采用该Bias进行偏置。

  • 源仓库代码介绍：现有代码包含 带 bias 和不带 bias 参数两个工程。都支持分离架构（mmad_custom_cube_only.h）和耦合架构（mmad_custom.h）实现。基于 Ascend C 低级API实现，手动实现数据搬运。
  • 迁移目标：对分离架构带 bias 参数的实现增加 tiling 结构的数据搬运划分。

2. 实现方案：

  • Tiling 策略：对B矩阵进行列分块，因为 cube 计算单位为 1616，所以设置tiling 参数列方向shape 为16；对单核计算数据的切分不应太大。本测试单 cube 上计算矩阵shape 为3232，数据较少，A矩阵一次性读入；避免对A 矩阵 tiling只要是为了避免多次计算结果最后还要加和计算。
  • 单 cube 核 tiling 计算结果汇总策略：分离架构数据流向为 GM -> L1 -> L0A/L0B -> Cube -> L0C -> Fixpipe -> GM 和 GM -> L1 -> L0A/L0B -> Cube -> L0C -> L1;tiling 计算结果考虑暂存，单核计算完成汇总后经过 Fixpipe 通路搬运到 GM。
  • 后者通路使用 L1 缓存空间，会占用输入 L1空间，不可行；
  • 矩阵计算临时空间 TPosition 有 TSCM，但是查阅文档，TSCM 只支持标量数据 VECIN/VECOUT/VECCALC -> TSCM,矩阵计算无法使用，不可行；
  • 结论：暂存数据汇总的方案不可行，只能每次计算完仍然走 L0C -> Fixpipe -> GM 的数据通路。

4. 核心代码

  • op_host\mmad_tiling.h


#ifndef MMAD_TILING_H
#define MMAD_TILING_H
#include "register/tilingdata_base.h"

namespace optiling {
  BEGIN_TILING_DATA_DEF(TilingData)
  //对 B 矩阵 tiling 参数,split N
  TILING_DATA_FIELD_DEF(uint32_t, tileBBlockShape);
  END_TILING_DATA_DEF;
  
  REGISTER_TILING_DATA_CLASS(Mmad, TilingData)
}
#endif  //MMAD_TILING_H
  • op_kernel\mmad.cpp
#include "kernel_operator.h"

class KernelMmad {
public:
    __aicore__ inline KernelMmad()
    {
        aSize = m * k;
        bSize = k * n;
        cSize = m * n;
    }
    __aicore__ inline void Init(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, uint32_t tileBBlockShape)
    {
        // set cube only
        KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIC_ONLY);
        cubeBlockShape = tileBBlockShape;
        CubeBlockSize = cubeBlockShape * cubeBlockShape;

        aGM.SetGlobalBuffer((__gm__ half *)a);
        bGM.SetGlobalBuffer((__gm__ half *)b);
        cGM.SetGlobalBuffer((__gm__ float *)c);
        biasGM.SetGlobalBuffer((__gm__ float *)bias);
        pipe.InitBuffer(inQueueA1, 1, aSize * sizeof(half));
        pipe.InitBuffer(inQueueA2, 1, aSize * sizeof(half));
        pipe.InitBuffer(inQueueB1, 1, bSize * sizeof(half));
        pipe.InitBuffer(inQueueB2, 1, k * cubeBlockShape * sizeof(half));
        pipe.InitBuffer(outQueueCO1, 1, m * cubeBlockShape * sizeof(float));
        pipe.InitBuffer(inQueueC1, 1, n * sizeof(float));
        pipe.InitBuffer(inQueueC2, 1, cubeBlockShape * sizeof(float));
    }

    __aicore__ inline void Process()
    {
        CopyIn();
        SplitA();
        AscendC::LocalTensor<half> a2Local = inQueueA2.DeQue<half>();
        AscendC::LocalTensor<half> b1Local = inQueueB1.DeQue<half>();
        AscendC::LocalTensor<float> bias1Local = inQueueC1.DeQue<float>();
        for(int i=0;i<n/cubeBlockShape;i++)
        {
            SplitB(b1Local,i);
            SplitBias(bias1Local,i);
            Compute(a2Local);
            CopyOut(i);
        }
        inQueueA2.FreeTensor(a2Local);
        inQueueB1.FreeTensor(b1Local);
        inQueueC1.FreeTensor(bias1Local);
    }

private:
    __aicore__ inline uint32_t CeilCubeBlock(uint32_t len) {
        return (len + cubeBlockShape - 1) / cubeBlockShape;
    }

    __aicore__ inline void CopyIn()
    {
        AscendC::LocalTensor<half> a1Local = inQueueA1.AllocTensor<half>();
        AscendC::LocalTensor<half> b1Local = inQueueB1.AllocTensor<half>();
        AscendC::LocalTensor<float> bias1Local = inQueueC1.AllocTensor<float>();

        AscendC::Nd2NzParams nd2nzA1Params;
        nd2nzA1Params.ndNum = 1;
        nd2nzA1Params.nValue = m;
        nd2nzA1Params.dValue = k;
        nd2nzA1Params.srcNdMatrixStride = 0;
        nd2nzA1Params.srcDValue = k;
        nd2nzA1Params.dstNzC0Stride = CeilCubeBlock(m) * cubeBlockShape;
        nd2nzA1Params.dstNzNStride = 1;
        nd2nzA1Params.dstNzMatrixStride = 0;
        AscendC::DataCopy(a1Local, aGM, nd2nzA1Params);

        AscendC::Nd2NzParams nd2nzB1Params;
        nd2nzB1Params.ndNum = 1;
        nd2nzB1Params.nValue = k;
        nd2nzB1Params.dValue = n;
        nd2nzB1Params.srcNdMatrixStride = 0;
        nd2nzB1Params.srcDValue = n;
        nd2nzB1Params.dstNzC0Stride = CeilCubeBlock(k) * cubeBlockShape;
        nd2nzB1Params.dstNzNStride = 1;
        nd2nzB1Params.dstNzMatrixStride = 0;
        AscendC::DataCopy(b1Local, bGM, nd2nzB1Params);

        AscendC::DataCopy(bias1Local, biasGM, n);
        inQueueA1.EnQue(a1Local);
        inQueueB1.EnQue(b1Local);
        inQueueC1.EnQue(bias1Local);
    }

    __aicore__ inline void SplitA()
    {
        AscendC::LocalTensor<half> a1Local = inQueueA1.DeQue<half>();
        AscendC::LocalTensor<half> a2Local = inQueueA2.AllocTensor<half>();

        uint32_t dstOffset = CeilCubeBlock(k) * CubeBlockSize;
        uint32_t srcOffset = CubeBlockSize;
 
        //nz to zz
        AscendC::LoadData2DParams loadDataParams;
        loadDataParams.repeatTimes = CeilCubeBlock(k);
        loadDataParams.srcStride = CeilCubeBlock(m);
        loadDataParams.dstGap = 0;
        loadDataParams.ifTranspose = false;
        for (int i = 0; i < CeilCubeBlock(m); ++i) {
            AscendC::LoadData(a2Local[i * dstOffset], a1Local[i * srcOffset], loadDataParams);
        }

        inQueueA2.EnQue<half>(a2Local);
        inQueueA1.FreeTensor(a1Local);
    }
    __aicore__ inline void SplitB(const AscendC::LocalTensor<half>& b1Local,const uint32_t bSplitIdx)
    {
        AscendC::LocalTensor<half> b2Local = inQueueB2.AllocTensor<half>();

        // Nz -> Zn
        AscendC::LoadData2DParams loadDataParams;
        loadDataParams.repeatTimes = CeilCubeBlock(k);
        loadDataParams.srcStride = 1;
        loadDataParams.ifTranspose = true;
        AscendC::LoadData(b2Local, b1Local[bSplitIdx * CeilCubeBlock(n) * CubeBlockSize], loadDataParams);

        inQueueB2.EnQue<half>(b2Local);
    }
    __aicore__ inline void SplitBias(const AscendC::LocalTensor<float>& bias1Local,const uint32_t bSplitIdx)
    {
        AscendC::LocalTensor<float> bias2Local = inQueueC2.AllocTensor<float>();
        AscendC::DataCopy(bias2Local, bias1Local[bSplitIdx*cubeBlockShape], cubeBlockShape);
        inQueueC2.EnQue<float>(bias2Local);
    }
    __aicore__ inline void Compute(const AscendC::LocalTensor<half> a2Local)
    {
        AscendC::LocalTensor<half> b2Local = inQueueB2.DeQue<half>();
        AscendC::LocalTensor<float> bias2Local = inQueueC2.DeQue<float>();
        AscendC::LocalTensor<float> c1Local = outQueueCO1.AllocTensor<float>();
        AscendC::MmadParams mmadParams;
        mmadParams.m = m;
        mmadParams.n = cubeBlockShape;
        mmadParams.k = k;
        mmadParams.cmatrixInitVal = false;
        AscendC::Mmad(c1Local, a2Local, b2Local, bias2Local, mmadParams);
        outQueueCO1.EnQue<float>(c1Local);
        inQueueB2.FreeTensor(b2Local);
        inQueueC2.FreeTensor(bias2Local);
    }
    __aicore__ inline void CopyOut(const uint32_t bSplitIdx )
    {
        AscendC::LocalTensor<float> c1Local = outQueueCO1.DeQue<float>();
        // FixpipeParamsV220 : CO1 -> gm
        AscendC::FixpipeParamsV220 fixpipeParams;
        fixpipeParams.nSize = cubeBlockShape;
        fixpipeParams.mSize = m;
        fixpipeParams.srcStride = cubeBlockShape*sizeof(float); //表示源NZ矩阵中相邻Z排布的起始地址偏移
        fixpipeParams.dstStride = n;

        fixpipeParams.ndNum = 1;
        fixpipeParams.srcNdStride = 0;
        fixpipeParams.dstNdStride = 0;
        // 默认设置 nz -> nd
        AscendC::Fixpipe(cGM[bSplitIdx*cubeBlockShape], c1Local, fixpipeParams);
        outQueueCO1.FreeTensor(c1Local);
    }

private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::TPosition::A1, 1> inQueueA1;
    AscendC::TQue<AscendC::TPosition::A2, 1> inQueueA2;
    AscendC::TQue<AscendC::TPosition::B1, 1> inQueueB1;
    AscendC::TQue<AscendC::TPosition::B2, 1> inQueueB2;
    AscendC::TQue<AscendC::TPosition::CO1, 1> outQueueCO1;
    // AscendC::TQue<AscendC::TPosition::CO1, 1> outQueueCO1_;  //分离架构无 CO2
    AscendC::TQue<AscendC::TPosition::C1, 1> inQueueC1;
    AscendC::TQue<AscendC::TPosition::C2, 1> inQueueC2;

    AscendC::GlobalTensor<half> aGM;
    AscendC::GlobalTensor<half> bGM;
    AscendC::GlobalTensor<float> cGM;
    AscendC::GlobalTensor<float> biasGM;
    uint16_t m = 32, k = 32, n = 32;
    uint16_t aSize, bSize, cSize;
    uint32_t cubeBlockShape,CubeBlockSize;
};


extern "C" __global__ __aicore__ void mmad(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace, GM_ADDR tiling) {

    GET_TILING_DATA(tiling_data,tiling);
    KernelMmad op;
    op.Init(a,b,bias,c,tiling_data.tileBBlockShape);
    op.Process();
}

3. 项目实施中遇到的问题及解决方案

  • 运行环境问题：在华为云测试服务器运行环运行境脚本 init_env.sh 安装 cann-toolkit 时，经常遇到网络代理错误，安装包下载失败的问题。

解决方案：先不要运行环境脚本，本地下载 cann-toolkit 安装包上传到服务器，再运行安装环境脚本。

  • 执行权限问题
Permission denied

解决方案：项目文件权限对其他用户设置为 0。

  • tiling 注册失败
/home/user/cann/03mmad/mmad-2/MmadCustom/op_kernel/mmad.cpp:12:21: error: use of undeclared identifier 'tiling_data'
    GET_TILING_DATA(tiling_data,tiling);
                    ^
/home/user/cann/03mmad/mmad-2/MmadCustom/op_kernel/mmad.cpp:14:24: error: use of undeclared identifier 'tiling_data'
    op.Init(a,b,bias,c,tiling_data.tileBBlockShape);
...
[ERROR] [ascend910b] Mmad do not registe tiling struct!!!

分析：以上错误为本地编译错误，推测本地不是昇腾硬件安装的软件环境导致。
解决方案：在服务器上编译正常。

  • ACLNN 算子调用运行错误
aclnnMmadmGetWorkspaceSize failed. ERROR: 161002

分析：似乎与数据分配的空间有关。
解决：检查 ACLNN 调用示例代码，发现输入数据 bias 数据类型错误，修正后解决。

  • ACLNN 算子调用运行错误
aclrtSynchronizeStream failed. ERROR: 507015

分析：检查错误代码为 AICore 运行错误，通过检查 ~/ascend/log/debug/plog 中的日志文件发现错误代码 rtKernelLaunchWithHandleV2:ErrCode=107000, 为 UB 内存错误问题。
解决：检查发现使用了 CO2 数据空间，分离架构不支持该 TPosition，应该使用分离架构的数据通路，L0C -> Fixpipe -> GM。

  • ACLNN 算子调用测试结果错误

分析：测试结果与标杆结果数据比对不一致，考虑每次计算完 Fixpipe 向 GM 搬运数据 NZ2ND 的过程数据组织错误。
解决：目标 GM 数据位置应为第一行元素的索引。

AscendC::Fixpipe(cGM[bSplitIdx*m*cubeBlockShape], c1Local, fixpipeParams); 
更正为
AscendC::Fixpipe(cGM[bSplitIdx*cubeBlockShape], c1Local, fixpipeParams);
  • ST 测试未找到 Mmad op
b"ATC run failed, Please check the detail log, Try 'atc --help' for more information"
b'EZ3003: [PID: 174824] 2025-05-22-16:46:44.846.249 No supported Ops kernel and engine are found for [Mmad], optype [Mmad].'
b'Possible Cause: The operator is not supported by the system. Therefore, no hit is found in any operator information library.'

分析：使用 cann-ops 仓其进行其他单算子编译安装的方式，出现不同的错误，环境应有问题。
解决：仓库 CI 检查流程会自动进行 ST，使用 CI ST 测试环境。

  • CI 门禁x86编译错误

分析：经管理员判定服务器环境配置规格问题。
解决：约一个月后CI部门解决了该问题，代码已合入。

posted @ 2025-05-27 10:16  安洛8  阅读(233)  评论(0)    收藏  举报