动态shape场景

在算子实现章节，已经介绍了简单的固定shape矢量算子的kernel侧实现，算子的shape、数据类型都是固定的；在实际的算子开发场景中，这些信息是支持动态变化的，场景会更加灵活和复杂。下文重点进行动态shape与固定shape差异点的介绍。

最主要的区别是：动态Shape场景下，输入的Shape是未知的。一些与输入Shape相关的变量（比如每次搬运的块大小等），需要通过Tiling计算出来，然后传递到kernel侧，kernel侧使用该参数进行后续的计算。

• 算子实现章节中固定shape的算子样例中，TILE_NUM（每个核上总计算数据分块个数）、BLOCK_LENGTH（每个核上总计算数据大小）、TILE_LENGTH（每个分块大小）等是固定的数值。

  constexpr int32_t TOTAL_LENGTH = 8 * 2048;                            // total length of data
  constexpr int32_t USE_CORE_NUM = 8;                                   // num of core used
  constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;         // length computed of each core
  constexpr int32_t TILE_NUM = 8;                                       // split data into 8 tiles for each core
  constexpr int32_t BUFFER_NUM = 2;                                     // tensor num for each queue
  constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // each tile length is separated to 2 part, due to double buffer

• 如果需要将上述代码转换为动态shape，需要在核函数定义中增加Tiling参数，在host侧计算Tiling参数并传入，然后基于Tiling参数计算得到singleCoreSize（每个核上总计算数据大小）、tileNum（每个核上总计算数据分块个数）、tileLength（每个分块大小）等变量。

  __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
  {
      this->blockLength = totalLength / GetBlockNum();
      this->tileNum = tileNum;
      this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
      // ...
  
  }
  
  extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, AddCustomTilingData tiling)
  {
      KernelAdd op;
      op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
      op.Process();
  }