DataCopy

功能说明

数据搬运接口，支持的数据传输通路如下：

GM->A1

GM->B1

CO1->CO2

CO2->GM

GM->VECIN

VECIN->VECOUT

VECOUT->GM

定义原型

普通数据搬运0级接口，适用于连续和不连续数据搬运：

• 通路：GM->VECIN, GM->A1, GM->B1

  template <typename T>
  void DataCopy(const LocalTensor<T>& dstLocal, const GlobalTensor<T>& srcGlobal, const DataCopyParams& intriParams);

• 通路：VECIN->VECOUT

  template <typename T>
  void DataCopy(const LocalTensor<T>& dstLocal, const LocalTensor <T>& srcLocal, const DataCopyParams& intriParams);

• 通路：VECOUT->GM, CO2->GM

  template <typename T>
  void DataCopy(const GlobalTensor <T>& dstGlobal, const LocalTensor <T>& srcLocal, const DataCopyParams& intriParams);

普通数据搬运2级接口，适用于连续数据搬运：

• 通路：GM->VECIN, GM->A1, GM->B1

  template <typename T>
  void DataCopy(const LocalTensor<T>& dstLocal, const GlobalTensor<T>& srcGlobal, const uint32_t calCount);

• 通路：VECIN->VECOUT

  template <typename T>
  void DataCopy(const LocalTensor<T>& dstLocal, const LocalTensor <T>& srcLocal, const uint32_t calCount);

• 通路：VECOUT->GM, CO2->GM

  template <typename T>
  void DataCopy(const GlobalTensor <T>& dstGlobal, const LocalTensor <T>& srcLocal, const uint32_t calCount);

增强数据搬运0级接口，相比于普通数据搬运0级接口，搬运时增加了随路计算：

• 通路：GM->VECIN, GM->A1, GM->B1

  template <typename T>
  __aicore__ inline void DataCopy(const LocalTensor<T>& dstLocal, const GlobalTensor<T>& srcGlobal, const DataCopyParams& intriParams, const DataCopyEnhancedParams& enhancedParams)

• 通路：VECIN->VECOUT, CO1->CO2

  __aicore__ inline void DataCopy(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const DataCopyParams& intriParams, const DataCopyEnhancedParams& enhancedParams)

• 通路：VECOUT->GM, CO2->GM

  template <typename T>
  __aicore__ inline void DataCopy(const GlobalTensor<T>& dstGlobal, const LocalTensor<T>& srcLocal, const DataCopyParams& intriParams, const DataCopyEnhancedParams& enhancedParams)

切片数据搬运，主要适用于非连续vector数据搬运：

• 通路：GM->VECIN

  template <typename T>
  void DataCopy(const LocalTensor<T> &dstLocal, const GlobalTensor<T> &srcGlobal, const SliceInfo dstSliceInfo[], const SliceInfo srcSliceInfo[], const uint32_t dimValue = 1)

• 通路：VECOUT->GM, CO2->GM

  template <typename T>
  void DataCopy(const GlobalTensor<T> &dstGlobal, const LocalTensor<T> &srcLocal, const SliceInfo dstSliceInfo[], const SliceInfo srcSliceInfo[], const uint32_t dimValue = 1)

随路格式转换数据搬运，适用于cube数据搬运，在搬运时进行格式转换：

• 通路：GM->A1, GM->B1(ND2NZ)

  template <typename T>
  void DataCopy(const LocalTensor<T>& dstLocal, const GlobalTensor<T>& srcGlobal, const Nd2NzParams& intriParams);

• 通路：VECOUT->GM(NZ2ND)

  template <typename T>
  void DataCopy(const GlobalTensor <T>& dstGlobal, const LocalTensor <T>& srcLocal, const Nz2NdParamsFull &intriParams);

参数说明（普通数据搬运）

下面的样例呈现了DataCopyParams结构体参数的使用方法，样例中完成了2个连续传输数据块的搬运，每个数据块含有8个block，源操作数相邻数据块之间无间隔，目的操作数相邻数据块尾与头之间间隔1个block。

参数说明（增强数据搬运）

参数说明（切片数据搬运）

参数说明（随路格式转换）

ND2NZ转换示意图如下，为了方便展示和理解，这里我们假设分形矩阵的大小是2x2，ND矩阵的大小是4x4，请注意这与实际的情况并不相符。样例中参数设置值和解释说明如下：

