Documentation
Download

Using the Standard C++ Syntax to Define the Tiling Structure

Procedure

When defining a tiling structure, you can use the standard C++ syntax to define a Plain Old Data (POD) type, that is, a data type compatible with the C language. The procedure is described below. For details, refer to the sample of defining the Tiling structure using standard C++ syntax.

  1. Use the C++ syntax to define the tiling structure.

    The header file of the structure definition must be stored in the op_kernel directory of the operator project. Only files in this directory are packed into the operator package for online compilation. If required files are stored in other directories, online compilation may fail because related files cannot be found.

    When using the tiling structure of a high-level API, you can reference the predefined tiling structure in kernel_tiling/kernel_tiling.h through the AscendC::tiling namespace, as shown in the following code. 
     1
 2
 3
 4
 5
 6
 7
 8
 9
10
    		 
    #ifndef MATMUL_CUSTOM_TILING_H
#define MATMUL_CUSTOM_TILING_H
#include <cstdint>
#include "kernel_tiling/kernel_tiling.h"    // for TCubeTiling

struct MatmulCustomTilingData {
    uint64_t localMemSize;
    AscendC::tiling::TCubeTiling cubeTilingData;
};
#endif  // MATMUL_CUSTOM_TILING_H
  2. Assign values to the tiling structure in the tiling function on the host.
    • Contain the header file of the tiling structure definition.
    • Call GetTilingData to obtain the pointer to the tiling structure and assign values to its member variables.
     1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
    		 
    #include "../op_kernel/matmul_custom_tiling.h" // Contain the header file of the tiling structure definition.
...

namespace optiling {
static ge::graphStatus TilingFunc(gert::TilingContext *context)
{
    ...
    MultiCoreMatmulTiling cubeTiling(ascendcPlatform);
    ...
    // Obtain the pointer to the tiling structure.
    MatmulCustomTilingData *tiling = context->GetTilingData<MatmulCustomTilingData>();
    // Assign values to the member variables of the tiling structure.
    if (cubeTiling.GetTiling(tiling->cubeTilingData) == -1) {
        return ge::GRAPH_FAILED;
    }
    uint64_t localMemSize;
    ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, localMemSize);
    tiling->localMemSize = localMemSize;
    ...
    return ge::GRAPH_SUCCESS;
}
} // namespace optiling
  3. Register the tiling structure on the kernel and parse the tiling data to the TilingData structure for use. 
     1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
    		 
    #include "kernel_operator.h"
#include "matmul_custom_tiling.h" // Contain the header file of the tiling structure definition.

extern "C" __global__ __aicore__ void matmul_custom(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace, GM_ADDR tiling)
{
    REGISTER_TILING_DEFAULT(MatmulCustomTilingData);
    GET_TILING_DATA(tilingData, tiling);
    MatmulKernel<half, half, float, float> matmulKernel;
    AscendC::TPipe pipe;
    REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulKernel.matmulObj, &tilingData.cubeTilingData); // Initialize the matmul object.
    matmulKernel.Init(a, b, bias, c, workspace, tilingData.localMemSize, tilingData.cubeTilingData);
    ...
}

Advantages of Using the Standard C++ Syntax to Define the Tiling Structure

Compared with the method of using macros such as BEGIN_TILING_DATA_DEF for definition, this method is more in line with the development habits of C++ developers and provides powerful flexibility.

  • The bool type, array, structure array, and list initialization are supported. 
    class A {
public:
    bool xxx;
    uint32_t xxx[2][128] = {0};
};

class B {
public:
    bool xxx = false;
    uint8_t xxx[2][2]{0};
    A[10];
};
  • Different operators can define tiling structures with the same name but different layouts, which can be distinguished by referencing the corresponding header files in the operators. This method allows each operator to use a tiling structure definition tailored to its own requirements, without conflicting with other operators.

    In contrast, when macros such as BEGIN_TILING_DATA_DEF are used to define tiling structures with the same name but different layouts, these structures are registered with a global tiling structure management variable. In this case, accessing a structure by name may fail to retrieve the one intended for the current operator, leading to undefined behavior.

    Operator A: 
    class TilingData {
public:
    uint32_t length;
};

    Operator B:

    class TilingData {
public:
    uint32_t length;
    uint16_t coreNum;
};
  • Custom tiling value assignment is supported. You can directly assign values by accessing the member variables of the tiling structure, or customize a tiling value assignment function. (In macro definition mode, you can only assign or access the values using the set_xx/get_xx method generated by the framework.)
    Tiling structure definition 
    class TilingData {
public:
    uint32_t xxx1;
    uint32_t xxx2;
    uint8_t xxx3;
    bool xxx4;
};

    Tiling function on the host:

    #include "../op_kernel/xxx_custom_tiling.h" // Contain the header file of the tiling structure definition.
...

namespace optiling {
static void ComputeTiling(TilingData* tiling, ...)
{
    // Calculate the tiling logic.
    ...
    tiling->xxx1 = xxx;
    tiling->xxx2 = xxx;
    tiling->xxx3 = xxx;
    tiling->bool = xxx;
}
static ge::graphStatus TilingFunc(gert::TilingContext *context)
{    
    ...
    TilingData *tiling = context->GetTilingData<TilingData>();
    ...
    ComputeTiling(tiling, ...)
    ...

    return ge::GRAPH_SUCCESS;
}
} // namespace optiling

Restrictions

