Documentation

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

Procedure

When defining the 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 as follows: For details about the complete example, click sample of defining the Tiling structure using standard C++ syntax.

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

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

    When using the Tiling structure of the 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.
    • The header file that defines the Tiling structure must be included.
    • 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" // Include the header file that defines the Tiling structure.
...

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. The tiling structure is registered on the kernel side, and the tiling data is parsed to the TilingData structure and used. 
     1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
    		 
    #include "kernel_operator.h"
#include "matmul_custom_tiling.h" // Include the header file that defines the tiling structure.

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 defining the tiling structure using standard C++ syntax

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];
};
  • Operators can define tiling structures with the same name but different structures. You can distinguish them by referencing the corresponding header files. This method allows each operator to use a tiling structure definition that meets its own requirements, without conflicting with other operators.

    In contrast, when macros such as BEGIN_TILING_DATA_DEF are used to define tiling structs with the same name but different structures, these structs are registered with the global tiling struct management variable. As a result, the tiling struct actually used by the current operator may fail to be obtained when the struct name is accessed, leading to undefined behavior.

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

    Operator B

    class TilingData {
public:
    uint32_t length;
    uint16_t coreNum;
};
  • Customized tiling value assignment is supported. You can directly assign values by accessing the member variables of the tiling struct, or customize the tiling value assignment function. (In macro definition mode, you can only assign values or access the tiling struct through 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" // Include the header file that defines the tiling struct.
...

namespace optiling {
static void ComputeTiling(TilingData* tiling, ...)
{
    // Compute 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 the tiling structure, the following restrictions apply:

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

    __aicore__ funcA() { ... } // is incorrect. The __aicore__ modifier is not supported during compilation on the host side, and a compilation error occurs.
    void func() { ... } // is incorrect. The __aicore__ modifier is missing on the device side, and the function cannot be executed.
};
  • The member variables of the tiling structure do not support pointer or reference types. These data types will cause data parsing exceptions 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 and 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 obtained by GetTilingData does not contain the initial value. You need to explicitly assign a value to the tiling, or define the tiling in the tiling structure and call the tiling assignment function. 
     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, the values of totalLength and tileNum are 0, without the initial values.
    ......
    //Explicit assignment is required.
    tiling->totalLength = totalLength; // Assign the member variable of the tiling struct.
    tiling->tileNum = TILE_NUM; // Assign the member variable of the tiling struct.
    ......
    return ge::GRAPH_SUCCESS;
}

How Do I Change the Macro-Defined Tiling Structure to the Standard C++ Syntax?

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

  1. First, move the header file that defines the tiling structure from the op_host directory to the op_kernel directory. The following shows the comparison of the content before and after the change. Note that the included header files are changed, and the header files related to macro definitions are no longer required.
    Table 1 Comparison between the two ways

    Macro definition

    Standard C++ syntax definition

    		 
    #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. Then, 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 by the macro definition.
    Table 2 Comparison between the two ways

    Macro definition

    Standard C++ syntax definition

    		 
    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 in macro definition mode is required.

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

    return ge::GRAPH_SUCCESS;
}
} // namespace optiling
    		 
    #include "../op_kernel/matmul_custom_tiling.h" // Include the header file that defines the tiling structure.
...

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; // Use the user-friendly C++ pointer to assign values to member variables.

    ...

    return ge::GRAPH_SUCCESS;
}
} // namespace optiling
  3. Finally, add the REGISTER_TILING_DEFAULT call at the kernel function entry to register the Tiling structure. This registration operation is used to inform the framework that the user has defined the Tiling structure using the standard C++ syntax and specified 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