REGISTER_TILING_FOR_TILINGKEY

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

Registers a custom TilingData structure that matches the TilingKey on the kernel. This API needs to provide a logical expression, which uses the string TILING_KEY_VAR to indicate the actual TilingKey and the range that the TilingKey meets.

Prototype

REGISTER_TILING_FOR_TILINGKEY(EXPRESSION, TILING_STRUCT)

Parameters

Parameter	Input/Output	Description
EXPRESSION	Input	EXPRESSION is a logical operation, where TILING_KEY_VAR refers to TilingKey.
TILING_STRUCT	Input	User-registered custom TilingData structure that matches TilingKey.

Constraints

Before using this API, ensure that the default custom TilingData structure has been registered using REGISTER_TILING_DEFAULT to instruct the framework user to use the standard C++ syntax to define TilingData.
Currently, EXPRESSION supports the following operators: bitwise operators (&, |, ~, and ^), shift operators (<< and >>), arithmetic operators (+, -, *, /, and %), relational operators (> and <), logical operators (&& and ||), and parentheses. The priority is the same as that of C++.
If the TilingData structure is in the namespace, the corresponding namespace scope resolution operator must be carried during registration.
The same TilingKey cannot point to different TilingData structures. Otherwise, an interception error is reported.
Currently, the kernel launch project is not supported.

Example

extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *tiling)
{
    REGISTER_TILING_DEFAULT(optiling::TilingData);  // Register the default custom TilingData structure defined by the user.
    REGISTER_TILING_FOR_TILINGKEY("TILING_KEY_VAR == 1", optiling::TilingDataA); // Register the TilingData structure with TilingKey being 1.
    REGISTER_TILING_FOR_TILINGKEY("(TILING_KEY_VAR >= 10) && (TILING_KEY_VAR <= 15)", optiling::TilingDataB); // Register the TilingData structure with TilingKey between 10 and 15.
    REGISTER_TILING_FOR_TILINGKEY("TILING_KEY_VAR & 0xFF", optiling::TilingDataC); // Register the TilingData structure whose lower 16 bits of TilingKey are 1.
    if (TILING_KEY_IS(1)) {
        GET_TILING_DATA_WITH_STRUCT(optiling::TilingDataA, tilingData, tiling);
        ......
    } else if (TILING_KEY_IS(11)) {
        GET_TILING_DATA_WITH_STRUCT(optiling::TilingDataB, tilingData, tiling);
        ......
    } else if (TILING_KEY_IS(14)) {
        GET_TILING_DATA_WITH_STRUCT(optiling::TilingDataB, tilingData, tiling);
        ......
    } else if (TILING_KEY_IS(255)) {
        GET_TILING_DATA_WITH_STRUCT(optiling::TilingDataC, tilingData, tiling);
        ......
    } else {
        GET_TILING_DATA(tilingData, tiling);
        ......
    }
}

class TilingDataA{
public:
    ...
};
class TilingDataB{
public:
    ...
};
class TilingDataC{
public:
    ...
};

The following is an example of a matching tiling function on the host:

ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    // Other code logic
    ...
    if(condition1){
        context->SetTilingKey(1);
        optiling::TilingDataA *Addtiling = context->GetTilingData<optiling::TilingDataA>();
        ...
    } else if (condition2){
        context->SetTilingKey(11);
        optiling::TilingDataB *Addtiling = context->GetTilingData<optiling::TilingDataB >();
        ...
    } else if (condition3){
        context->SetTilingKey(14);
        optiling::TilingDataB *Addtiling = context->GetTilingData<optiling::TilingDataB >();
        ...
    } else if (condition4){
        context->SetTilingKey(255);
        optiling::TilingDataC *Addtiling = context->GetTilingData<optiling::TilingDataC >();
        ...
    }
    ...
    // Other code logic
}

Parent topic: Kernel Tiling