When the standard C++ syntax is used to define a tiling structure, the following restrictions apply:

  • Member functions cannot be defined in the tiling structure because the member functions on the device are different from those on the host (functions on the device require the __aicore__ modifier). However, the tiling structure is shared by the device and host. As a result, an error may occur during compilation or execution. 
    class TilingData {
public:
    uint32_t xxx;

    __aicore__ funcA() { ... } // Error. The __aicore__ modifier is not supported during compilation on the host, resulting in a compilation error.
    void func() { ... } // Error. The __aicore__ modifier is missing on the device, resulting in execution failure.
};
  • The member variables of the tiling structure do not support pointer or reference types. These data types can cause parsing exceptions when data is transferred from the host to the device. 
    1
2
3
4
5
    		 
    class TilingData {
public:
    uint32_t* totalLength; // The pointer type is not supported. The host cannot pass the pointer to the device.
    uint32_t& tileNum;       // The reference type is not supported. The host cannot pass the pointer to the device.
};
  • The tiling structure only supports the POD type. It does not support the object-oriented features such as virtual functions and virtual inheritance, or the template class. 
     1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
    		 
    class A {
public:
    uint32_t totalLength;
    uint32_t tileNum;
};
class B: public A {
public:
    uint32_t xxx;
    uint32_t xxx;
};
static ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    // Incorrect use.
    B *tiling = context->GetTilingData<A>(); // Not supported. Unknown issues will be triggered.
    // Correct use.
    B *tiling = context->GetTilingData<B>();
    ......
    return ge::GRAPH_SUCCESS;
}
  • The tiling data obtained through GetTilingData does not contain any initial value. You need to explicitly assign values, or define and call a tiling value assignment function in the tiling structure. 
     1
 2
 3
 4
 5
 6
 7
 8
 9
10
    		 
    static ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    TilingData *tiling = context->GetTilingData<TilingData>(); // Obtain the tiling structure. In this case, totalLength and tileNum are 0, and no initial values exist.
    ......
    // Explicit value assignment is required.
    tiling->totalLength = totalLength; // Assign value to a member variable of the tiling structure.
    tiling->tileNum = TILE_NUM;         // Assign value to a member variable of the tiling structure.
    ......
    return ge::GRAPH_SUCCESS;
}

How to Change the Tiling Structure Definition Mode from Macros to Standard C++ Syntax

This section describes how to change the tiling structure definition mode from using macros such as BEGIN_TILING_DATA_DEF to using the standard C++ syntax.

  1. Move the header file of tiling structure definition from the op_host directory to the op_kernel directory. The following table compares the content before and after the change. Note that the included header file has been changed, and the header files related to macro definitions are no longer required.
    Table 1 Comparison between the two modes

    Macro

    Standard C++ Syntax

    		 
    #include "register/tilingdata_base.h"
#include "tiling/tiling_api.h" // The TCubeTiling structure is defined using macros.

namespace optiling {
BEGIN_TILING_DATA_DEF(MatmulCustomTilingData)
TILING_DATA_FIELD_DEF(uint64_t, localMemSize);
TILING_DATA_FIELD_DEF_STRUCT(TCubeTiling, cubeTilingData);
END_TILING_DATA_DEF;
REGISTER_TILING_DATA_CLASS(MatmulCustom, MatmulCustomTilingData)
} // namespace optiling
    		 
    #include <cstdint>
#include "kernel_tiling/kernel_tiling.h" // The TCubeTiling structure is defined using C++ syntax.

struct MatmulCustomTilingData {
    uint64_t localMemSize;
    AscendC::tiling::TCubeTiling cubeTilingData;
};
  2. Modify the tiling function implementation on the host. In this case, you can assign values to the member variables of the tiling structure using the C++ pointer assignment method, instead of using the set method generated through macros.
    Table 2 Comparison between the two modes

    Macro

    Standard C++ Syntax

    		 
    namespace optiling {
static ge::graphStatus TilingFunc(gert::TilingContext *context)
{
    ...
    MultiCoreMatmulTiling cubeTiling(ascendcPlatform);
    ...
    MatmulCustomTilingData tiling;
    if (cubeTiling.GetTiling(tiling.cubeTilingData) == -1) { // Get matmul tiling.
        return ge::GRAPH_FAILED;
    }

    uint64_t localMemSize;
    ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, localMemSize);
    tiling.set_localMemSize(localMemSize); // The set method generated through macros is required.

    ...
    // Save local variables to the context.
    tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());  
    ...

    return ge::GRAPH_SUCCESS;
}
} // namespace optiling
    		 
    #include "../op_kernel/matmul_custom_tiling.h" // Contain the header file of the tiling structure definition.
...

namespace optiling {
static ge::graphStatus TilingFunc(gert::TilingContext *context)
{
    ...
    MultiCoreMatmulTiling cubeTiling(ascendcPlatform);
    ...
    MatmulCustomTilingData *tiling = context->GetTilingData<MatmulCustomTilingData>();
    if (cubeTiling.GetTiling(tiling->cubeTilingData) == -1) {
        return ge::GRAPH_FAILED;
    }

    uint64_t localMemSize;
    ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, localMemSize);
    tiling->localMemSize = localMemSize; // Assign value to a member variable using user-friendly C++ pointers.

    ...

    return ge::GRAPH_SUCCESS;
}
} // namespace optiling
  3. Add the REGISTER_TILING_DEFAULT call at the entry point of the kernel function to register the tiling structure. This registration operation informs the framework that the tiling structure has been defined using the standard C++ syntax and specifies its type, so that the framework can correctly identify and use the structure during tiling data parsing.
    #include "matmul_custom_tiling.h"
...

extern "C" __global__ __aicore__ void matmul_custom(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace, GM_ADDR tiling)
{
    REGISTER_TILING_DEFAULT(MatmulCustomTilingData); // Add the REGISTER_TILING_DEFAULT call to register the tiling structure.
    ...
}
This site uses cookies to improve your browsing experience. By continuing to browse this site, you agree to our use of cookies.View Details