昇腾社区首页
中文
注册
开发者
下载
支持
积分兑换 NEW

如何使用更底层编程方式提升算子性能

Ascend C编程范式提供了基于Pipe进行算子开发的方式，Pipe（TPipe类）统一管理Device端内存等资源，开发者无需感知内存管理、DoubleBuffer流水、同步等处理，只需要按照计算流编写算子即可，但由此也带来了一些运行时开销（如TPipe创建、InitBuffer等）。

Ascend C提供更底层的编程方式，开发者可以直接构造指定地址和存储位置的LocalTensor，使用该Tensor传入计算、搬运等API进行编程，减少了Pipe编程方式中带来的运行时开销，更有助于开发者实现极致性能。两种编程方式的对比如下：

图1 两种编程方式的对比（伪码）

使用约束和限制

相比于Pipe编程方式，使用更底层编程时需要遵循以下约束和限制：

  • 在更底层编程场景中，开发者不能使用TPipe/TQue/TQueBind/TBufPool等框架接口，仅能使用LocalTensor对象以及部分API（和上述框架接口混用可能会出现未定义行为）。LocalTensor的构造方式有两种，通过LocalTensor构造函数或者LocalMemAllocator构造。
  • 在更底层编程场景中只能使用部分API。具体支持的API列表见支持的API范围。因为不在列表范围内的API内部依赖TPipe分配事件ID，可能会和开发者定义的事件ID产生冲突。
  • 同步事件需要由开发者使用SetFlag/WaitFlag(ISASI)和PipeBarrier(ISASI)手动插入，事件的类型/事件ID由开发者自行管理，但需要注意事件ID不能使用6/7（可能与内部使用的事件ID出现冲突，进而出现未定义行为）。
  • 由于需要使用SetFlag/WaitFlag/PipeBarrier底层同步接口（属于ISASI硬件体系结构相关的接口），无法保证跨硬件版本兼容。
  • Kernel入口处需要开发者手动调用InitSocState接口用来初始化全局状态寄存器（在TPipe框架编程中，初始化过程由TPipe完成，无需开发者关注）。

使用示例

这段代码定义了一个用于执行矢量加法的核函数add_custom和算子类KernelAdd。算子类KernelAdd的核心实现是数据的搬入、加法计算和搬出。

  • Kernel入口处手动调用InitSocState接口来初始化全局状态寄存器。
  • 使用LocalTensor接口直接构造LocalTensor对象。
  • 开发者手动进行同步控制，插入的同步事件如下：
    • 正向同步

      在本次数据搬入和计算之间，插入MTE2_V（矢量计算流水等待MT2搬运流水）同步事件，确保数据搬入之后再进行计算；在本次数据计算和搬出之间，插入V_MTE3（MTE3搬运流水等待矢量计算流水）同步事件，确保数据计算完成后再进行搬出。

    • 反向同步

      在上一次的数据计算和本次数据搬入之间，插入V_MTE2（MT2搬运流水等待矢量计算流水）同步事件，确保上一次的数据计算完成后，本次的数据再进行搬入。防止本次的数据会覆盖掉上一次未计算完成的数据；在上一次的数据搬出和本次数据计算之间，插入MTE3_V（矢量计算流水等待MT3搬运流水）同步事件，确保上一次的数据搬出后，再进行本次数据的计算。防止本次的数据会覆盖掉上一次未搬出的数据。

      上述的同步逻辑在使用Pipe编程框架时，框架会使用EnQue/DeQue/AllocTensor/FreeTensor进行封装。您可以通过编程模型设计原理来了解应该如何在使用更底层编程方式时手动进行同步控制。

 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
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
 
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);
    }
    __aicore__ inline void Process()
    {
        int32_t loopCount = this->tileNum * BUFFER_NUM;
        uint32_t xAddr = 0;
        uint32_t yAddr = this->tileLength * sizeof(DTYPE_X);
        uint32_t zAddr = this->tileLength * (sizeof(DTYPE_X) + sizeof(DTYPE_Y));
        AscendC::LocalTensor<DTYPE_X> xLocal(AscendC::TPosition::VECIN, xAddr, this->tileLength);
        AscendC::LocalTensor<DTYPE_Y> yLocal(AscendC::TPosition::VECIN, yAddr, this->tileLength);
        AscendC::LocalTensor<DTYPE_Z> zLocal(AscendC::TPosition::VECOUT, zAddr, this->tileLength);
        for (int32_t i = 0; i < loopCount; i++) {
            if (i != 0) {
                // xLocal/yLocal的反向依赖，后一次搬入需要等上一次的计算结束
                AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(0);
            }
            AscendC::DataCopy(xLocal, xGm[i * this->tileLength], this->tileLength);
            AscendC::DataCopy(yLocal, yGm[i * this->tileLength], this->tileLength);
            AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(0);
            AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(0);
            if (i != 0) {
                // zLocal的反向依赖，后一次计算需要等上一次的搬出结束
                AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(0);
            }
            AscendC::Add(zLocal, xLocal, yLocal, this->tileLength);
            if (i != loopCount -1) {
                // xLocal/yLocal的反向依赖，后一次的搬入需要等上一次的计算结束
                AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(0);
            }
            AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(0);
            AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(0);
            DataCopy(zGm[i * this->tileLength], zLocal, this->tileLength);
            if (i != loopCount -1) {
                // zLocal的反向依赖，后一次计算需要等上一次的搬出结束
                AscendC::SetFlag<AscendC::HardEvent::MTE3_V>(0);
            }
        }
    }
