GET_TILING_DATA_PTR_WITH_STRUCT

Product Support

Function

Allows developers to retrieve the corresponding tiling information by specifying the structure name and fill data into the respective tiling structure. Returns a pointer to the tiling structure after information filling and uses the __tiling_data_ptr__ modifier to modify the pointer. This modification method ensures 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 needs to be used in the operator kernel code, and the type of the input dst_ptr parameter does not need to be declared.
• In the dynamic shape scenario, the retrieved dst_ptr is a pointer to the global memory variable. In the static shape scenario, the retrieved dst_ptr is a pointer to the local variable. Ensure that the pointer is used within a valid scope.
• Currently, the kernel launch project is not supported.

Example

 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. Illegal access to a local variable whose lifecycle has ended.

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