ArithProgression
功能说明
给定起始值,等差值和长度,返回一个等差数列。
实现原理
以float类型,ND格式,firstValue和diffValue输入Scalar为例,描述ArithProgression高阶API内部算法框图,如下图所示。
计算过程分为如下几步,均在Vector上进行:
- 等差数列长度8以内步骤:按照firstValue和diffValued的值,使用SetValue实现等差数列扩充,扩充长度最大为8,如果等差数列长度小于8,算法结束;
- 等差数列长度8至64的步骤:对第一步中的等差数列结果使用Adds进行扩充,最大循环7次扩充至64,如果等差数列长度小于64,算法结束;
- 等差数列长度64以上的步骤:对第二步中的等差数列结果使用Adds进行扩充,不断循环直至达到等差数列长度为止。
函数原型
template <typename T> __aicore__ inline void ArithProgression(const LocalTensor<T> &dstLocal, const T firstValue, const T diffValue, const int32_t count)
参数说明
|
接口 |
功能 |
|---|---|
|
T |
操作数的数据类型。 |
|
参数名 |
输入/输出 |
描述 |
|---|---|---|
|
dstLocal |
输出 |
目的操作数。dstTensor的大小应大于等于 count * sizeof(T)。 类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。 Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/float/int16_t/int32_t Atlas推理系列产品AI Core,支持的数据类型为:half/float/int16_t/int32_t |
|
firstValue |
输入 |
等差数列的首个元素值。 Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/float/int16_t/int32_t Atlas推理系列产品AI Core,支持的数据类型为:half/float/int16_t/int32_t |
|
diffValue |
输入 |
等差数列元素之间的差值,应大于等于0。 Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:half/float/int16_t/int32_t Atlas推理系列产品AI Core,支持的数据类型为:half/float/int16_t/int32_t |
|
count |
输入 |
等差数列的长度。count>0。 |
返回值
无
支持的型号
Atlas A2训练系列产品/Atlas 800I A2推理产品
Atlas推理系列产品AI Core
注意事项
当前仅支持ND格式的输入,不支持其他格式。
调用示例
#ifdef ASCENDC_CPU_DEBUG
#include "tikicpulib.h"
#endif
#include "kernel_operator.h"
using namespace AscendC;
template <typename T> class KernelArithProgression {
public:
__aicore__ inline KernelArithProgression() {}
__aicore__ inline void Init(GM_ADDR dst_gm, int64_t firstValue, int64_t diffValue, uint32_t count)
{
firstValue_ = firstValue;
diffValue_ = diffValue;
count_ = count;
dst_global.SetGlobalBuffer(reinterpret_cast<__gm__ T *>(dst_gm), count_);
pipe.InitBuffer(outDst, 1, (sizeof(T) * count_ + 32 - 1) / 32 * 32);
}
__aicore__ inline void Process()
{
CopyIn();
Compute();
CopyOut();
}
private:
__aicore__ inline void CopyIn()
{
;
}
__aicore__ inline void Compute()
{
LocalTensor<T> dstLocal = outDst.AllocTensor<T>();
ArithProgression<T>(dstLocal, static_cast<T>(firstValue_), static_cast<T>(diffValue_), count_);
outDst.EnQue<T>(dstLocal);
}
__aicore__ inline void CopyOut()
{
LocalTensor<T> dstLocal = outDst.DeQue<T>();
const int32_t BLOCK_NUM = 32 / sizeof(T);
DataCopy(dst_global, dstLocal, (count_ + BLOCK_NUM - 1) / BLOCK_NUM * BLOCK_NUM);
outDst.FreeTensor(dstLocal);
}
private:
TPipe pipe;
TQue<QuePosition::VECOUT, 1> outDst;
GlobalTensor<T> dst_global;
int64_t firstValue_;
int64_t diffValue_;
uint32_t count_;
};
extern "C" __global__ __aicore__ void kernel_arith_progression_operator(GM_ADDR dstLocal)
{
KernelArithProgression<half> op;
op.Init(dstLocal, 1, 2, 15);
op.Process();
}