COPY_TILING_WITH_STRUCT

Product Support

Function

Copies the tiling structure and returns a pointer to the copied tiling structure. This macro applies to nested structures and can be used to copy substructure member variables of a structure. This macro copies a specified structure to the stack, making it ideal for scenarios with frequent access to tiling data, thereby accelerating data access speed.

Prototype

1

COPY_TILING_WITH_STRUCT(tiling_struct, src_ptr, dst_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.
• This macro must be used together with GET_TILING_DATA_PTR_WITH_STRUCT. The input src_ptr parameter is the pointer retrieved by GET_TILING_DATA_PTR_WITH_STRUCT.
• The tiling structure that the dst_ptr pointer retrieved by this macro points to is a local variable. Ensure that the tiling structure 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
11
12
13
14
15

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);

    if ASCEND_IS_AIV {
        Using VectorTilingStruct on the COPY_TILING_WITH_STRUCT(VectorTilingStruct, tilingDataPtr->vectorTiling, vTilingStructPtr); // VectorTilingStruct used on the Vector side.
        op.Init(x, y, z, vTilingStructPtr->totalLength, vTilingStructPtr->tileNum);
        op.Process();
    } else {
        Using CubeTilingStruct on the COPY_TILING_WITH_STRUCT(CubeTilingStruct, tilingDataPtr->cubeTiling, cTilingStructPtr); // CubeTilingStruct used on the Cube side.
	op.Init(x, y, z, *cTilingStructPtr);
        op.Process();
    }
}