• ndNum = 1，表示传输nd矩阵的数目为1。nValue=4，矩阵的行数，也就是矩阵的高度为4。dValue=4，矩阵的列数，也就是矩阵的宽度为4。
• srcNdMatrixStride = 0，当前只有一个nd矩阵，不存在相邻nd矩阵起始地址间的偏移，设置为0。同理，dstNzMatrixStride =0。
• srcDValue= dValue, 源操作数同一nd矩阵的相邻行起始地址间的偏移，也就是图中addr1和addr2之间的偏移。
• dstNzC0Stride = nValue *C0_SIZE/C0_SIZE。ND转换到NZ格式后，源操作数中的一行会转换为目的操作数的多行，也就是图中的line_src会转换为line_dst1和line_dst2。多行数据起始地址之间的偏移就是addr3和addr4之间的偏移，偏移为nValue *C0_SIZE，因为该参数的单位为C0_SIZE，所以，dstNzC0Stride = nValue *C0_SIZE/C0_SIZE。
• dstNzNStride=C0_SIZE/C0_SIZE，即图中addr3和addr5之间的偏移C0_SIZE，因为该参数的单位为C0_SIZE，所以，dstNzNStride=C0_SIZE/C0_SIZE。

图1 Nd2Nz转换示意图

支持的型号

Atlas 训练系列产品

Atlas推理系列产品AI Core

Atlas A2训练系列产品

注意事项

• 增强搬运主要分为CO1->CO2及其他情况两种。
• DataCopy的搬运量要求为32byte的倍数，因此使用2级接口时，calCount * sizeof(T)需要32byte对齐，若不对齐，搬运量将对32byte做向下取整。
• 配置relu与deqScale的某bit配置冲突时，以relu为准。
• 切片数据搬运中的横向burstLen大小设置，需要用户自己通过计算：横向切片元素个数* sizeof(T)/32byte。横向切片元素个数* sizeof(T)的大小必须32byte的倍数。
• 切片数据搬运中的SliceInfo结构体数组大小和dimValue需要保持一致，并且不超过8。
• 切片数据搬运中的srcSliceInfo数组大小的和dstSliceInfo的大小需要保持一致，两者的结构体中的burstLen需要相等（srcSliceInfo[i].burstLen = dstSliceInfo[i].burstLen）。
• 切片数据搬运对参数有一定要求，建议使用者参考调用示例，并在CPU上仿真结果无误后，再到NPU侧执行。
• 如果 CO1->CO2有随路精度转换，通路为 UB 的操作数的 blockCount 单位需要减半。
• 随路格式转换nd2nz只支持GM->A1,GM->B1； nz2nd只支持VECOUT->GM。
• 硬件在执行数据搬运时会以block作为基本单位，而 1 block = 32 Byte，故使用者可以尝试通过每次指令处理32Byte整数倍大小的数据来提高指令的执行效率。

返回值

无

调用示例

• 普通数据搬运2级接口

  #include "kernel_operator.h"
  namespace AscendC {
  class KernelDataCopy {
  public:
      __aicore__ inline KernelDataCopy() {}
      __aicore__ inline void Init(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm)
      {
          src0Global.SetGlobalBuffer((__gm__ half*)src0Gm);
          src1Global.SetGlobalBuffer((__gm__ half*)src1Gm);
          dstGlobal.SetGlobalBuffer((__gm__ half*)dstGm);
          pipe.InitBuffer(inQueueSrc0, 1, 512 * sizeof(half));
          pipe.InitBuffer(inQueueSrc1, 1, 512 * sizeof(half));
          pipe.InitBuffer(outQueueDst, 1, 512 * sizeof(half));
      }
      __aicore__ inline void Process()
      {
          CopyIn();
          Compute();
          CopyOut();
      }
  private:
      __aicore__ inline void CopyIn()
      {
          LocalTensor<half> src0Local = inQueueSrc0.AllocTensor<half>();
          LocalTensor<half> src1Local = inQueueSrc1.AllocTensor<half>();
  
          DataCopy(src0Local, src0Global, 512);
          DataCopy(src1Local, src1Global, 512);
  
          inQueueSrc0.EnQue(src0Local);
          inQueueSrc1.EnQue(src1Local);
      }
      __aicore__ inline void Compute()
      {
          LocalTensor<half> src0Local = inQueueSrc0.DeQue<half>();
          LocalTensor<half> src1Local = inQueueSrc1.DeQue<half>();
          LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();
          Add(dstLocal, src0Local, src1Local, 512);
          outQueueDst.EnQue<half>(dstLocal);
          inQueueSrc0.FreeTensor(src0Local);
          inQueueSrc1.FreeTensor(src1Local);
      }
      __aicore__ inline void CopyOut()
      {
          LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
  
          DataCopy(dstGlobal, dstLocal, 512);
  
          outQueueDst.FreeTensor(dstLocal);
      }
  private:
      TPipe pipe;
      TQue<QuePosition::VECIN, 1> inQueueSrc0, inQueueSrc1;
      TQue<QuePosition::VECOUT, 1> outQueueDst;
      GlobalTensor<half> src0Global, src1Global, dstGlobal;
  };
  } // namespace AscendC
  extern "C" __global__ __aicore__ void data_copy_kernel(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm, __gm__ uint8_t* dstGm)
  {
      AscendC::KernelDataCopy op;
      op.Init(src0Gm, src1Gm, dstGm);
      op.Process();
  }

  结果示例：

  输入数据(src0Global): [1 2 3 ... 512]
  输入数据(src1Global): [1 2 3 ... 512]
  输出数据(dstGlobal):[2 4 6 ... 1024]

