COPY_TILING_WITH_ARRAY

Product Support

Function

Copies the content of an array of a specified size to the destination array and returns a pointer to the copied array. This macro can be used to array member variables within a structure. It copies the specified array to the stack, making it ideal for scenarios with frequent access to tiling data, thereby accelerating data access speed.

Prototype

1

COPY_TILING_WITH_ARRAY(arr_type, arr_count, 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 array that the dst_ptr pointer retrieved by this macro points to is a 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
11
12
13
14
15
16

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 {
        COPY_TILING_WITH_ARRAY(uint64_t, 2, tilingDataPtr->vectorTilingArray, vTilingArrayPtr);        
        op.Init(x, y, z, (*vTilingArrayPtr)[0], (*vTilingArrayPtr)[1]);
        op.Process();
    } else {
        COPY_TILING_WITH_ARRAY(uint64_t, 2, tilingDataPtr->cubeTilingArray, cTilingArrayPtr);
	op.Init(x, y, z, (*cTilingArrayPtr)[0], (*cTilingArrayPtr)[1]);
	op.Process();
    }
}