GET_TILING_DATA

Product Support

Product

Supported

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

Function

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

Parameter

Input/Output

Description

tiling_data

Output

Default TilingData structure variable to be returned.

tiling_arg

Input

Tiling parameter input by the entrypoint function of the operator.

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);
    KernelAdd op;
    op.Init(x, y, z, tilingData.blkDim, tilingData.totalSize, tilingData.splitTile);
    op.Process();
}
The following is an example of a matching 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.
}