GET_TILING_DATA
Supported Products
|
Product |
Supported/Unsupported |
|---|---|
|
|
√ |
|
|
√ |
|
|
√ |
|
|
√ |
|
|
√ |
|
|
x |
Functions
Obtains the tiling information passed by the kernel entry function of the operator and fills the information into the registered TilingData structure. This function is compiled in macro expansion mode. The TilingData structure needs to be defined in the corresponding operator implementation on the host to implement and register the Tiling function for TilingData computation. For details, see Tiling Implementation on the Host. If a user has registered multiple TilingData structures in TilingData Structure Registration, the default registered structures will be returned when this API is called.
Prototype
1
|
GET_TILING_DATA(tiling_data, tiling_arg) |
Parameters
|
Parameter |
Input/Output |
Description |
|---|---|---|
|
tiling_data |
Output |
Returns the default TilingData structure variable. |
|
tiling_arg |
Input |
Tiling parameter passed to the operator entry function. |
Constraints
- This function needs to be used in the operator kernel code, and the type of the input tiling_data parameter does not need to be declared.
- Currently, the kernel launch project is not supported.
Example
1 2 3 4 5 6 7 |
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *tiling) { GET_TILING_DATA(tilingData, tiling); KernelAdd op; op.Init(x, y, z, tilingData.blkDim, tilingData.totalSize, tilingData.splitTile); op.Process(); } |
1 2 3 4 5 6 7 8 9 10 11 12 |
ge::graphStatus TilingFunc(gert::TilingContext* context) { // Other code logic ... TilingData tiling; // Correspond to the TilingData structure defined in the operator host implementation. tiling.set_blkDim(blockDim); // Correspond to the members in the TilingData structure defined in the operator host implementation. tiling.set_totalSize(totalSize); tiling.set_splitTile(splitTile); tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity()); ... // Other code logic } |