Skip to content
版 本

GET_TILING_DATA

产 品 支 持 情 况

产 品

是 否 支 持

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入 口 函 数 传 入 的Tiling信 息,并 填 入 注 册 的TilingData结 构 体 中,此 函 数 会 以 宏 展 开 的 方 式 进 行 编 译。对 应 的 算 子host实 现 中 需 要 定 义TilingData结 构 体,实 现 并 注 册 计 算TilingData的Tiling函 数。如 果 用 户 通 过TilingData结 构 注 册注 册 了 多 个TilingData结 构 体,使 用 该 接 口 返 回 默 认 注 册 的 结 构 体。

函 数 原 型

Text
GET_TILING_DATA(tiling_data, tiling_arg)

参 数 说 明

参 数

输 入/输 出

说 明

tiling_data

输 出

返 回 默 认TilingData结 构 体 变 量。

tiling_arg

输 入

此 参 数 为 算 子 入 口 函 数 处 传 入 的Tiling参 数。

约 束 说 明

  • 本 函 数 需 在 算 子kernel代 码 处 使 用,并 且 传 入 的tiling_data参 数 不 需 要 声 明 类 型。
  • 暂 不 支 持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)
{
    GET_TILING_DATA(tilingData, tiling);// 反 序 列 化SaveToBuffer生 成 的 数 据,并 填 入 注 册 的TilingData结 构 体 中
    KernelAdd op;
    op.Init(x, y, z, tilingData.numBlocks, tilingData.totalSize, tilingData.splitTile);
    op.Process();
}

配 套 的host侧Tiling函 数 示 例:

Text
ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    // 其 他 代 码 逻 辑
    ...
    TilingData tiling;  // 与 算 子host实 现 中 定 义TilingData结 构 体 的 对 应
    tiling.set_blkDim(numBlocks);  // 与 算 子host实 现 中 定 义TilingData结 构 体 中 的 成 员 的 对 应
    tiling.set_totalSize(totalSize);
    tiling.set_splitTile(splitTile);
    tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
    ...
    // 其 他 代 码 逻 辑
}

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