private:
    AscendC::GlobalTensor<half> xGm, yGm, zGm;
};

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    // 初始化全局状态寄存器（在TPipe框架编程中，初始化过程由TPipe完成，无需开发者关注；更底层编程方式中需要开发者手动调用）
    AscendC::InitSocState();
    KernelAdd op;
    op.Init(x, y, z);
    op.Process(); 
}

支持的API范围

表1 针对 Atlas 推理系列产品 AI Core，支持的API范围

接口分类

接口名称

基础API > 标量计算

ScalarGetCountOfValue、ScalarCountLeadingZero、ScalarCast、CountBitsCntSameAsSignBit、ScalarGetSFFValue

基础API > 矢量计算 > 基础算术

Exp、Ln、Abs、Reciprocal、Sqrt、Rsqrt、Not、Relu、VectorPadding(ISASI)、Add、Sub、Mul、Div、Max、Min、

BilinearInterpolation(ISASI)、Adds、Muls、Maxs、Mins、LeakyRelu、Axpy

基础API > 矢量计算 > 逻辑计算

And、Or

基础API > 矢量计算 > 复合计算

AddRelu、AddReluCast、AddDeqRelu、SubRelu、SubReluCast、MulAddDst、MulCast、FusedMulAdd、FusedMulAddRelu

基础API > 矢量计算 > 比较指令

Compare、Compare（结果存入寄存器）、CompareScalar、GetCmpMask(ISASI)、SetCmpMask(ISASI)

基础API > 矢量计算 > 选择指令

Select、GatherMask

基础API > 矢量计算 > 精度转换指令

Cast、CastDeq

基础API > 矢量计算 > 归约指令

WholeReduceMax、WholeReduceMin、WholeReduceSum、BlockReduceMax、BlockReduceMin、BlockReduceSum、PairReduceSum、RepeatReduceSum、GetReduceMaxMinCount

基础API > 矢量计算 > 数据转换

Transpose、TransDataTo5HD

基础API > 矢量计算 > 数据填充

Duplicate

基础API > 矢量计算 > 排序组合（ISASI）

ProposalConcat、ProposalExtract、RpSort16、MrgSort4、GetMrgSortResult

基础API > 矢量计算 > 数据分散/数据收集

Gather、Scatter(ISASI)

基础API > 矢量计算 > 掩码操作

SetMaskCount、SetMaskNorm、SetVectorMask、ResetMask

基础API > 矢量计算 > 量化设置

SetDeqScale

基础API > 数据搬运 > DataCopy

普通数据搬运

基础API > 数据搬运

InitConstValue(ISASI)、LoadData(ISASI)、SetAippFunctions(ISASI)、LoadImageToLocal(ISASI)、LoadUnzipIndex(ISASI)、LoadDataUnzip(ISASI)、SetLoadDataBoundary(ISASI)、SetLoadDataPaddingValue(ISASI)

基础API > 内存管理与同步控制 > 核内同步

SetFlag/WaitFlag(ISASI)、PipeBarrier(ISASI)

基础API > 缓存处理

DataCachePreload、DataCacheCleanAndInvalid、ICachePreLoad(ISASI)

基础API > 系统变量访问

GetBlockNum、GetBlockIdx、GetDataBlockSizeInBytes、GetArchVersion、GetTaskRation、InitSocState、GetProgramCounter(ISASI)、CheckLocalMemoryIA(ISASI)

基础API > 原子操作

SetAtomicAdd、SetAtomicNone

基础API > 矩阵计算

Mmad(ISASI)
表2 针对 Atlas A2 训练系列产品 / Atlas 800I A2 推理产品 /A200I A2 Box 异构组件，支持的API范围

接口分类

接口名称

备注

基础API > 标量计算

ScalarGetCountOfValue、ScalarCountLeadingZero、ScalarCast、CountBitsCntSameAsSignBit、ScalarGetSFFValue、ToBfloat16、ToFloat

-

基础API > 矢量计算 > 基础算术

