Kernel侧算子实现

自动生成kernel侧算子实现模板

在算子工程目录下的“op_kernel/xxx.cpp”文件中实现算子的核函数。核函数的定义模板已通过msOpGen工具自动生成，样例如下所示。注意这里参数的顺序按照“输入、输出、workspace、tiling”的顺序排布，开发者不要调整其顺序。

#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);// 获取Tiling参数，详见下文介绍
    // TODO: user kernel impl
}

算子原型定义中的输入和输出同名的情况下，自动生成的核函数中，输出参数增加ref后缀予以区分。示例如下：

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR x_ref, GM_ADDR workspace, GM_ADDR tiling) {
    ...
}

GET_TILING_DATA获取Tiling参数

提供GET_TILING_DATA，用于获取算子kernel入口函数传入的tiling信息，并填入注册的Tiling结构体中，此函数会以宏展开的方式进行编译。注意，对应的算子host实现中需要定义TilingData结构体，实现并注册计算TilingData的Tiling函数。具体请参考Host侧tiling实现。

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(tilingData, tiling);
    KernelAdd op;
    op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum);
    if (TILING_KEY_IS(1)) {
        op.Process();
    }
}

核函数内推导输入数据类型和格式

template<class T> func() {}
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
{
    DTYPE_X temp;
    func<DTYPE_Z>();
    if (FORMAT_Y == FORMAT_ND) {
        ...
    }
}

输出shape依赖计算的算子kernel侧实现

某些算子，比如NonZero（统计tensor中非零值的个数），计算完成前无法得知算子输出的shape信息，算子计算完成后才能获取。该类算子在原型定义时，需要使用OutputShapeDependOnCompute接口进行标识，同时在算子核函数中将实际输出shape写入到出参中，便于框架侧基于该信息进行输出内存的管理。

在核函数所有输出的最后增加一个GM_ADDR类型的输出参数，并在核函数计算完成后，将输出shape信息写入到该出参中。shape信息的排布格式如下，大小为n * (8 + 1)，每个元素的数据类型为uint64_t。其中n表示待刷新shape信息的输出个数，每个输出的shape信息都通过第1个元素来保存实际的shape维度（dim），后续的8个元素来保存具体每个维度的shape信息。

• 输出的顺序和原型定义中输出的顺序保持一致。
• 对于uint64_t的输出数据类型（对于tensor而言），需要将dim的uint32_t的高位设置为1，表示以uint64_t类型解析该tensor。

• 如下示例中，算子中有一个输出依赖计算得出，输出tensor的数据类型为uint32_t，计算完成后，得到输出的shape为（32, 64），出参shape_out用于存放该shape信息，值为（2, 32, 64）。代码示例如下：

  extern "C" __global__ __aicore__ void xxx_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR shape_out, GM_ADDR workspace, GM_ADDR tiling) {
  ...
      constexpr uint32_t SHAPEOUT_SIZE = 9;
      // 输出数据为2维([32, 64])，tensor类型为uint32_t
      GlobalTensor<uint64_t> shapeoutGlobal_uint32;
      shapeoutGlobal_uint32.SetGlobalBuffer((__gm__ uint64_t*)shape_out, SHAPEOUT_SIZE);
      shapeoutGlobal_uint32.SetValue(0, 2);
      shapeoutGlobal_uint32.SetValue(1, 32);
      shapeoutGlobal_uint32.SetValue(2, 64);
  ...
  }

• 如下示例中，算子中有一个输出依赖计算得出，输出tensor的数据类型为uint64_t，计算完成后，得到输出的shape为（32, 64），出参shape_out用于存放该shape信息，值为（0x0000000010000000 | 2, 32, 64）。代码示例如下：

  extern "C" __global__ __aicore__ void xxx_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR shape_out, GM_ADDR workspace, GM_ADDR tiling) {
  ...
      constexpr uint32_t SHAPEOUT_SIZE = 9;
      // 输出数据为4维([1, 64, 128, 128])，tensor类型为uint64_t
      GlobalTensor<uint64_t> shapeoutGlobal_uint64;
      shapeoutGlobal_uint64.SetGlobalBuffer((__gm__ uint64_t*)shape_out, SHAPEOUT_SIZE);
      shapeoutGlobal_uint64.SetValue(0, 0x0000000010000000 | 4);
      shapeoutGlobal_uint64.SetValue(1, 1);
      shapeoutGlobal_uint64.SetValue(2, 64);
      shapeoutGlobal_uint64.SetValue(3, 128);
      shapeoutGlobal_uint64.SetValue(4, 128);
  ...
  }

• 如下示例中，算子中有两个输出依赖计算得出，输出tensor的数据类型为uint64_t，计算完成后，得到输出的shape为（16, 32）和 （1, 16, 16, 32），出参shape_out用于存放该shape信息。示例如下：

  extern "C" __global__ __aicore__ void xxx_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR shape_out, GM_ADDR workspace, GM_ADDR tiling) {
      ...
      // 有两个输出需要刷新shape，一个维度为2维[16, 32]，一个维度为4维[1, 16, 16, 32]
      // tensor类型为uint64_t
      constexpr uint32_t SHAPEOUT_SIZE_2 = 18;
      GlobalTensor<uint64_t> shapeoutGlobal_uint64_2;
      shapeoutGlobal_uint64_2.SetGlobalBuffer((__gm__ uint64_t*)shape_out, SHAPEOUT_SIZE_2 );
      shapeoutGlobal_uint64_2.SetValue(0, 0x0000000010000000 | 2);
      shapeoutGlobal_uint64_2.SetValue(1, 16);
      shapeoutGlobal_uint64_2.SetValue(2, 32);
      // index[3]~index[8]数据为占位
      shapeoutGlobal_uint64_2.SetValue(9, 0x0000000010000000 | 4);
      shapeoutGlobal_uint64_2.SetValue(10, 1);
      shapeoutGlobal_uint64_2.SetValue(11, 16);
      shapeoutGlobal_uint64_2.SetValue(12, 16);
      shapeoutGlobal_uint64_2.SetValue(13, 32);
      ...
  }