GET_TILING_DATA

Supported Products

Product

Supported/Unsupported

Atlas A3 training products / Atlas A3 inference products

Atlas A2 training products / Atlas A2 inference products

Atlas 200I/500 A2 inference products

Atlas inference product 's AI Core

Atlas inference product 's Vector Core

Atlas training products

x

Functions

Obtains the tiling information passed by the kernel entry function of the operator and fills the information into the registered TilingData structure. 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

Parameter

Input/Output

Description

tiling_data

Output

Returns the default TilingData structure variable.

tiling_arg

Input

Tiling parameter passed to the operator entry function.

Constraints

  • 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);
    KernelAdd op;
    op.Init(x, y, z, tilingData.blkDim, tilingData.totalSize, tilingData.splitTile);
    op.Process();
}
The following is an example of the tiling function on the host:
 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(blockDim);  // 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
}