Div

功能说明

按element求商，公式表达如下，其中PAR表示矢量计算单元一个迭代能够处理的元素个数：

函数原型

• 整个tensor参与计算

  1	dstLocal = src0Local / src1Local;

• tensor前n个数据计算

  1 2	template <typename T> __aicore__ inline void Div(const LocalTensor<T>& dstLocal, const LocalTensor<T>& src0Local, const LocalTensor<T>& src1Local, const int32_t& calCount)

• tensor高维切分计算
  • mask逐bit模式

    1 2	template <typename T, bool isSetMask = true> __aicore__ inline void Div(const LocalTensor<T>& dstLocal, const LocalTensor<T>& src0Local, const LocalTensor<T>& src1Local, uint64_t mask[2], const uint8_t repeatTimes, const BinaryRepeatParams& repeatParams)

  • mask连续模式

    1 2	template <typename T, bool isSetMask = true> __aicore__ inline void Div(const LocalTensor<T>& dstLocal, const LocalTensor<T>& src0Local, const LocalTensor<T>& src1Local, uint64_t mask, const uint8_t repeatTimes, const BinaryRepeatParams& repeatParams)

参数说明

返回值

无

支持的型号

Atlas 训练系列产品

Atlas推理系列产品AI Core

Atlas A2训练系列产品/Atlas 800I A2推理产品

Atlas 200/500 A2推理产品

注意事项

• 使用tensor高维切分计算接口时，节省地址空间，开发者可以定义一个Tensor，供源操作数与目的操作数同时使用（即地址重叠），相关约束如下：
  • 单次迭代内，要求源操作数和目的操作数之间100%重叠，不支持部分重叠。
  • 多次迭代间，第N次目的操作数是第N+1次源操作数的情况下，是不支持地址重叠的，因为第N+1次依赖第N次的结果。

• 使用整个tensor参与计算接口符号重载时，运算量为目的LocalTensor的总长度。
• 注意除零错误。
• 操作数地址偏移对齐要求请参见通用约束。

调用示例

• tensor高维切分计算样例-mask连续模式

  #include "kernel_operator.h"
   
  class KernelDiv {
  public:
      __aicore__ inline KernelDiv() {}
      __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()
      {
          AscendC::LocalTensor<half> src0Local = inQueueSrc0.AllocTensor<half>();
          AscendC::LocalTensor<half> src1Local = inQueueSrc1.AllocTensor<half>();
          AscendC::DataCopy(src0Local, src0Global, 512);
          AscendC::DataCopy(src1Local, src1Global, 512);
          inQueueSrc0.EnQue(src0Local);
          inQueueSrc1.EnQue(src1Local);
      }
      __aicore__ inline void Compute()
      {
          AscendC::LocalTensor<half> src0Local = inQueueSrc0.DeQue<half>();
          AscendC::LocalTensor<half> src1Local = inQueueSrc1.DeQue<half>();
          AscendC::LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();
          
          uint64_t mask = 128;
          AscendC::Div(dstLocal, src0Local, src1Local, mask, 4, { 1, 1, 1, 8, 8, 8 });
   
          outQueueDst.EnQue<half>(dstLocal);
          inQueueSrc0.FreeTensor(src0Local);
          inQueueSrc1.FreeTensor(src1Local);
      }
      __aicore__ inline void CopyOut()
      {
          AscendC::LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
          AscendC::DataCopy(dstGlobal, dstLocal, 512);
          outQueueDst.FreeTensor(dstLocal);
      }
  private:
      AscendC::TPipe pipe;
      AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueSrc0, inQueueSrc1;
      AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueDst;
      AscendC::GlobalTensor<half> src0Global, src1Global, dstGlobal;
  };
   
  extern "C" __global__ __aicore__ void div_simple_kernel(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm,
      __gm__ uint8_t* dstGm)
  {
      KernelDiv op;
      op.Init(src0Gm, src1Gm, dstGm);
      op.Process();
  }

• tensor高维切分计算样例-mask逐bit模式（本样例中只展示Compute流程中的部分代码。如果您需要运行样例代码，请将该代码段拷贝并替换上方样例的Compute函数中粗体部分即可。）

  uint64_t mask[2] = { UINT64_MAX, UINT64_MAX };
  // repeatTimes = 4, 一次迭代计算128个数, 共计算512个数
  // dstBlkStride, src0BlkStride, src1BlkStride = 1, 单次迭代内数据连续读取和写入
  // dstRepStride, src0RepStride, src1RepStride = 8, 相邻迭代间数据连续读取和写入
  AscendC::Div(dstLocal, src0Local, src1Local, mask, 4, { 1, 1, 1, 8, 8, 8 });

• tensor前n个数据计算样例（本样例中只展示Compute流程中的部分代码。如果您需要运行样例代码，请将该代码段拷贝并替换上方样例的Compute函数中粗体部分即可。）

  AscendC::Div(dstLocal, src0Local, src1Local, 512);

• 整个tensor参与计算样例（本样例中只展示Compute流程中的部分代码。如果您需要运行样例代码，请将该代码段拷贝并替换上方样例的Compute函数中粗体部分即可。）

  dstLocal = src0Local / src1Local;

输入数据(src0Local): [1.0 2.0 3.0 ... 512.0]
输入数据(src1Local): [2.0 2.0 2.0 ... 2.0]
输出数据(dstLocal): [0.5 1.0 1.5 ... 256.0]