• 随路格式转换数据搬运，nz2nd

  #include "kernel_operator.h"
  namespace AscendC {
  class KernelDataCopyUb2GmNz2Nd {
  public:
      __aicore__ inline KernelDataCopyUb2GmNz2Nd()
      {}
      __aicore__ inline void Init(__gm__ uint8_t* dstGm, __gm__ uint8_t* srcGm)
      {
          Nz2NdParamsFull intriParamsIn{1, 32, 32, 1, 32, 32, 1};
          intriParams = intriParamsIn;
          srcGlobal.SetGlobalBuffer((__gm__ half *)srcGm);
          dstGlobal.SetGlobalBuffer((__gm__ half *)dstGm);
          pipe.InitBuffer(inQueueSrcVecIn, 1, intriParams.nValue * intriParams.dValue * sizeof(half));
          pipe.InitBuffer(inQueueSrcVecOut, 1, intriParams.nValue * intriParams.dValue * sizeof(half));
      }
      __aicore__ inline void Process()
      {
          CopyIn();
          Compute();
          CopyOut();
      }
  
  private:
      __aicore__ inline void CopyIn()
      {
          LocalTensor<half> srcLocal = inQueueSrcVecIn.AllocTensor<half>();
          DataCopy(srcLocal, srcGlobal, intriParams.nValue * intriParams.dValue);
          inQueueSrcVecIn.EnQue(srcLocal);
      }
      __aicore__ inline void Compute()
      {
          LocalTensor<half> dstLocal = inQueueSrcVecIn.DeQue<half>();
          LocalTensor<half> srcOutLocal = inQueueSrcVecOut.AllocTensor<half>();
          DataCopy(srcOutLocal, dstLocal, intriParams.nValue * intriParams.dValue);
          inQueueSrcVecOut.EnQue(srcOutLocal);
          inQueueSrcVecIn.FreeTensor(dstLocal);
      }
      __aicore__ inline void CopyOut()
      {
          LocalTensor<half> srcOutLocalDe = inQueueSrcVecOut.DeQue<half>();
          DataCopy(dstGlobal, srcOutLocalDe, intriParams);
          inQueueSrcVecOut.FreeTensor(srcOutLocalDe);
      }
  
  private:
      TPipe pipe;
      TQue<QuePosition::VECIN, 1> inQueueSrcVecIn;
      TQue<QuePosition::VECOUT, 1> inQueueSrcVecOut;
      GlobalTensor<half> srcGlobal;
      GlobalTensor<half> dstGlobal;
      Nz2NdParamsFull intriParams;
  };
  }  // namespace AscendC
  
  extern "C" __global__ __aicore__ void kernel_data_copy_nz2nd_ub2out(__gm__ uint8_t* src_gm, __gm__ uint8_t* dst_gm)
  {
      AscendC::KernelDataCopyUb2GmNz2Nd op;
      op.Init(dst_gm, src_gm);
      op.Process();
  }

  结果示例：

  输入数据(srcGlobal): [1 2 3 ... 1024]
  输出数据(dstGlobal):[1 2 ... 15 16 513 514 ... 527 528 17 18 ... 31 32 529 530 ... 543 544 ...497 498 ...  511 512  1009 1010... 1023 1024]

