REGISTER_TILING_FOR_TILINGKEY
产 品 支 持 情 况
功 能 说 明
用 于 在kernel侧 注 册 与TilingKey相 匹 配 的TilingData自 定 义 结 构 体;该 接 口 需 提 供 一 个 逻 辑 表 达 式,逻 辑 表 达 式 以 字 符 串“TILING_KEY_VAR”代 指 实 际TilingKey,表 达TilingKey所 满 足 的 范 围。
函 数 原 型
Text
REGISTER_TILING_FOR_TILINGKEY(EXPRESSION, TILING_STRUCT)
参 数 说 明
约 束 说 明
- 使 用 该 接 口 时,需 确 保 已 使 用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 >();
...
}
...
// 其 他 代 码 逻 辑
}