LeakyRelu

功能说明

按element做带泄露线性整流Leaky ReLU：

带泄露线性整流函数（Leaky Rectified Linear Unit, Leaky ReLU激活函数），是一种人工神经网络中常用的激活函数，其数学表达式为：

和ReLU的区别是：ReLU是将所有的负值都设为零，而Leaky Relu 是给所有负值赋予一个斜率。下图表示了Relu和vlrelu的区别：

对于Leaky ReLU函数，如果src的值小于零，dst的值等于src的值乘以scalar的值。如果src大于等于零，则dst的值等于src的值。

函数原型

• tensor前n个数据计算

  1 2	template <typename T, bool isSetMask = true> __aicore__ inline void LeakyRelu(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const T& scalarValue, const int32_t& calCount)

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

    1 2	template <typename T, bool isSetMask = true> __aicore__ inline void LeakyRelu(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const T& scalarValue, uint64_t mask[], const uint8_t repeatTimes, const UnaryRepeatParams& repeatParams)

  • mask连续模式

    1 2	template <typename T, bool isSetMask = true> __aicore__ inline void LeakyRelu(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const T& scalarValue, uint64_t mask, const uint8_t repeatTimes, const UnaryRepeatParams& repeatParams)

dstLocal和srcLocal使用TensorTrait类型时，其数据类型TensorTrait和scalarValue的数据类型（对应TensorTrait中的LiteType类型）不一致。因此新增模板类型U表示scalarValue的数据类型，并通过std::enable_if检查T中萃取出的LiteType和U是否完全一致，一致则接口通过编译，否则编译失败。接口原型定义如下：

• tensor前n个数据计算

  1 2	template <typename T, typename U, bool isSetMask = true, typename std::enable_if<IsSameType<PrimT<T>, U>::value, bool>::type = true> __aicore__ inline void LeakyRelu(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const U& scalarValue, const int32_t& calCount)

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

    1 2

    template <typename T, typename U, bool isSetMask = true, typename std::enable_if<IsSameType<PrimT<T>, U>::value, bool>::type = true> __aicore__ inline void LeakyRelu(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const U& scalarValue, uint64_t mask[], const uint8_t repeatTimes, const UnaryRepeatParams& repeatParams)

  • mask连续模式

    1 2

    template <typename T, typename U, bool isSetMask = true, typename std::enable_if<IsSameType<PrimT<T>, U>::value, bool>::type = true> __aicore__ inline void LeakyRelu(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const U& scalarValue, uint64_t mask, const uint8_t repeatTimes, const UnaryRepeatParams& repeatParams)

参数说明

返回值

无

支持的型号

Atlas推理系列产品AI Core

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

Atlas 200/500 A2推理产品

注意事项

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

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

调用示例

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

  #include "kernel_operator.h"
  class KernelBinaryScalar {
  public:
      __aicore__ inline KernelBinaryScalar() {}
      __aicore__ inline void Init(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
      {
          srcGlobal.SetGlobalBuffer((__gm__ half*)src);
          dstGlobal.SetGlobalBuffer((__gm__ half*)dstGm);
          pipe.InitBuffer(inQueueSrc, 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> srcLocal = inQueueSrc.AllocTensor<half>();
          AscendC::DataCopy(srcLocal, srcGlobal, 512);
          inQueueSrc.EnQue(srcLocal);
      }
      __aicore__ inline void Compute()
      {
          AscendC::LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
          AscendC::LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();
   
          uint64_t mask = 128;
          half scalar = 2;
          // repeatTimes = 4, 128 elements one repeat, 512 elements total
          // dstBlkStride, srcBlkStride = 1, no gap between blocks in one repeat
          // dstRepStride, srcRepStride =8, no gap between repeats
          AscendC::LeakyRelu(dstLocal, srcLocal, scalar, mask, 4, {1, 1, 8, 8});
          
          outQueueDst.EnQue<half>(dstLocal);
          inQueueSrc.FreeTensor(srcLocal);
      }
      __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> inQueueSrc;
      AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueDst;
      AscendC::GlobalTensor<half> srcGlobal, dstGlobal;
  };
  extern "C" __global__ __aicore__ void binary_scalar_simple_kernel(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
  {
      KernelBinaryScalar op;
      op.Init(src, dstGm);
      op.Process();
  }

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

  uint64_t mask[2] = { UINT64_MAX, UINT64_MAX };
  int16_t scalar = 2;
  // repeatTimes = 4, 单次迭代处理128个数，计算512个数需要迭代4次
  // dstBlkStride, srcBlkStride = 1, 每个迭代内src0参与计算的数据地址间隔为1个datablock，表示单次迭代内数据连续读取和写入
  // dstRepStride, srcRepStride = 8, 相邻迭代间的地址间隔为8个datablock，表示相邻迭代间数据连续读取和写入
  AscendC::LeakyRelu(dstLocal, srcLocal, scalar, mask, 4, {1, 1, 8, 8});

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

  half scalar = 2;
  AscendC::LeakyRelu(dstLocal, srcLocal, scalar, 512);

输入数据(src0Local): [1. 2. 3. ... 512.]
输入数据 scalar = 2.
输出数据(dstLocal): [1. 2. 3. ... 512.]