GET_TILING_DATA

函数功能

用于获取算子kernel入口函数传入的tiling信息，并填入注册的Tiling结构体中，此函数会以宏展开的方式进行编译。对应的算子host实现中需要定义TilingData结构体，实现并注册计算TilingData的Tiling函数，具体请参考Host侧tiling实现。如果用户通过TilingData结构注册注册了多个TilingData结构体，使用该接口返回默认注册的结构体。

函数原型

1

GET_TILING_DATA(tiling_data, tiling_arg)

参数说明

支持的型号

Atlas 训练系列产品

Atlas 推理系列产品 AI Core

Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件

Atlas 200I/500 A2 推理产品

约束说明

• 本函数需在算子kernel代码处使用，并且传入的tiling_data参数不需要声明类型。
• 暂不支持kernel直调工程。

调用示例

1
2
3
4
5
6
7

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);
    KernelAdd op;
    op.Init(x, y, z, tilingData.blkDim, tilingData.totalSize, tilingData.splitTile);
    op.Process();
}

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12

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