Tiling Information Obtained by the Kernel Is Incorrect

Symptom

An operator is used to add PRINTF printing to the kernel implementation code, but the tiling information obtained by the kernel is incorrect.

The added printing code is as follows:

PRINTF("tiling_data.totalLength: %d tiling_data.tileNum: %d.\n", tiling_data.totalLength, tiling_data.tileNum);

The printed tiling data is all 0s:

1

tiling_data.totalLength: 0 tiling_data.tileNum: 0.

Cause Analysis

The possible causes are as follows:

• The logic for computing tiling on the host is incorrect.
• The parameters of the kernel function on the kernel are not set in the correct sequence.

Solution

1. Print TilingData and check whether it is correctly serialized on the host as follows. If the printed value is incorrect, the tiling computation logic may be incorrect. In this case, check the tiling implementation code on the host to check whether the computation logic is correct.

   std::out<<*reinterpret_cast<uint32_t *>(context->GetRawTilingData()->GetData())<<std::endl; // Print the first parameter value of TilingData based on the actual data type. To confirm other values, offset the value pointer backward.

2. If TilingData printed in the previous step is correct, check whether the parameters of the kernel function on the kernel are set in the correct sequence.

   When you use the msOpGen tool to create an operator project and develop operators in the kernel based on the project, the definition template of the kernel function has been automatically generated by the msOpGen tool. The parameters are arranged in the sequence of "input, output, workspace, tiling". Check whether the parameter sequence is changed.

   #include "kernel_operator.h"
   extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) {
       GET_TILING_DATA(tiling_data, tiling);// Obtain tiling parameters.
       // TODO: user kernel impl
   }