Skip to content
版 本

REGISTER_TILING_FOR_TILINGKEY

产 品 支 持 情 况

产 品

是 否 支 持

Ascend 950PR/Ascend 950DT

Atlas A3 训 练 系 列 产 品/Atlas A3 推 理 系 列 产 品

Atlas A2 训 练 系 列 产 品/Atlas A2 推 理 系 列 产 品

Atlas 200I/500 A2 推 理 产 品

Atlas 推 理 系 列 产 品AI Core

Atlas 推 理 系 列 产 品Vector Core

Atlas 训 练 系 列 产 品

x

功 能 说 明

用 于 在kernel侧 注 册 与TilingKey相 匹 配 的TilingData自 定 义 结 构 体;该 接 口 需 提 供 一 个 逻 辑 表 达 式,逻 辑 表 达 式 以 字 符 串“TILING_KEY_VAR”代 指 实 际TilingKey,表 达TilingKey所 满 足 的 范 围。

函 数 原 型

Text
REGISTER_TILING_FOR_TILINGKEY(EXPRESSION, TILING_STRUCT)

参 数 说 明

参 数

输 入/输 出

说 明

EXPRESSION

输 入

EXPRESSION为 逻 辑 运 算,其 中 用TILING_KEY_VAR指 代TilingKey。

TILING_STRUCT

输 入

用 户 注 册 的 与TilingKey相 匹 配 的TilingData自 定 义 结 构 体。

约 束 说 明

  • 使 用 该 接 口 时,需 确 保 已 使 用REGISTER_TILING_DEFAULT注 册 默 认 的 用 户 自 定 义TilingData结 构 体,用 于 告 知 框 架 侧 用 户 使 用 标 准C++语 法 来 定 义TilingData。
  • EXPRESSION当 前 支 持 位 运 算:&、|、~、^;移 位 运 算 符:<<、>>;算 术 运 算:+、-、*、/、%;条 件 运 算 符:==、!=、>、<、>=、<=;逻 辑 与&&、或||以 及()。优 先 级 同C++。
  • 若TilingData结 构 体 在 命 名 空 间 内,注 册 时 需 要 携 带 对 应 的 命 名 空 间 作 用 域 符。
  • 不 支 持 同 个TilingKey指 向 不 同TilingData结 构 体,会 出 现 拦 截 报 错。
  • 暂 不 支 持kernel直 调 工 程。

调 用 示 例

Text
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *tiling)
{
    REGISTER_TILING_DEFAULT(optiling::TilingData);  // 注 册 用 户 默 认 自 定 义TilingData结 构 体
    REGISTER_TILING_FOR_TILINGKEY("TILING_KEY_VAR == 1", optiling::TilingDataA); // 注 册TilingKey为1的TilingData结 构 体
    REGISTER_TILING_FOR_TILINGKEY("(TILING_KEY_VAR >= 10) && (TILING_KEY_VAR <= 15)", optiling::TilingDataB); // 注 册TilingKey在[10,15]之 间 的TilingData结 构 体
    REGISTER_TILING_FOR_TILINGKEY("TILING_KEY_VAR & 0xFF", optiling::TilingDataC); // 注 册TilingKey低16位 为1的TilingData结 构 体
    if (TILING_KEY_IS(1)) {
        GET_TILING_DATA_WITH_STRUCT(optiling::TilingDataA, tilingData, tiling);
        ......
    } else if (TILING_KEY_IS(11)) {
        GET_TILING_DATA_WITH_STRUCT(optiling::TilingDataB, tilingData, tiling);
        ......
    } else if (TILING_KEY_IS(14)) {
        GET_TILING_DATA_WITH_STRUCT(optiling::TilingDataB, tilingData, tiling);
        ......
    } else if (TILING_KEY_IS(255)) {
        GET_TILING_DATA_WITH_STRUCT(optiling::TilingDataC, tilingData, tiling);
        ......
    } else {
        GET_TILING_DATA(tilingData, tiling);
        ......
    }
}

使 用 标 准C++语 法 注 册tiling结 构 体:

Text
class TilingDataA{
public:
    ...
};
class TilingDataB{
public:
    ...
};
class TilingDataC{
public:
    ...
};

配 套 的host侧tiling函 数 示 例:

Text
ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    // 其 他 代 码 逻 辑
    ...
    if(condition1){
        context->SetTilingKey(1);
        optiling::TilingDataA *Addtiling = context->GetTilingData<optiling::TilingDataA>();
        ...
    } else if (condition2){
        context->SetTilingKey(11);
        optiling::TilingDataB *Addtiling = context->GetTilingData<optiling::TilingDataB >();
        ...
    } else if (condition3){
        context->SetTilingKey(14);
        optiling::TilingDataB *Addtiling = context->GetTilingData<optiling::TilingDataB >();
        ...
    } else if (condition4){
        context->SetTilingKey(255);
        optiling::TilingDataC *Addtiling = context->GetTilingData<optiling::TilingDataC >();
        ...
    }
    ...
    // 其 他 代 码 逻 辑
}

免 责 声 明:本 站 内 容 由 asc-devkit 仓 master 分 支 自 动 编 译 生 成,属 于 持 续 开 发 版 本,可 能 存 在 缺 陷,仅 供 预 览 与 参 考。如 需 稳 定 及 商 用 资 料,请 查 阅 官 方 昇 腾 社 区