Power
功能说明
实现按元素做幂运算功能,提供3类接口,处理逻辑分别为:
函数原型
- Power(dstTensor, srcTensor1, srcTensor2)
- 通过sharedTmpBuffer入参传入临时空间
- 源操作数Tensor全部/部分参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const LocalTensor<T>& src1Tensor, const LocalTensor<uint8_t>& sharedTmpBuffer, uint32_t calCount)
- 源操作数Tensor全部参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const LocalTensor<T>& src1Tensor, const LocalTensor<uint8_t>& sharedTmpBuffer)
- 源操作数Tensor全部/部分参与计算
- 接口框架申请临时空间
- 源操作数Tensor全部/部分参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const LocalTensor<T>& src1Tensor, uint32_t calCount)
- 源操作数Tensor全部参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const LocalTensor<T>& src1Tensor)
- 源操作数Tensor全部/部分参与计算
- 通过sharedTmpBuffer入参传入临时空间
- Power(dstTensor, srcTensor1, scalarValue)
- 通过sharedTmpBuffer入参传入临时空间
- 源操作数Tensor全部/部分参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const T& src1Scalar, const LocalTensor<uint8_t>& sharedTmpBuffer, uint32_t calCount)
- 源操作数Tensor全部参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const T& src1Scalar, const LocalTensor<uint8_t>& sharedTmpBuffer)
- 源操作数Tensor全部/部分参与计算
- 接口框架申请临时空间
- 源操作数Tensor全部/部分参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const T& src1Scalar, uint32_t calCount)
- 源操作数Tensor全部参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const LocalTensor<T>& src0Tensor, const T& src1Scalar)
- 源操作数Tensor全部/部分参与计算
- 通过sharedTmpBuffer入参传入临时空间
- Power(dstTensor, scalarValue, srcTensor2)
- 通过sharedTmpBuffer入参传入临时空间
- 源操作数Tensor全部/部分参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const T& src0Scalar, const LocalTensor<T>& src1Tensor, const LocalTensor<uint8_t>& sharedTmpBuffer, uint32_t calCount)
- 源操作数Tensor全部参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const T& src0Scalar, const LocalTensor<T>& src1Tensor, const LocalTensor<uint8_t>& sharedTmpBuffer)
- 源操作数Tensor全部/部分参与计算
- 接口框架申请临时空间
- 源操作数Tensor全部/部分参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const T& src0Scalar, const LocalTensor<T>& src1Tensor, uint32_t calCount)
- 源操作数Tensor全部参与计算
1 2
template <typename T, bool isReuseSource = false> __aicore__ inline void Power(const LocalTensor<T>& dstTensor, const T& src0Scalar, const LocalTensor<T>& src1Tensor)
- 源操作数Tensor全部/部分参与计算
- 通过sharedTmpBuffer入参传入临时空间
由于该接口的内部实现中涉及复杂的数学计算,需要额外的临时空间来存储计算过程中的中间变量。临时空间支持接口框架申请和开发者通过sharedTmpBuffer入参传入两种方式。
- 接口框架申请临时空间,开发者无需申请,但是需要预留临时空间的大小。
- 通过sharedTmpBuffer入参传入,使用该tensor作为临时空间进行处理,接口框架不再申请。该方式开发者可以自行管理sharedTmpBuffer内存空间,并在接口调用完成后,复用该部分内存,内存不会反复申请释放,灵活性较高,内存利用率也较高。
接口框架申请的方式,开发者需要预留临时空间;通过sharedTmpBuffer传入的情况,开发者需要为sharedTmpBuffer申请空间。临时空间大小BufferSize的获取方式如下:通过GetPowerMaxMinTmpSize中提供的GetPowerMaxMinTmpSize接口获取需要预留空间的范围大小。
参数说明
参数名 |
描述 |
---|---|
T |
操作数的数据类型。 |
isReuseSource |
是否允许修改源操作数。该参数预留,传入默认值false即可。 |
参数名 |
输入/输出 |
描述 |
---|---|---|
dstTensor |
输出 |
目的操作数。 类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/float/int32_t Atlas推理系列产品(Ascend 310P处理器)AI Core,支持的数据类型为:half/float/int32_t |
src0Tensor |
输入 |
源操作数。 类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。 源操作数的数据类型需要与目的操作数保持一致。 Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/float/int32_t Atlas推理系列产品(Ascend 310P处理器)AI Core,支持的数据类型为:half/float/int32_t |
src1Tensor |
输入 |
源操作数。 类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。 源操作数的数据类型需要与目的操作数保持一致。 Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/float/int32_t Atlas推理系列产品(Ascend 310P处理器)AI Core,支持的数据类型为:half/float/int32_t |
src0Scalar/src1Scalar |
输入 |
源操作数,类型为Scalar。源操作数的数据类型需要与目的操作数保持一致。 Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/float/int32_t Atlas推理系列产品(Ascend 310P处理器)AI Core,支持的数据类型为:half/float/int32_t |
sharedTmpBuffer |
输入 |
临时内存空间。 类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:uint8_t Atlas推理系列产品(Ascend 310P处理器)AI Core,支持的数据类型为:uint8_t 针对3个power接口,不同输入数据类型情况下,临时空间大小BufferSize的获取方式请参考GetPowerMaxMinTmpSize。 |
返回值
无
支持的型号
Atlas A2训练系列产品/Atlas 800I A2推理产品
Atlas推理系列产品(Ascend 310P处理器)AI Core
约束说明
- 不支持源操作数与目的操作数地址重叠。
- 对于Atlas推理系列产品(Ascend 310P处理器)AI Core,幂运算的指数必须小于231-1。
- 操作数地址偏移对齐要求请参见通用约束。
调用示例
本样例中只展示Compute流程中的部分代码。如果您需要运行样例代码,请将该代码段拷贝并替换样例模板中Compute函数的部分代码即可。
- Power(dstTensor, srcTensor1, srcTensor2)
Power(dstLocal, srcLocal1, srcLocal2)
结果示例如下:输入数据(srcLocal1): [1.4608411 4.344736 ... 0.46437776] 输入数据(srcLocal2): [-5.4534287 4.5122147 ... -0.9344089] 输出数据(dstLocal): [0.12657544 756.1846 ... 2.0477564]
- Power(dstTensor, srcTensor1, scalarValue)
Power(dstLocal, srcLocal1, scalarValue)
结果示例如下:输入数据(srcLocal1): [2.263972 2.902264 ... 0.40299487] 输入数据(scalarValue): 1.2260373 输出数据(dstLocal): [2.7232351 3.6926038 ... 0.32815763]
- Power(dstTensor, scalarValue, srcTensor2)
Power(dstLocal, scalarValue, srcLocal2)
结果示例如下:输入数据(scalarValue): 4.382112 输入数据(srcLocal2): [5.504859 2.0677629 ... 1.053188] 输出数据(dstLocal): [3407.0386 21.225077 ... 4.7403817]
样例模板
#include "kernel_operator.h" template <typename srcType> class KernelPower { public: __aicore__ inline KernelPower() {} __aicore__ inline void Init(GM_ADDR src1Gm, GM_ADDR src2Gm, GM_ADDR dstGm, uint32_t srcSize) { src1Global.SetGlobalBuffer(reinterpret_cast<__gm__ srcType *>(src1Gm), srcSize); src2Global.SetGlobalBuffer(reinterpret_cast<__gm__ srcType *>(src2Gm), srcSize); dstGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ srcType *>(dstGm), srcSize); pipe.InitBuffer(inQueueX1, 1, srcSize * sizeof(srcType)); pipe.InitBuffer(inQueueX2, 1, srcSize * sizeof(srcType)); pipe.InitBuffer(outQueue, 1, srcSize * sizeof(srcType)); bufferSize = srcSize; } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensor<srcType> srcLocal1 = inQueueX1.AllocTensor<srcType>(); AscendC::DataCopy(srcLocal1, src1Global, bufferSize); inQueueX1.EnQue(srcLocal1); AscendC::LocalTensor<srcType> srcLocal2 = inQueueX2.AllocTensor<srcType>(); AscendC::DataCopy(srcLocal2, src2Global, bufferSize); inQueueX2.EnQue(srcLocal2); } __aicore__ inline void Compute() { AscendC::LocalTensor<srcType> dstLocal = outQueue.AllocTensor<srcType>(); AscendC::LocalTensor<srcType> srcLocal1 = inQueueX1.DeQue<srcType>(); AscendC::LocalTensor<srcType> srcLocal2 = inQueueX2.DeQue<srcType>(); AscendC::LocalTensor<srcType> tmpLocal; srcType scalarValue1 = srcLocal1.GetValue(0); srcType scalarValue2 = srcLocal2.GetValue(0); AscendC::Power<srcType, false>(dstLocal, scalarValue1, srcLocal2); outQueue.EnQue<srcType>(dstLocal); inQueueX1.FreeTensor(srcLocal1); inQueueX2.FreeTensor(srcLocal2); } __aicore__ inline void CopyOut() { AscendC::LocalTensor<srcType> dstLocal = outQueue.DeQue<srcType>(); AscendC::DataCopy(dstGlobal, dstLocal, bufferSize); outQueue.FreeTensor(dstLocal); } private: AscendC::GlobalTensor<srcType> src1Global; AscendC::GlobalTensor<srcType> src2Global; AscendC::GlobalTensor<srcType> dstGlobal; AscendC::TPipe pipe; AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX1; AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX2; AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueue; uint32_t bufferSize = 0; }; template <typename dataType> __aicore__ void kernel_power_operator(GM_ADDR src1Gm, GM_ADDR src2Gm, GM_ADDR dstGm, uint32_t srcSize) { KernelPower<dataType> op; op.Init(src1Gm, src2Gm, dstGm, srcSize); op.Process(); } extern "C" __global__ __aicore__ void power_operator_custom(GM_ADDR src1Gm, GM_ADDR src2Gm, GM_ADDR dstGm, uint32_t srcSize) { kernel_power_operator<half>(src1Gm, src2Gm, dstGm, srcSize); }