tensor前n个数据计算:
template <typename T> void Duplicate(const LocalTensor<T>& dstLocal, const T& scalarValue, const int32_t& calCount)
表1 参数说明
| 参数名称 | 输入/输出 | 含义 |
|---|---|---|
| dstLocal | 输出 | 目的操作数。 类型为LocalTensor,支持的TPosition为VECIN、VECCALC、VECOUT。 LocalTensor的起始地址需要32字节对齐。 Kirin9020系列处理器、Kirin9030系列处理器、KirinX90系列处理器,支持的数据类型为:uint8、int8、uint16_t、int16_t、half、uint32_t、int32_t、float16、float。 |
| scalarValue | 输入 | 被复制的源操作数,支持输入变量和立即数,数据类型需与dstLocal中元素的数据类型保持一致。 |
| calCount | 输入 | 输入数据元素个数。 |
本样例中只展示Compute流程中的部分代码。如果开发者需要运行样例代码,请将该代码段拷贝并替换样例模板中的Compute函数粗体部分即可。
tensor前n个数据计算样例:
half inputVal(18.0); AscendC::Duplicate<half>(dstLocal, inputVal, srcDataSize);
结果示例如下。
输入数据:[0 1.0 2.0 ... 254.0 255.0] // 不关心输入数据,会被Duplicate盖掉 输出数据:[18.0 18.0 18.0 ... 18.0 18.0]
#include "kernel_operator.h"
class KernelDuplicate {
public:
__aicore__ inline KernelDuplicate() {}
__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, srcDataSize * sizeof(half));
pipe.InitBuffer(outQueueDst, 1, dstDataSize * 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, srcDataSize);
inQueueSrc.EnQue(srcLocal);
}
__aicore__ inline void Compute()
{
AscendC::LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
AscendC::LocalTensor<half> dstLocal = outQueueDst.AllocTensor<half>();
**half inputVal(18.0);**
**AscendC::Duplicate<half>(dstLocal, inputVal, srcDataSize);**
outQueueDst.EnQue<half>(dstLocal);
inQueueSrc.FreeTensor(srcLocal);
}
__aicore__ inline void CopyOut()
{
AscendC::LocalTensor<half> dstLocal = outQueueDst.DeQue<half>();
AscendC::DataCopy(dstGlobal, dstLocal, dstDataSize);
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;
int srcDataSize = 256;
int dstDataSize = 256;
};
extern "C" __global__ __aicore__ void duplicate_kernel(__gm__ uint8_t* src, __gm__ uint8_t* dstGm)
{
KernelDuplicate op;
op.Init(src, dstGm);
op.Process();
}