GET_TILING_DATA_PTR_WITH_STRUCT

产 品 支 持 情 况

产 品	是 否 支 持
Ascend 950PR/Ascend 950DT	√
Atlas A3 训 练 系 列 产 品/Atlas A3 推 理 系 列 产 品	√
Atlas A2 训 练 系 列 产 品/Atlas A2 推 理 系 列 产 品	√
Atlas 200I/500 A2 推 理 产 品	√
Atlas 推 理 系 列 产 品AI Core	√
Atlas 推 理 系 列 产 品Vector Core	√
Atlas 训 练 系 列 产 品	x

功 能 说 明

在 使 用 该 宏 时，开 发 者 可 以 通 过 指 定 结 构 体 名 称 来 获 取 相 应 的Tiling信 息，并 将 其 填 入 对 应 的Tiling结 构 体 中。完 成 填 充 后，该 宏 将 返 回 一 个 指 向 该Tiling结 构 体 的 指 针，并 使 用__tiling_data_ptr__修 饰 符 对 该 指 针 进 行 修 饰。这 种 修 饰 方 式 能 够 确 保 在 动 静 态Shape场 景 下 代 码 的 统 一 性 和 兼 容 性。

函 数 原 型

Text

GET_TILING_DATA_PTR_WITH_STRUCT(tiling_struct, dst_ptr, tiling_ptr)

参 数 说 明

参 数	输 入/输 出	说 明
tiling_struct	输 入	指 定 的 结 构 体 名 称。
dst_ptr	输 出	返 回 指 定 的Tiling结 构 体 指 针。
tiling_ptr	输 入	算 子 入 口 函 数 处 传 入 的Tiling参 数。

约 束 说 明

• 该 宏 需 在 算 子Kernel代 码 处 使 用，并 且 传 入 的dst_ptr参 数 无 需 声 明 类 型。
• 动 态Shape场 景 下，获 取 到 的dst_ptr是 指 向Global Memory变 量 的 指 针；静 态Shape场 景 下，获 取 到 的dst_ptr是 指 向 局 部 变 量 的 指 针，需 确 保 在 合 理 的 作 用 域 范 围 内 使 用。
• 暂 不 支 持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)
{
    KernelAdd op;

    GET_TILING_DATA_PTR_WITH_STRUCT(AddCustomTilingData, tilingDataPtr, tiling);

    op.Init(x, y, z, tilingDataPtr->totalLength, tilingDataPtr->tileNum);
    op.Process();
    
}

以 下 是 错 误 调 用 的 示 例：

Text

__aicore__ __tiling_data_ptr__ AddCustomTilingData* foo(__gm__ uint8_t *tiling)
{
    GET_TILING_DATA_PTR_WITH_STRUCT(AddCustomTilingData, tilingDataPtr, tiling);
    return tilingDataPtr;
}

extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *tiling)
{
    KernelAdd op;

    auto tilingDataPtr = foo(tiling);  // 错 误，foo函 数 已 经 执 行 完 成，非 法 访 问 生 命 周 期 已 经 结 束 的 局 部 变 量

    op.Init(x, y, z, tilingDataPtr->totalLength, tilingDataPtr->tileNum);
    op.Process();
}