GET_TILING_DATA
产 品 支 持 情 况
功 能 说 明
用 于 获 取 算 子kernel入 口 函 数 传 入 的Tiling信 息,并 填 入 注 册 的TilingData结 构 体 中,此 函 数 会 以 宏 展 开 的 方 式 进 行 编 译。对 应 的 算 子host实 现 中 需 要 定 义TilingData结 构 体,实 现 并 注 册 计 算TilingData的Tiling函 数。如 果 用 户 通 过TilingData结 构 注 册注 册 了 多 个TilingData结 构 体,使 用 该 接 口 返 回 默 认 注 册 的 结 构 体。
函 数 原 型
Text
GET_TILING_DATA(tiling_data, tiling_arg)
参 数 说 明
约 束 说 明
- 本 函 数 需 在 算 子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());
...
// 其 他 代 码 逻 辑
}