Transpose Tiling

功能说明

用于获取Transpose Tiling参数。

函数原型

说明

GetConfusionTransposeMaxMinTmpSize、GetConfusionTransposeTilingInfo、GetConfusionTransposeOnlyTilingInfo接口废弃，并将在后续版本移除，请不要使用该接口。请使用GetTransposeMaxMinTmpSize、GetTransposeTilingInfo接口。

获取最小临时空间大小

Text

void GetTransposeMaxMinTmpSize(const AscendC::TensorShape& srcShape, const uint32_t typeSize, const uint32_t transposeTypeIn, uint32_t& maxValue, uint32_t& minValue)

Text

void GetConfusionTransposeMaxMinTmpSize(const AscendC::TensorShape& srcShape, const uint32_t typeSize, const uint32_t transposeTypeIn, uint32_t& maxValue, uint32_t& minValue)

获取Transpose Tiling

Text

void GetTransposeTilingInfo(const AscendC::TensorShape& srcShape, const uint32_t stackBufferSize, const uint32_t typeSize, const uint32_t transposeTypeIn, optiling::ConfusionTransposeTiling& tiling)

Text

void GetTransposeTilingInfo(const AscendC::TensorShape& srcShape, const uint32_t stackBufferSize, const uint32_t typeSize, const uint32_t transposeTypeIn, AscendC::tiling::ConfusionTransposeTiling& tiling)

Text

void GetConfusionTransposeOnlyTilingInfo(const AscendC::TensorShape& srcShape, const uint32_t stackBufferSize, const uint32_t typeSize, optiling::ConfusionTransposeTiling& tiling)

Text

void GetConfusionTransposeOnlyTilingInfo(const AscendC::TensorShape& srcShape, const uint32_t stackBufferSize, const uint32_t typeSize, AscendC::tiling::ConfusionTransposeTiling& tiling)

Text

void GetConfusionTransposeTilingInfo(const AscendC::TensorShape& srcShape, const uint32_t stackBufferSize, const uint32_t typeSize, const uint32_t transposeTypeIn, optiling::ConfusionTransposeTiling& tiling)

Text

void GetConfusionTransposeTilingInfo(const AscendC::TensorShape& srcShape, const uint32_t stackBufferSize, const uint32_t typeSize, const uint32_t transposeTypeIn, AscendC::tiling::ConfusionTransposeTiling& tiling)

参数说明

表1 GetTransposeMaxMinTmpSize接口参数说明

参数名	输入/输出	含义
srcShape	输入	输入Tensor的shape信息，参数类型为AscendC::TensorShape，具体srcShape传入格式为：场景1：[B, N, S, H/N] 场景2：[B, N, S, H/N] 场景3：[B, N, S, H/N] 场景4：[B, N, S, H/N] 场景5：[B, N, S, H/N] 场景6：[B, N, S, H/N] 场景7：[H, W] 场景13：[H, W]或者[N, H, W]，仅支持Ascend 950PR/Ascend 950DT。场景14：[N, H, W]，仅支持Ascend 950PR/Ascend 950DT。场景15：[N, H, W]，仅支持Ascend 950PR/Ascend 950DT。场景16：[H, W]，仅支持Ascend 950PR/Ascend 950DT。
typeSize	输入	输入的数据类型大小，单位为字节。比如输入的数据类型为half，此处应传入2。
transposeTypeIn	输入	选择数据排布及reshape的类型，根据输入数字选择对应的场景，参数范围为[1,7]。针对Ascend 950PR/Ascend 950DT，该参数取值范围为[1, 7]和[13, 16]。场景1（NZ2ND，1、2轴互换）：1 场景2（NZ2NZ，1、2轴互换）：2 场景3（NZ2NZ，尾轴切分）：3 场景4（NZ2ND，尾轴切分）：4 场景5（NZ2ND，尾轴合并）：5 场景6（NZ2NZ，尾轴合并）：6 场景7（二维转置）：7 场景13 （二维转置或者三维转置中后两维转置）：13，仅支持Ascend 950PR/Ascend 950DT。场景14 （三维转置中第一维和第二维互换）：14，仅支持Ascend 950PR/Ascend 950DT。场景15 （三维转置中第一维和第三维互换）：15，仅支持Ascend 950PR/Ascend 950DT。场景16 （使用交织指令进行两维ND2NZ转置）：16，仅支持Ascend 950PR/Ascend 950DT。
maxValue	输出	Transpose接口能完成计算所需的最大临时空间大小，超出该值的空间不会被该接口使用。在最小临时空间-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。 maxValue仅作为参考值，有可能大于Unified Buffer剩余空间的大小，该场景下，开发者需要根据Unified Buffer剩余空间的大小来选取合适的临时空间大小。
minValue	输出	Transpose接口能完成计算所需最小临时空间大小。为保证功能正确，接口计算时预留/申请的临时空间不能小于该数值。

