Div

产品支持情况

功能说明

按元素求商，公式表达如下：

函数原型

• 整个tensor参与计算

  dst = src0 / src1;
  

• tensor前n个数据计算

  template <typename T, const DivConfig& config = DEFAULT_DIV_CONFIG>
  __aicore__ inline void Div(const LocalTensor<T>& dst, const LocalTensor<T>& src0, const LocalTensor<T>& src1, const int32_t& count)
  

• tensor高维切分计算

  • mask逐bit模式

    template <typename T, bool isSetMask = true, const DivConfig& config = DEFAULT_DIV_CONFIG>
    __aicore__ inline void Div(const LocalTensor<T>& dst, const LocalTensor<T>& src0, const LocalTensor<T>& src1, uint64_t mask[], const uint8_t repeatTime, const BinaryRepeatParams& repeatParams)
    

  • mask连续模式

    template <typename T, bool isSetMask = true, const DivConfig& config = DEFAULT_DIV_CONFIG>
    __aicore__ inline void Div(const LocalTensor<T>& dst, const LocalTensor<T>& src0, const LocalTensor<T>& src1, uint64_t mask, const uint8_t repeatTime, const BinaryRepeatParams& repeatParams)
    

参数说明

表 1 模板参数说明

enum class DivAlgo {
    INTRINSIC = 0,
    DIFF_COMPENSATION,
    PRECISION_1ULP_FTZ_TRUE,
    PRECISION_0ULP_FTZ_TRUE,
    PRECISION_0ULP_FTZ_FALSE,
    PRECISION_1ULP_FTZ_FALSE
};
struct DivConfig {
    DivAlgo algo = DivAlgo::INTRINSIC;
};

constexpr DivConfig DEFAULT_DIV_CONFIG = { DivAlgo::INTRINSIC };

表 2 参数说明

返回值说明

无

约束说明

• 操作数地址对齐要求请参见通用地址对齐约束。

• 操作数地址重叠约束请参考通用地址重叠约束。

• 使用整个tensor参与计算接口符号重载时，运算量为目的LocalTensor的总长度。

• 注意除零错误。

• 针对Ascend 950PR/Ascend 950DT，uint64_t/int64_t/complex32/complex64数据类型仅支持tensor前n个数据计算接口和整个tensor参与计算的运算符重载。

调用示例

更多样例可参考LINK。

• 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::TPosition::VECIN, 1> inQueueSrc0, inQueueSrc1;
      AscendC::TQue<AscendC::TPosition::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模式

  uint64_t mask[2] = { UINT64_MAX, UINT64_MAX };
  // repeatTime = 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个数据计算样例

  AscendC::Div(dstLocal, src0Local, src1Local, 512);
  // Div 0ulp
  static constexpr DivConfig config = { DivAlgo::DIFF_COMPENSATION };
  Div<T, config>(dstLocalX, srcLocalX, srcLocalY, calCount);
  // Div Subnormal
  static constexpr DivConfig config = { DivAlgo::PRECISION_0ULP_FTZ_FALSE };
  Div<T, config>(dstLocalX, srcLocalX, srcLocalY, calCount);
  

• 整个tensor参与计算样例

  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]