文档

REGISTER_TILING_FOR_TILINGKEY

Function Usage

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

1
REGISTER_TILING_FOR_TILINGKEY(EXPRESSION, TILING_STRUCT)

Parameters

Parameter

Input/Output

Description

EXPRESSION

Input

Logical operation, where TILING_KEY_VAR refers to TilingKey.

TILING_STRUCT

Input

User-registered custom TilingData structure that matches TilingKey.

Availability

Atlas Training Series Product

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.
  • EXPRESSION currently supports bitwise operations: &, |, ~, and ^; shift operators: << and >>; arithmetic operations: +, –, *, /, and %; conditional operators: ==, !=, > , <, >= , and <=; and logical operators: &&, ||, and (). 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

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
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(0xFF)) {
        GET_TILING_DATA_WITH_STRUCT(optiling::TilingDataC, tilingData, tiling);
        ......
    } else {
        GET_TILING_DATA(tilingData, tiling);
        ......
    }
}
Parent topic: Kernel Tiling