GET_TILING_DATA

Applicability

Function Usage

Retrieves tiling information passed to a kernel entry point function and fills a registered TilingData structure with this data. This function is compiled in macro expansion mode. The TilingData structure needs to be defined in the corresponding operator implementation on the host to implement and register the Tiling function for TilingData computation. For details, see Tiling Implementation on the Host. If a user has registered multiple TilingData structures in TilingData Structure Registration, the default registered structures will be returned when this API is called.

Prototype

1

GET_TILING_DATA(tiling_data, tiling_arg)

Parameters

Restrictions

• This function needs to be used in the operator kernel code, and the type of the input tiling_data parameter does not need to be declared.
• Currently, the kernel launch project is not supported.

Example

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);// Deserialize the data generated by SaveToBuffer and fill it in the registered TilingData structure.
    KernelAdd op;
    op.Init(x, y, z, tilingData.numBlocks, tilingData.totalSize, tilingData.splitTile);
    op.Process();
}

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

ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    // Other code logic
    ...
    TilingData tiling;  // Correspond to the TilingData structure defined in the operator host implementation.
    tiling.set_blkDim(numBlocks);  // Correspond to the members in the TilingData structure defined in the operator host implementation.
    tiling.set_totalSize(totalSize);
    tiling.set_splitTile(splitTile);
    tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
    ...
    // Other code logic
}