LayerNormGradBeta【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit产品支持情况产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√功能说明LayerNormGradBeta接口用于获取反向beta/gamma的数值和LayerNormGrad共同输出pdx, gamma和beta算法公式为:函数原型由于该接口的内部实现中涉及复杂的计算需要额外的临时空间来存储计算过程中的中间变量。临时空间大小BufferSize的获取方法通过LayerNormGradBeta Tiling中提供的GetLayerNormGradBetaMaxMinTmpSize接口获取所需最大和最小临时空间大小最小空间可以保证功能正确最大空间用于提升性能。临时空间支持接口框架申请和开发者通过sharedTmpBuffer入参传入两种方式因此LayerNormGradBeta接口的函数原型有两种通过sharedTmpBuffer入参传入临时空间template typename T, bool isReuseSource false __aicore__ inline void LayerNormGradBeta(const LocalTensorT outputPdGamma, const LocalTensorT outputPdBeta, const LocalTensorT resForGamma, const LocalTensorT inputDy, const LocalTensoruint8_t sharedTmpBuffer, const LayerNormGradBetaTiling tiling)该方式下开发者需自行申请并管理临时内存空间并在接口调用完成后复用该部分内存内存不会反复申请释放灵活性较高内存利用率也较高。接口框架申请临时空间template typename T, bool isReuseSource false __aicore__ inline void LayerNormGradBeta(const LocalTensorT outputPdGamma, const LocalTensorT outputPdBeta, const LocalTensorT resForGamma, const LocalTensorT inputDy, LayerNormGradBetaTiling tiling)该方式下开发者无需申请但是需要预留临时空间的大小。参数说明表 1模板参数说明参数名描述T操作数的数据类型。Ascend 950PR/Ascend 950DT支持的数据类型为half、float。Atlas A3 训练系列产品 / Atlas A3 推理系列产品支持的数据类型为half、float。Atlas A2 训练系列产品 / Atlas A2 推理系列产品支持的数据类型为half、float。isReuseSource是否允许修改源操作数默认值为false。如果开发者允许源操作数被改写可以使能该参数使能后能够节省部分内存空间。设置为true则本接口内部计算时复用inputDy的内存空间节省内存空间设置为false则本接口内部计算时不复用inputDy的内存空间。对于float数据类型输入支持开启该参数half数据类型输入不支持开启该参数。isReuseSource的使用样例请参考更多样例。表 2接口参数说明参数名称输入/输出含义outputPdGamma输出目的操作数shape为[H]LocalTensor数据结构的定义请参考LocalTensor。尾轴长度需要32B对齐类型为LocalTensor支持的TPosition为VECIN/VECCALC/VECOUT。outputPdBeta输出目的操作数shape为[H]LocalTensor数据结构的定义请参考LocalTensor。尾轴长度需要32B对齐类型为LocalTensor支持的TPosition为VECIN/VECCALC/VECOUT。resForGamma输入源操作数shape为[B, S, H]LocalTensor数据结构的定义请参考LocalTensor。resForGamma的数据类型需要与目的操作数保持一致尾轴长度需要32B对齐。需提前调用LayerNormGrad接口获取resForGamma参数值。类型为LocalTensor支持的TPosition为VECIN/VECCALC/VECOUT。inputDy输入源操作数shape为[B, S, H]LocalTensor数据结构的定义请参考LocalTensor。inputDy的数据类型需要与目的操作数保持一致尾轴长度需要32B对齐。类型为LocalTensor支持的TPosition为VECIN/VECCALC/VECOUT。sharedTmpBuffer输入共享缓冲区用于存放API内部计算产生的临时数据。该方式开发者可以自行管理sharedTmpBuffer内存空间并在接口调用完成后复用该部分内存内存不会反复申请释放灵活性较高内存利用率也较高。共享缓冲区大小的获取方式请参考LayerNormGradBeta Tiling。类型为LocalTensor支持的TPosition为VECIN/VECCALC/VECOUT。tiling输入LayerNormGradBeta计算所需Tiling信息Tiling信息的获取请参考LayerNormGradBeta Tiling。返回值说明无约束说明操作数地址对齐要求请参见通用地址对齐约束。源操作数和目的操作数的Tensor空间可以复用。仅支持输入shape为ND格式。输入数据不满足对齐要求时开发者需要进行补齐补齐的数据应设置为0防止出现异常值从而影响网络计算。不支持对尾轴H轴的切分。调用示例#include kernel_operator.h template typename T, bool isReuseSource false class KernelLayernormGradBeta { public: __aicore__ inline KernelLayernormGradBeta() {} __aicore__ inline void Init(__gm__ uint8_t *resForGammaGm, __gm__ uint8_t *inputDyGm, __gm__ uint8_t *outputPdGammaGm, __gm__ uint8_t *outputPdBetaGm, const LayerNormGradBetaTiling tiling) { this-bLength tiling.bLength; this-sLength tiling.sLength; this-hLength tiling.hLength; this-tiling tiling; bshLength bLength * sLength * hLength; bsLength bLength * sLength; resForGammaGlobal.SetGlobalBuffer(reinterpret_cast__gm__ T *(resForGammaGm), bshLength); inputDyGlobal.SetGlobalBuffer(reinterpret_cast__gm__ T *(inputDyGm), bshLength); outputPdGammaGlobal.SetGlobalBuffer(reinterpret_cast__gm__ T *(outputPdGammaGm), hLength); outputPdBetaGlobal.SetGlobalBuffer(reinterpret_cast__gm__ T *(outputPdBetaGm), hLength); pipe.InitBuffer(inQueueResForGamma, 1, sizeof(T) * bshLength); pipe.InitBuffer(inQueueDy, 1, sizeof(T) * bshLength); pipe.InitBuffer(outQueuePdGamma, 1, sizeof(T) * hLength); pipe.InitBuffer(outQueuePdBeta, 1, sizeof(T) * hLength); } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensorT resForGammaLocal inQueueResForGamma.AllocTensorT(); AscendC::LocalTensorT inputDyLocal inQueueDy.AllocTensorT(); AscendC::DataCopy(resForGammaLocal, resForGammaGlobal, bshLength); AscendC::DataCopy(inputDyLocal, inputDyGlobal, bshLength); inQueueResForGamma.EnQue(resForGammaLocal); inQueueDy.EnQue(inputDyLocal); } __aicore__ inline void Compute() { AscendC::LocalTensorT resForGammaLocal inQueueResForGamma.DeQueT(); AscendC::LocalTensorT inputDyLocal inQueueDy.DeQueT(); AscendC::LocalTensorT outputPdGammaLocal outQueuePdGamma.AllocTensorT(); AscendC::LocalTensorT outputPdBetaLocal outQueuePdBeta.AllocTensorT(); AscendC::LayerNormGradBetaT, isReuseSource( outputPdGammaLocal, outputPdBetaLocal, resForGammaLocal, inputDyLocal, tiling); outQueuePdGamma.EnQueT(outputPdGammaLocal); outQueuePdBeta.EnQueT(outputPdBetaLocal); inQueueResForGamma.FreeTensor(resForGammaLocal); inQueueDy.FreeTensor(inputDyLocal); } __aicore__ inline void CopyOut() { AscendC::LocalTensorT outputPdGammaLocal outQueuePdGamma.DeQueT(); AscendC::LocalTensorT outputPdBetaLocal outQueuePdBeta.DeQueT(); AscendC::DataCopy(outputPdGammaGlobal, outputPdGammaLocal, hLength); AscendC::DataCopy(outputPdBetaGlobal, outputPdBetaLocal, hLength); outQueuePdGamma.FreeTensor(outputPdGammaLocal); outQueuePdBeta.FreeTensor(outputPdBetaLocal); } private: AscendC::TPipe pipe; AscendC::TQueAscendC::TPosition::VECIN, 1 inQueueResForGamma, inQueueDy; AscendC::TQueAscendC::TPosition::VECOUT, 1 outQueuePdGamma, outQueuePdBeta; AscendC::GlobalTensorT resForGammaGlobal; AscendC::GlobalTensorT inputDyGlobal; AscendC::GlobalTensorT outputPdGammaGlobal; AscendC::GlobalTensorT outputPdBetaGlobal; uint32_t bLength; uint32_t sLength; uint32_t hLength; uint32_t bshLength; uint32_t bsLength; LayerNormGradBetaTiling tiling; }; extern C __global__ __aicore__ void kernel_layernorm_grad_beta_operator( GM_ADDR outputPdGammaGm, GM_ADDR outputPdBetaGm, GM_ADDR resForGammaGm, GM_ADDR inputDyGm, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); KernelLayernormGradBetahalf, false op; op.Init(resForGammaGm, inputDyGm, outputPdGammaGm, outputPdBetaGm, tilingData.layerNormGradBetaTiling); op.Process(); }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考