• 切片数据搬运，非连续转为连续

  #include "kernel_operator.h"
  namespace AscendC {
  class KernelDataCopySliceGM2UB {
  public:
      __aicore__ inline KernelDataCopySliceGM2UB()
      {}
      __aicore__ inline void Init(__gm__ uint8_t* dstGm, __gm__ uint8_t* srcGm)
      {
          SliceInfo srcSliceInfoIn = {{16, 70, 7, 3}, {0, 2, 1, 1}};
          SliceInfo dstSliceInfoIn = {{0, 47, 0, 3}, {0, 1, 0, 1}};// UB空间相对紧张，建议设置stride为0。
          uint32_t dimValueIn = 2;
          uint32_t dstDataSize = 96;
          uint32_t srcDataSize = 261;
          dimValue = dimValueIn;
  
          for (uint32_t i = 0; i < dimValueIn; i++) {
              srcSliceInfo[i].startIndex = srcSliceInfoIn[i].startIndex;
              srcSliceInfo[i].endIndex = srcSliceInfoIn[i].endIndex;
              srcSliceInfo[i].stride = srcSliceInfoIn[i].stride;
              srcSliceInfo[i].burstLen = srcSliceInfoIn[i].burstLen;
  
              dstSliceInfo[i].startIndex = dstSliceInfoIn[i].startIndex;
              dstSliceInfo[i].endIndex = dstSliceInfoIn[i].endIndex;
              dstSliceInfo[i].stride = dstSliceInfoIn[i].stride;
              dstSliceInfo[i].burstLen = dstSliceInfoIn[i].burstLen;
          }
          uint32_t srsGmShape = {87, 3};  // W:87, H:3 需要注意shape的顺序
          uint32_t dstGmShape = {48, 2};  // W:48, H:2 需要注意shape的顺序
          ShapeInfo srcShapeInfo{2, srsGmShape, 2, srsGmShape, DataFormat::ND};
          ShapeInfo dstShapeInfo{2, dstGmShape, 2, dstGmShape, DataFormat::ND};
  
          srcGlobal.SetGlobalBuffer((__gm__ float *)srcGm);
          dstGlobal.SetGlobalBuffer((__gm__ float *)dstGm);
          srcGlobal.SetShapeInfo(srcShapeInfo);
          dstGlobal.SetShapeInfo(dstShapeInfo);
  
          pipe.InitBuffer(inQueueSrcVecIn, 1, dstDataSize * sizeof(float));
      }
      __aicore__ inline void Process()
      {
          CopyIn();
          CopyOut();
      }
  
  private:
      __aicore__ inline void CopyIn()
      {
          LocalTensor<float> srcLocal = inQueueSrcVecIn.AllocTensor<float>();
          ShapeInfo localShapeInfo = {2, {48, 2}, 2, {48, 2}, DataFormat::ND}; // 设置ub上的localtensor的shape
          srcLocal.SetShapeInfo(localShapeInfo);
          DataCopy(srcLocal, srcGlobal, srcSliceInfo, dstSliceInfo, dimValue);
          inQueueSrcVecIn.EnQue(srcLocal);
      }
      __aicore__ inline void CopyOut()
      {
          LocalTensor<float> srcOutLocal = inQueueSrcVecIn.DeQue<float>();
          DataCopy(dstGlobal, srcOutLocal, dstSliceInfo, dstSliceInfo, dimValue);
          inQueueSrcVecIn.FreeTensor(srcOutLocal);
      }
  
  private:
      TPipe pipe;
      TQue<QuePosition::VECIN, 1> inQueueSrcVecIn;
      GlobalTensor<float> srcGlobal;
      GlobalTensor<float> dstGlobal;
      SliceInfo dstSliceInfo[K_MAX_DIM];
      SliceInfo srcSliceInfo[K_MAX_DIM]; // K_MAX_DIM = 8
      uint32_t dimValue;
  };
  }  // namespace AscendC
  
  extern "C" __global__ __aicore__ void kernel_data_copy_slice_out2ub(__gm__ uint8_t* src_gm, __gm__ uint8_t* dst_gm)
  {
      AscendC::KernelDataCopySliceGM2UB op;
      op.Init(dst_gm, src_gm);
      op.Process();
  }

  结果示例：

  输入数据(srcGlobal)：

  00000000	00000000	11111111	11111111	11111111	0000000（7个0）	11111111	11111111	11111111	00000000	00000000
  00000000	00000000	00000000	00000000	00000000	0000000（7个0）	00000000	00000000	00000000	00000000	00000000
  00000000	00000000	11111111	11111111	11111111	0000000（7个0）	11111111	11111111	11111111	00000000	00000000

  输出数据(dstGlobal)：

  11111111	11111111	11111111	11111111	11111111	11111111
  11111111	11111111	11111111	11111111	11111111	11111111