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 about the complete sample, see 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.

Restrictions

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

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 the 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 entrypoint 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.
    ...
}

Parent topic: Tiling Implementation on the Host