表2 GetTransposeTilingInfo接口参数列表

参数名	输入/输出	含义
srcShape	输入	输入的shape信息，参数类型为AscendC::TensorShape，具体srcShape传入格式为：场景1：[B, N, S, H/N] 场景2：[B, N, S, H/N] 场景3：[B, N, S, H/N] 场景4：[B, N, S, H/N] 场景5：[B, N, S, H/N] 场景6：[B, N, S, H/N] 场景7：[H, W] 场景13：[H, W]或者[N, H, W]，仅支持Ascend 950PR/Ascend 950DT。场景14：[N, H, W]，仅支持Ascend 950PR/Ascend 950DT。场景15：[N, H, W]，仅支持Ascend 950PR/Ascend 950DT。场景16 ：[H, W]，仅支持Ascend 950PR/Ascend 950DT。
stackBufferSize	输入	可供Transpose接口计算的空间大小，单位Byte。
typeSize	输入	输入的数据类型大小，单位为字节。比如输入的数据类型为half，此处应传入2。
transposeTypeIn	输入	选择数据排布及reshape的类型，根据输入数字选择对应的场景，参数范围为[1,7]。 Ascend 950PR/Ascend 950DT，该参数取值范围为[1, 7]和[13, 16]。场景1（NZ2ND，1、2轴互换）：1 场景2（NZ2NZ，1、2轴互换）：2 场景3（NZ2NZ，尾轴切分）：3 场景4（NZ2ND，尾轴切分）：4 场景5（NZ2ND，尾轴合并）：5 场景6（NZ2NZ，尾轴合并）：6 场景7（二维转置）：7 场景13 （二维转置或者三维转置中后两维转置）：13，仅支持Ascend 950PR/Ascend 950DT。场景14 （三维转置中第一维和第二维互换）：14，仅支持Ascend 950PR/Ascend 950DT。场景15 （三维转置中第一维和第三维互换）：15，仅支持Ascend 950PR/Ascend 950DT。场景16（使用交织指令进行两维ND2NZ转置）：16，仅支持Ascend 950PR/Ascend 950DT。
tiling	输出	输入数据的切分信息。

返回值说明

无

约束说明

无

调用示例

如下样例介绍了使用Transpose高阶API时host侧获取Tiling参数的流程以及该参数如何在kernel侧使用。样例中为场景1，输入Tensor的shape大小为[1, 2, 64, 32]，输入的数据类型为half。

将ConfusionTransposeTiling结构体参数增加至TilingData结构体，作为TilingData结构体的一个字段。

Text

BEGIN_TILING_DATA_DEF(TilingData)         // 注册一个tiling的类，以tiling的名字作为入参
    TILING_DATA_FIELD_DEF(uint32_t, tileNum); // 添加tiling字段，每个核上总计算数据分块个数
    ...                                       // 添加其他tiling字段
    TILING_DATA_FIELD_DEF_STRUCT(
        ConfusionTransposeTiling,
        confusionTransposeTilingData); // 将ConfusionTransposeTiling结构体参数增加至TilingData结构体
END_TILING_DATA_DEF;

Tiling实现函数中，首先调用GetTransposeMaxMinTmpSize接口获取Transpose接口能完成计算所需最大和最小临时空间大小，根据该范围结合实际的内存使用情况设置合适的空间大小；然后根据输入shape、剩余的可供计算的空间大小等信息获取Transpose kernel侧接口所需Tiling参数。

Text

namespace optiling {
const uint32_t NUM_BLOCKS = 8;
const uint32_t TILE_NUM = 8;
static ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    TilingData tiling;
    uint32_t totalLength = context->GetInputTensor(0)->GetShapeSize();
    context->SetSimdNumBlocks(NUM_BLOCKS);
    tiling.set_tileNum(TILE_NUM);
    // 设置其他Tiling参数
    ...
    std::vector<int64_t> shapeVec = {1, 2, 64, 32};
    AscendC::TensorShape srcShape(shapeVec);

