GET_TILING_DATA
Product Support
|
Product |
Supported |
|---|---|
|
|
√ |
|
|
√ |
|
|
√ |
|
|
√ |
|
|
√ |
|
|
x |
Function
Retrieves tiling information passed to a kernel entry point function and fills a registered TilingData structure with this data. 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 |
Default TilingData structure variable to be returned. |
|
tiling_arg |
Input |
Tiling parameter input by the entrypoint function of the operator. |
Restrictions
- 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. } |