GET_TILING_DATA_WITH_STRUCT
函数功能
使用该接口指定结构体名称,可获取指定的tiling信息,并填入对应的Tiling结构体中,此函数会以宏展开的方式进行编译。与GET_TILING_DATA的区别是:GET_TILING_DATA只能获取默认注册的结构体,该接口可以根据指定的结构体名称获取对应的结构体,常用于针对不同的TilingKey注册了不同结构体的情况下。
函数原型
1
|
GET_TILING_DATA_WITH_STRUCT(struct_name, tiling_data, tiling_arg) |
参数说明
参数 |
输入/输出 |
说明 |
---|---|---|
struct_name |
输入 |
指定的结构体名称。 |
tiling_data |
输出 |
返回指定Tiling结构体变量。 |
tiling_arg |
输入 |
此参数为算子入口函数处传入的tiling参数。 |
支持的型号
Atlas 训练系列产品
Atlas推理系列产品(Ascend 310P处理器)AI Core
Atlas A2训练系列产品/Atlas 800I A2推理产品
Atlas 200/500 A2推理产品
约束说明
- 本函数需在算子kernel代码处使用,并且传入的tiling_data参数不需要声明类型。
- 暂不支持kernel直调工程。
调用示例
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 |
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; if (TILING_KEY_IS(1)) { GET_TILING_DATA_WITH_STRUCT(Add_Struct_Special, tilingData, tiling); // 使用算子指定注册的结构体 op.Init(x, y, z, tilingData.totalLengthSpecial, tilingData.tileNumSpecial); } else { GET_TILING_DATA(tilingData, tiling); // 使用算子默认注册的结构体 op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum); } if (TILING_KEY_IS(1)) { op.Process(); } else if (TILING_KEY_IS(2)) { op.Process(); } else if (TILING_KEY_IS(3)) { op.Process(); } } |
父主题: Kernel Tiling