Exp、Ln、Abs、Reciprocal、Sqrt、Rsqrt、Not、Relu、Add、Sub、Mul、Div、Max、Min、BilinearInterpolation(ISASI)、Adds、Muls、Maxs、Mins、LeakyRelu、Axpy

-

基础API > 矢量计算 > 逻辑计算

And、Or、ShiftLeft、ShiftRight

-

基础API > 矢量计算 > 复合计算

AddRelu、AddReluCast、AddDeqRelu、SubRelu、SubReluCast、MulAddDst、MulCast、FusedMulAdd、FusedMulAddRelu

-

基础API > 矢量计算 > 比较指令

Compare、Compare（结果存入寄存器）、CompareScalar、GetCmpMask(ISASI)、SetCmpMask(ISASI)

-

基础API > 矢量计算 > 选择指令

Select、GatherMask

-

基础API > 矢量计算 > 精度转换指令

Cast、CastDeq

-

基础API > 矢量计算 > 归约指令

WholeReduceMax、WholeReduceMin、WholeReduceSum、BlockReduceMax、BlockReduceMin、BlockReduceSum、PairReduceSum、RepeatReduceSum、GetAccVal(ISASI)、GetReduceMaxMinCount

-

基础API > 矢量计算 > 数据转换

Transpose、TransDataTo5HD

-

基础API > 矢量计算 > 数据填充

Duplicate、Brcb

-

基础API > 矢量计算 > 排序组合(ISASI)

Sort32、MrgSort、GetMrgSortResult

-

基础API > 矢量计算 > 数据分散/数据收集

Gather、Gatherb(ISASI)

-

基础API > 矢量计算 > 掩码操作

SetMaskCount、SetMaskNorm、SetVectorMask、ResetMask

-

基础API > 矢量计算 > 量化设置

SetDeqScale

-

基础API > 数据搬运 > DataCopy

普通数据搬运

不支持VECIN/VECCALC/VECOUT -> TSCM通路的数据搬运。

增强数据搬运

不支持VECIN/VECCALC/VECOUT -> TSCM通路的数据搬运。

切片数据搬运

-

随路格式转换

不支持VECIN/VECCALC/VECOUT -> TSCM通路的数据搬运。

基础API > 数据搬运

Copy、DataCopyPad(ISASI)、SetPadValue(ISASI)、SetFixPipeConfig(ISASI)、SetFixpipeNz2ndFlag(ISASI)、SetFixpipePreQuantFlag(ISASI)、InitConstValue(ISASI)、LoadData(ISASI)、LoadDataWithTranspose(ISASI)、SetAippFunctions(ISASI)、LoadImageToLocal(ISASI)、LoadDataWithSparse(ISASI)、SetFmatrix(ISASI)、SetLoadDataBoundary(ISASI)、SetLoadDataRepeat(ISASI)、SetLoadDataPaddingValue(ISASI)、Fixpipe(ISASI)

-

基础API > 内存管理与同步控制 > 核内同步

SetFlag/WaitFlag(ISASI)、PipeBarrier(ISASI)、DataSyncBarrier(ISASI)

-

基础API > 内存管理与同步控制 > 核间同步

CrossCoreSetFlag(ISASI)、CrossCoreWaitFlag(ISASI)

-

基础API > 缓存处理

DataCachePreload、DataCacheCleanAndInvalid、ICachePreLoad(ISASI)、GetICachePreloadStatus(ISASI)

-

基础API > 系统变量访问

GetBlockNum、GetBlockIdx、GetDataBlockSizeInBytes、GetArchVersion、GetTaskRation、InitSocState、GetProgramCounter(ISASI)、GetSubBlockNum(ISASI)、GetSubBlockIdx(ISASI)、GetSystemCycle(ISASI)、

CheckLocalMemoryIA(ISASI)

-

基础API > 原子操作

SetAtomicAdd、SetAtomicType、SetAtomicNone、SetAtomicMax(ISASI)、SetAtomicMin(ISASI)、SetStoreAtomicConfig(ISASI)、GetStoreAtomicConfig(ISASI)

-

基础API > 矩阵计算（ISASI）

Mmad、MmadWithSparse、SetHF32Mode、SetHF32TransMode、SetMMLayoutTransform

-

高阶API > C++标准库 > 算法

max、min、index_sequence

-

高阶API > C++标准库 > 容器函数

tuple、get、make_tuple

-

高阶API > C++标准库 > 类型特性

is_convertible、is_base_of、is_same、enable_if、conditional

-
表3 针对 Atlas A3 训练系列产品 / Atlas A3 推理系列产品 ，支持的API范围

接口分类

接口名称

备注

基础API > 标量计算

ScalarGetCountOfValue、ScalarCountLeadingZero、ScalarCast、CountBitsCntSameAsSignBit、ScalarGetSFFValue、ToBfloat16、ToFloat

