C++语言拓展

预处理符号拓展

• __NPU_ARCH__
  __NPU_ARCH__是Device侧AI Core代码中的预处理宏，用于标识AI处理器的架构版本。该宏由四位数字组成，其中前三位数字用于标识AI Core的IP核(Intellectual Property Core)类型，第四位数字标识该AI Core同一个IP核的配置版本。通过该宏，开发者可以针对不同AI处理器，差异化进行代码适配和优化。产品型号和__NPU_ARCH__对应关系如下表所示：

  表1 产品型号和__NPU_ARCH__的对应关系

  产品型号	__NPU_ARCH__
  Atlas A3 训练系列产品 / Atlas A3 推理系列产品	2201
  Atlas A2 训练系列产品 / Atlas 800I A2 推理产品 /A200I A2 Box 异构组件	2201
  Atlas 200I/500 A2 推理产品	3002
  Atlas 推理系列产品	2002
  Atlas 训练系列产品	1001

  以下为通过__NPU_ARCH__控制在不同AI处理器上算子输出值舍入模式的示例。

  1 2 3 4 5 6 7 8 9 10 11 12

  __aicore__ static inline void CopyOut(uint64_t mulLen) { #if __NPU_ARCH__ == 2002 Cast(dstLocal, srcLocal, RoundMode::CAST_NONE, mulLen); // CAST_NONE表示舍入模式在转换有精度损失时使用CAST_RINT模式，在不涉及精度损失时不进行舍入 #elif __NPU_ARCH__ == 2201 Cast(dstLocal, srcLocal, RoundMode::CAST_RINT, mulLen); // CAST_RINT表示舍入模式为四舍六入五成双舍入 #endif event_t eventVToMTE3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3)); SetFlag<HardEvent::V_MTE3>(eventVToMTE3); WaitFlag<HardEvent::V_MTE3>(eventVToMTE3); CommonCopyOut<float>(dstLocal, mulLen); // 拷贝LocalTensor至GlobalTensor }

• ASCEND_IS_AIV、ASCEND_IS_AIC

  ASCEND_IS_AIV和ASCEND_IS_AIC是通过C++宏实现的条件判断语句，用于在__aicore__修饰的函数中实现代码的条件编译。基于分离架构（AIC核和AIV核分离）开发融合算子时，算子逻辑中同时涉及AIV核和AIV核的处理逻辑，并需要进行核间同步，此时需要通过ASCEND_IS_AIV/ ASCEND_IS_AIC进行AIV和AIC核代码的隔离。

  

  当使用高阶API Matmul时，其内部已通过REGIST_MATMUL_OBJ宏方式实现了AIV与AIC核代码的隔离，用户无需再使用该宏进行处理。

  以MatmulNzCustom算子为例，该算子在分离架构下需要分别在AIV核和AIC核上实现不同的逻辑。具体而言，AIV核负责将矩阵数据搬入Unified Buffer，完成数据的重排（将矩阵数据转换为NZ格式），并将其写入Global Memory。而AIC核则直接从Global Memory读取已经重排好的NZ格式数据，并执行矩阵乘法（Matmul）计算。由于AIV核和AIC核的代码逻辑不同，需要通过ASCEND_IS_AIV和ASCEND_IS_AIC宏进行代码隔离，确保在编译时分别生成适用于AIV核和AIC核的代码。

  示例伪码如下：

  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34

  template <typename AType, typename BType, typename CType, typename BiasType> __aicore__ inline void MatmulKernel<AType, BType, CType, BiasType>::Process(AscendC::TPipe *pipe) { // 利用AIV核的Vector计算单元实现ND2NZ格式转换。如下代码中MatrixBtoNZ为将B矩阵进行ND2NZ格式转换的函数。 if ASCEND_IS_AIV { pipe->InitBuffer(ubBuf, TOTAL_UB_SIZE); MatrixBtoNZ<typename B_TYPE::T>(tempGM, bGMNZ, tiling, isTransB, ubBuf, tiling.baseK, tiling.baseN); // Vector侧实现的ND2NZ函数 SyncAll(); // AIC核和AIV核同步 AscendC::CrossCoreSetFlag<0x2, PIPE_MTE3>(0x4); return; } if ASCEND_IS_AIC { AscendC::CrossCoreWaitFlag(0x4); // 等待AIV核完成ND2NZ格式转换 } ... ... // 设置左矩阵A、右矩阵B、Bias。 matmulObj.SetTail(tailM, tailN); matmulObj.SetTensorA(aGlobal, false); matmulObj.SetTensorB(bGlobal, false); if (tiling.isBias) { matmulObj.SetBias(biasGlobal); } // 完成矩阵乘操作 matmulObj.IterateAll(cGlobal); // 结束矩阵乘操作 matmulObj.End(); }

函数修饰符

示例代码如下：

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28

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__ 表示该函数在Device侧执行, __inline__强制该函数内联
    {
        this->blockLength = totalLength / AscendC::GetBlockNum();
        this->tileNum = tileNum;
        this->tileLength = this->blockLength / tileNum / BUFFER_NUM;

        xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
        yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
        zGm.SetGlobalBuffer((__gm__ DTYPE_Z *)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
        pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y));
        pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Z));
    }
    ... ...
};



extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) // 使用__global__表示该函数为核函数入口，__aicore__ 表示该函数在Device侧执行
{
    GET_TILING_DATA(tiling_data, tiling);
    KernelAdd op;
    op.Init(x, y, z, tiling_data.totalLength, tiling_data.tileNum);
    op.Process();
}

地址空间修饰符

内置常量

内置变量