Copy
函数功能
VECIN, VECCALC, VECOUT之间的搬运指令,支持mask操作和datablock间隔操作。
函数原型
- mask逐bit模式
1 2
template <typename T, bool isSetMask = true> __aicore__ inline void Copy(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const uint64_t mask[], const uint8_t repeatTimes, const CopyRepeatParams& repeatParams)
- mask连续模式
1 2
template <typename T, bool isSetMask = true> __aicore__ inline void Copy(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, const uint64_t mask, const uint8_t repeatTimes, const CopyRepeatParams& repeatParams)
参数说明
参数名 |
描述 |
---|---|
T |
操作数数据类型。 Atlas A2训练系列产品/Atlas 800I A2推理产品,支持的数据类型为:uint16_t/int16_t/half/bfloat16_t/uint32_t/int32_t/float Atlas 200/500 A2推理产品,支持的数据类型为:uint16_t/int16_t/half/bfloat16_t/uint32_t/int32_t/float |
isSetMask |
是否在接口内部设置mask。
|
参数名 |
输入/输出 |
描述 |
||
---|---|---|---|---|
dstLocal |
输出 |
目的操作数。 类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。起始地址需要保证32字节对齐。 |
||
srcLocal |
输入 |
源操作数。 类型为LocalTensor,支持的TPosition为VECIN/VECCALC/VECOUT。起始地址需要保证32字节对齐。 源操作数的数据类型需要与目的操作数保持一致。 |
||
mask |
输入 |
|
||
repeatTimes |
输入 |
重复迭代次数。矢量计算单元,每次读取连续的256 Bytes数据进行计算,为完成对输入数据的处理,必须通过多次迭代(repeat)才能完成所有数据的读取与计算。repeatTimes表示迭代的次数。 |
||
CopyRepeatParams |
输入 |
控制操作数地址步长的数据结构。CopyRepeatParams类型,定义如下:
dstRepeatSize 、srcRepeatSize用于设置相邻迭代间的地址步长,取值范围为[0,4095];dstStride 、srcStride用于设置同一迭代内datablock的地址步长,取值范围为[0,65535]。 相邻迭代间的地址步长参数说明请参考repeatStride(相邻迭代间相同datablock的地址步长);同一迭代内datablock的地址步长参数说明请参考dataBlockStride(同一迭代内不同datablock的地址步长)。 |
返回值
无
支持的型号
Atlas A2训练系列产品/Atlas 800I A2推理产品
Atlas 200/500 A2推理产品
约束说明
源操作数和目的操作数的起始地址需要保证32字节对齐。
调用示例
本样例中只展示Compute流程中的部分代码。如果您需要运行样例代码,请将该代码段拷贝并替换样例模板中Compute函数的部分代码即可。
- mask连续模式
1 2 3 4 5
uint64_t mask = 128; // repeatTimes = 4, 128 elements one repeat, 512 elements total // dstStride, srcStride = 1, no gap between blocks in one repeat // dstRepStride, srcRepStride = 8, no gap between repeats AscendC::Copy(dstLocal, srcLocal, mask, 4, { 1, 1, 8, 8 });
结果示例如下:
输入数据(srcLocal): [9 -2 8 ... 9 0] 输出数据(dstLocal): [9 -2 8 ... 9]
- mask逐bit模式
1 2 3 4 5
uint64_t mask[2] = { UINT64_MAX, UINT64_MAX }; // repeatTimes = 4, 128 elements one repeat, 512 elements total // dstStride, srcStride = 1, no gap between blocks in one repeat // dstRepStride, srcRepStride = 8, no gap between repeats AscendC::Copy(dstLocal, srcLocal, mask, 4, { 1, 1, 8, 8 });
结果示例如下:
输入数据(srcLocal): [9 -2 8 ... 9 0] 输出数据(dstLocal): [9 -2 8 ... 9]
样例模板
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 |
#include "kernel_operator.h" class KernelCopy { public: __aicore__ inline KernelCopy() {} __aicore__ inline void Init(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm) { srcGlobal.SetGlobalBuffer((__gm__ int32_t*)srcGm); dstGlobal.SetGlobalBuffer((__gm__ int32_t*)dstGm); pipe.InitBuffer(inQueueSrc, 1, 512 * sizeof(int32_t)); pipe.InitBuffer(outQueueDst, 1, 512 * sizeof(int32_t)); } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensor<int32_t> srcLocal = inQueueSrc.AllocTensor<int32_t>(); AscendC::DataCopy(srcLocal, srcGlobal, 512); inQueueSrc.EnQue(srcLocal); } __aicore__ inline void Compute() { AscendC::LocalTensor<int32_t> srcLocal = inQueueSrc.DeQue<int32_t>(); AscendC::LocalTensor<int32_t> dstLocal = outQueueDst.AllocTensor<int32_t>(); uint64_t mask = 64; AscendC::Copy(dstLocal, srcLocal, mask, 4, { 1, 1, 8, 8 }); outQueueDst.EnQue<int32_t>(dstLocal); inQueueSrc.FreeTensor(srcLocal); } __aicore__ inline void CopyOut() { AscendC::LocalTensor<int32_t> dstLocal = outQueueDst.DeQue<int32_t>(); AscendC::DataCopy(dstGlobal, dstLocal, 512); outQueueDst.FreeTensor(dstLocal); } private: AscendC::TPipe pipe; AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueSrc; AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueueDst; AscendC::GlobalTensor<int32_t> srcGlobal, dstGlobal; }; extern "C" __global__ __aicore__ void copy_simple_kernel(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm) { KernelCopy op; op.Init(srcGm, dstGm); op.Process(); } |