    uint32_t transposeTypeIn = 1;
    uint32_t maxValue = 0;
    uint32_t minValue = 0;
    AscendC::GetTransposeMaxMinTmpSize(srcShape, sizeof(half), transposeTypeIn, maxValue, minValue);
    // 本样例为保证功能正确，传入最小临时空间大小，开发者可以根据需要传入合适的空间大小
    const uint32_t stackBufferSize = minValue;
    // 获取Transpose接口所需的Tiling参数
    AscendC::GetTransposeTilingInfo(
        srcShape, stackBufferSize, sizeof(half), transposeTypeIn, tiling.confusionTransposeTilingData);
    ...
    // 其他逻辑
    tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
    context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
    context->SetTilingKey(1);
    return ge::GRAPH_SUCCESS;
}
} // namespace optiling

对应的kernel侧通过在核函数中调用GET_TILING_DATA获取TilingData，继而将TilingData中的ConfusionTransposeTiling信息传入Transpose接口参与计算。完整的kernel侧样例请参考Transpose。

Text

extern "C" __global__ __aicore__ void func_custom(GM_ADDR src_gm, GM_ADDR dst_gm, GM_ADDR workspace, GM_ADDR tiling)
{
    GET_TILING_DATA(TilingData, tiling);
    KernelTranspose<half> op;
    op.Init(src_gm, dst_gm, TilingData.confusionTransposeTilingData);
    op.Process();
}

2.2.2. 数据结构

2.2.2.3. LocalTensor

2.2.2.4. GlobalTensor

2.2.3. 数据结构(TensorAPI，试验特性)

2.2.3.1. Layout数据结构

2.2.3.2. Tensor数据结构

2.2.3.3. 工具函数

2.2.4. 辅助数据结构

2.2.4.1. Coordinate

2.2.4.2. Layout

2.2.4.3. TensorTrait

2.2.4.6. TensorDesc

2.2.5. 数据搬运导览

2.2.5.1. 概览

2.2.6. 矩阵计算（ISASI）

2.2.6.1. 概述

2.2.6.2. 矩阵计算分形介绍

2.2.6.3. 矩阵计算的搬入

2.2.6.4. 矩阵搬入辅助配置接口

2.2.6.5. Mmad计算

2.2.6.6. Mmad计算关键特性说明

2.2.6.7. Mmad计算辅助配置接口

2.2.6.8. 矩阵计算的搬出

2.2.6.9. 矩阵搬出关键特性说明

2.2.6.10. 矩阵搬出辅助配置接口

2.2.7. 矩阵计算（Tensor API，试验特性）

2.2.7.2. 矩阵计算的搬入

2.2.7.3. Mmad计算

2.2.7.4. Mmad计算关键特性说明

2.2.7.6. 矩阵计算的搬出

2.2.7.7. 矩阵搬出关键特性说明

2.2.8. Memory矢量计算

2.2.8.2. SIMD计算说明

2.2.8.3. 数据搬运

2.2.8.4. 数据搬运辅助配置接口

2.2.8.5. 基础算术

2.2.8.6. 逻辑计算

2.2.8.7. 复合计算

2.2.8.8. 比较与选择

2.2.8.9. 类型转换

2.2.8.10. 类型转换辅助配置接口

2.2.8.11. 归约计算

2.2.8.12. 归约计算辅助配置接口

2.2.8.13. 数据排布转换

2.2.8.14. 数据填充

2.2.8.15. 排序组合（ISASI）

2.2.8.16. 离散与聚合

2.2.8.17. 掩码操作

2.2.8.18. 数据重排（ISASI）

2.2.9. Reg矢量计算

2.2.9.2. 寄存器数据类型

2.2.9.3. 辅助数据类型

2.2.9.4. VF调用

2.2.9.6. Reg数据搬入

2.2.9.7. Reg数据搬出

2.2.9.8. MaskReg计算

2.2.9.9. 基础算术

2.2.9.10. 逻辑计算

2.2.9.11. 复合计算

2.2.9.12. 比较与选择

2.2.9.13. 类型转换

2.2.9.14. 归约计算

2.2.9.15. 数据填充

2.2.9.16. 离散操作

2.2.9.17. 数据重排

2.2.9.18. 数据压缩

2.2.9.19. 直方图计算

2.2.9.20. 索引操作

2.2.9.21. 同步控制

2.2.9.22. 系统变量访问

2.2.10. 标量计算

2.2.11. 资源管理

2.2.11.1. TPipe

2.2.11.3. TBufPool

2.2.11.4. 自定义TBufPool

2.2.11.5. TQue

2.2.11.6. TSCM

2.2.11.7. TQueBind

2.2.11.8. TBuf

2.2.11.9. workspace