GET_TILING_DATA_PTR_WITH_STRUCT

Supported Products

Functions

When using this macro, developers can obtain the corresponding tiling information by specifying the structure name and fill it into the corresponding tiling structure. After the filling is complete, the macro returns a pointer to the tiling structure and uses the __tiling_data_ptr__ modifier to modify the pointer. This modification method ensures the code consistency and compatibility in both dynamic and static shape scenarios.

Prototype

1

GET_TILING_DATA_PTR_WITH_STRUCT(tiling_struct, dst_ptr, tiling_ptr)

Parameters

Restrictions

• This macro must be used in the operator kernel code, and the input dst_ptr parameter does not need to be declared with a type.
• In the dynamic shape scenario, the obtained dst_ptr is a pointer to the global memory variable. In the static shape scenario, the obtained dst_ptr is a pointer to the local variable. Ensure that the pointer is used within a proper scope.
• Currently, the kernel launch project is not supported.

Examples

 1
 2
 3
 4
 5
 6
 7
 8
 9
10

extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *tiling)
{
    KernelAdd op;

    GET_TILING_DATA_PTR_WITH_STRUCT(AddCustomTilingData, tilingDataPtr, tiling);

    op.Init(x, y, z, tilingDataPtr->totalLength, tilingDataPtr->tileNum);
    op.Process();
    
}

The following is an example of incorrect calling:

__aicore__ __tiling_data_ptr__ AddCustomTilingData* foo(__gm__ uint8_t *tiling)
{
    GET_TILING_DATA_PTR_WITH_STRUCT(AddCustomTilingData, tilingDataPtr, tiling);
    return tilingDataPtr;
}

extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *tiling)
{
    KernelAdd op;

    auto tilingDataPtr = foo(tiling); // Error. The foo function has been executed, and the lifetime of the local variable that has been accessed is ended.

    op.Init(x, y, z, tilingDataPtr->totalLength, tilingDataPtr->tileNum);
    op.Process();
}