-

基础API > 矢量计算 > 基础算术

Exp、Ln、Abs、Reciprocal、Sqrt、Rsqrt、Not、Relu、Add、Sub、Mul、Div、Max、Min、BilinearInterpolation(ISASI)、Adds、Muls、Maxs、Mins、LeakyRelu、Axpy

-

基础API > 矢量计算 > 逻辑计算

And、Or、ShiftLeft、ShiftRight

-

基础API > 矢量计算 > 复合计算

AddRelu、AddReluCast、AddDeqRelu、SubRelu、SubReluCast、MulAddDst、MulCast、FusedMulAdd、FusedMulAddRelu

-

基础API > 矢量计算 > 比较指令

Compare、Compare（结果存入寄存器）、CompareScalar、GetCmpMask(ISASI)、SetCmpMask(ISASI)

-

基础API > 矢量计算 > 选择指令

Select、GatherMask

-

基础API > 矢量计算 > 精度转换指令

Cast、CastDeq

-

基础API > 矢量计算 > 归约指令

WholeReduceMax、WholeReduceMin、WholeReduceSum、BlockReduceMax、BlockReduceMin、BlockReduceSum、PairReduceSum、RepeatReduceSum、GetAccVal(ISASI)、GetReduceMaxMinCount

-

基础API > 矢量计算 > 数据转换

Transpose、TransDataTo5HD

-

基础API > 矢量计算 > 数据填充

Duplicate、Brcb

-

基础API > 矢量计算 > 排序组合(ISASI)

Sort32、MrgSort、GetMrgSortResult

-

基础API > 矢量计算 > 数据分散/数据收集

Gather、Gatherb(ISASI)

-

基础API > 矢量计算 > 掩码操作

SetMaskCount、SetMaskNorm、SetVectorMask、ResetMask

-

基础API > 矢量计算 > 量化设置

SetDeqScale

-

基础API > 数据搬运 > DataCopy

普通数据搬运

不支持VECIN/VECCALC/VECOUT -> TSCM通路的数据搬运。

基础API > 数据搬运

增强数据搬运

不支持VECIN/VECCALC/VECOUT -> TSCM通路的数据搬运。

切片数据搬运

-

随路格式转换

不支持VECIN/VECCALC/VECOUT -> TSCM通路的数据搬运。

Copy、DataCopyPad(ISASI)、SetPadValue(ISASI)、SetFixPipeConfig(ISASI)、SetFixpipeNz2ndFlag(ISASI)、SetFixpipePreQuantFlag(ISASI)、InitConstValue(ISASI)、LoadData(ISASI)、LoadDataWithTranspose(ISASI)、SetAippFunctions(ISASI)、LoadImageToLocal(ISASI)、LoadDataWithSparse(ISASI)、SetFmatrix(ISASI)、SetLoadDataBoundary(ISASI)、SetLoadDataRepeat(ISASI)、SetLoadDataPaddingValue(ISASI)、Fixpipe(ISASI)

-

基础API > 内存管理与同步控制 > 核内同步

SetFlag/WaitFlag(ISASI)、PipeBarrier(ISASI)、DataSyncBarrier(ISASI)

-

基础API > 内存管理与同步控制 > 核间同步

CrossCoreSetFlag(ISASI)、CrossCoreWaitFlag(ISASI)

-

基础API > 缓存处理

DataCachePreload、DataCacheCleanAndInvalid、ICachePreLoad(ISASI)、GetICachePreloadStatus(ISASI)

-

基础API > 系统变量访问

GetBlockNum、GetBlockIdx、GetDataBlockSizeInBytes、GetArchVersion、GetTaskRation、InitSocState、GetProgramCounter(ISASI)、GetSubBlockNum(ISASI)、GetSubBlockIdx(ISASI)、GetSystemCycle(ISASI)、

CheckLocalMemoryIA(ISASI)

-

基础API > 原子操作

SetAtomicAdd、SetAtomicType、SetAtomicNone、SetAtomicMax(ISASI)、SetAtomicMin(ISASI)、SetStoreAtomicConfig(ISASI)、GetStoreAtomicConfig(ISASI)

-

基础API > 矩阵计算（ISASI）

Mmad、MmadWithSparse、SetHF32Mode、SetHF32TransMode、SetMMLayoutTransform

-

高阶API > C++标准库 > 算法

max、min、index_sequence

-

高阶API > C++标准库 > 容器函数

tuple、get、make_tuple

-

高阶API > C++标准库 > 类型特性

is_convertible、is_base_of、is_same、enable_if、conditional

-

高阶API > 模板库函数 > type_traits

is_convertible、is_base_of、is_same、enable_if、conditional

-
父主题： 常用操作