Div

功能说明

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

函数原型

参数说明

返回值

无

支持的型号

Atlas 训练系列产品

Atlas推理系列产品AI Core

Atlas A2训练系列产品

注意事项

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

• 使用三级接口符号重载时，运算量为目的LocalTensor的总长度。
• 三级接口符号重载会进行连续矢量运算，无法调整mask, repeatTimes, stride等参数。
• 注意除零错误。
• 操作数地址偏移对齐要求请见通用约束。

调用示例

• 0级接口样例-mask连续模式

  #include "kernel_operator.h"
  namespace AscendC {
  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()
      {
          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>();
          
          uint64_t mask = 128;
         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()
      {
          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 div_simple_kernel(__gm__ uint8_t* src0Gm, __gm__ uint8_t* src1Gm,
      __gm__ uint8_t* dstGm)
  {
      AscendC::KernelDiv op;
      op.Init(src0Gm, src1Gm, dstGm);
      op.Process();
  }

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

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

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

  Div(dstLocal, src0Local, src1Local, 512);

• 3级接口样例（本样例中只展示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]