Axpy
产品支持情况
功能说明
源操作数src中每个元素与标量求积后和目的操作数dst中的对应元素相加，计算公式如下：
函数原型
-
tensor前n个数据计算
template <typename T, typename U> __aicore__ inline void Axpy(const LocalTensor<T>& dst, const LocalTensor<U>& src, const U& scalarValue, const int32_t& count) -
tensor高维切分计算
-
mask逐bit模式
template <typename T, typename U, bool isSetMask = true> __aicore__ inline void Axpy(const LocalTensor<T>& dst, const LocalTensor<U>& src, const U& scalarValue, uint64_t mask[], const uint8_t repeatTime, const UnaryRepeatParams& repeatParams) -
mask连续模式
template <typename T, typename U, bool isSetMask = true> __aicore__ inline void Axpy(const LocalTensor<T>& dst, const LocalTensor<U>& src, const U& scalarValue, uint64_t mask, const uint8_t repeatTime, const UnaryRepeatParams& repeatParams)
-
参数说明
表 1 模板参数说明
目的操作数数据类型。目的操作数和源操作数的数据类型约束请参考表3。 Atlas A2 训练系列产品/Atlas A2 推理系列产品，支持的数据类型为：half/float Atlas A3 训练系列产品/Atlas A3 推理系列产品，支持的数据类型为：half/float |
|
|
Atlas A2 训练系列产品/Atlas A2 推理系列产品，支持的数据类型为：half/float Atlas A3 训练系列产品/Atlas A3 推理系列产品，支持的数据类型为：half/float |
|
|
表 2 参数说明
|
类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 |
||
|
类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。 |
||
|
||
|
矢量计算单元，每次读取连续的256Bytes数据进行计算，为完成对输入数据的处理，必须通过多次迭代（repeat）才能完成所有数据的读取与计算。repeatTime表示迭代的次数。 |
||
控制操作数地址步长的参数。UnaryRepeatParams类型，包含操作数相邻迭代间相同DataBlock的地址步长，操作数同一迭代内不同DataBlock的地址步长等参数。 |
表 3 数据类型约束
Atlas A2 训练系列产品/Atlas A2 推理系列产品 |
||||
Atlas A2 训练系列产品/Atlas A2 推理系列产品 |
||||
Atlas A2 训练系列产品/Atlas A2 推理系列产品 |
返回值说明
无
约束说明
-
操作数地址对齐要求请参见通用地址对齐约束。
-
操作数地址重叠约束请参考通用地址重叠约束。
-
使用tensor高维切分计算接口时，src和scalar的数据类型为half、dst的数据类型为float的情况下，一个迭代处理的源操作数元素个数需要和目的操作数保持一致，所以每次迭代选取前4个datablock参与计算。设置Repeat Stride参数和mask参数以及地址重叠时，需要考虑该限制。
调用示例
本样例中只展示Compute流程中的部分代码。如果您需要运行样例代码，请将该代码段拷贝并替换更多样例完整样例模板中Compute函数的部分代码即可。
-
tensor高维切分计算样例-mask连续模式
// repeatTime = 4, mask = 128, 128 elements one repeat, 512 elements total // srcLocal数据类型为half，scalar数据类型为half，dstLocal数据类型为half // dstBlkStride, srcBlkStride = 1, no gap between blocks in one repeat // dstRepStride, srcRepStride = 8, no gap between repeats AscendC::Axpy(dstLocal, srcLocal, (half)2.0, 128, 4,{ 1, 1, 8, 8 }); // srcLocal数据类型为half，scalar数据类型为half，dstLocal数据类型为float // repeatTime = 8, mask = 64, 64 elements one repeat, 512 elements total // dstBlkStride, srcBlkStride = 1, no gap between blocks in one repeat // dstRepStride = 8, srcRepStride = 4, no gap between repeats AscendC::Axpy(dstLocal, srcLocal, (half)2.0, 64, 8,{ 1, 1, 8, 4 }); // 每次迭代选取源操作数前4个datablock参与计算 -
tensor高维切分计算样例-mask逐bit模式
uint64_t mask[2] = { 0xFFFFFFFFFFFFFFFF, 0xFFFFFFFFFFFFFFFF }; // repeatTime = 4, 128 elements one repeat, 512 elements total, half精度组合 // dstBlkStride, srcBlkStride = 1, no gap between blocks in one repeat // dstRepStride, srcRepStride = 8, no gap between repeats AscendC::Axpy(dstLocal, srcLocal, (half)2.0, mask, 4,{ 1, 1, 8, 8 }); -
tensor前n个数据计算样例
AscendC::Axpy(dstLocal, src0Local, (half)2.0, 512);// half精度组合
更多样例
-
完整样例一：srcLocal、scalar、dstLocal的数据类型均为half。
#include "kernel_operator.h" class KernelAxpy { public: __aicore__ inline KernelAxpy() {} __aicore__ inline void Init(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm) { srcGlobal.SetGlobalBuffer((__gm__ half*)srcGm); 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>(); AscendC::Duplicate(dstLocal, (half)0.0, 512); AscendC::Axpy(dstLocal, srcLocal, (half)2.0, 512); 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::TPosition::VECIN, 1> inQueueSrc; AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueDst; AscendC::GlobalTensor<half> srcGlobal, dstGlobal; }; extern "C" __global__ __aicore__ void kernel_vec_ternary_scalar_Axpy_half_2_half(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm) { KernelAxpy op; op.Init(srcGm, dstGm); op.Process(); }结果示例如下：
输入数据(srcGm): [1. 1. 1. 1. 1. 1. ... 1.] 输出数据(dstGm): [2. 2. 2. 2. 2. 2. ... 2.] -
完整样例二：srcLocal、scalar的数据类型为half，dstLocal的数据类型为float。
#include "kernel_operator.h" class KernelAxpy { public: __aicore__ inline KernelAxpy() {} __aicore__ inline void Init(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm) { srcGlobal.SetGlobalBuffer((__gm__ half*)srcGm); dstGlobal.SetGlobalBuffer((__gm__ float*)dstGm); pipe.InitBuffer(outQueueDst, 1, 512 * sizeof(float)); pipe.InitBuffer(inQueueSrc, 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<float> dstLocal = outQueueDst.AllocTensor<float>(); AscendC::Duplicate(dstLocal, 0.0f, 512); AscendC::Axpy(dstLocal, srcLocal, (half)2.0, 64, 8, { 1, 1, 8, 4 }); outQueueDst.EnQue<float>(dstLocal); inQueueSrc.FreeTensor(srcLocal); } __aicore__ inline void CopyOut() { AscendC::LocalTensor<float> dstLocal = outQueueDst.DeQue<float>(); AscendC::DataCopy(dstGlobal, dstLocal, 512); outQueueDst.FreeTensor(dstLocal); } private: AscendC::TPipe pipe; AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueSrc; AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueDst; AscendC::GlobalTensor<half> srcGlobal; AscendC::GlobalTensor<float> dstGlobal; }; extern "C" __global__ __aicore__ void kernel_vec_ternary_scalar_Axpy_half_2_float(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm) { KernelAxpy op; op.Init(srcGm, dstGm); op.Process(); }结果示例如下：
输入数据(srcGm): [1. 1. 1. 1. 1. 1. ... 1.] 输出数据(dstGm): [2. 2. 2. 2. 2. 2. ... 2.]