Copy
产品支持情况
功能说明
VECIN，VECCALC，VECOUT之间的搬运指令，支持mask操作和DataBlock间隔操作。
函数原型
-
tensor前n个数据计算
template <typename T, bool isSetMask = true> __aicore__ inline void Copy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const uint32_t count) -
tensor高维切分计算
-
mask逐bit模式
template <typename T, bool isSetMask = true> __aicore__ inline void Copy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const uint64_t mask[], const uint8_t repeatTime, const CopyRepeatParams& repeatParams) -
mask连续模式
template <typename T, bool isSetMask = true> __aicore__ inline void Copy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const uint64_t mask, const uint8_t repeatTime, const CopyRepeatParams& repeatParams)
-
参数说明
表 1 模板参数说明
|
Ascend 950PR/Ascend 950DT，支持的数据类型为：uint8_t/int8_t/hifloat8_t/fp8_e4m3fn_t/fp8_e5m2_t/fp4x2_e2m1_t/fp4x2_e1m2_t/fp8_e8m0_t/uint16_t/int16_t/half/bfloat16_t/float/uint32_t/int32_t/uint64_t/int64_t Atlas A3 训练系列产品/Atlas A3 推理系列产品，支持的数据类型为：uint16_t/int16_t/half/bfloat16_t/uint32_t/int32_t/float Atlas A2 训练系列产品/Atlas A2 推理系列产品，支持的数据类型为：uint16_t/int16_t/half/bfloat16_t/uint32_t/int32_t/float Kirin X90，支持的数据类型为：half/float/uint16_t/int16_t/uint32_t/int32_t Kirin 9030，支持的数据类型为：half/float/uint16_t/int16_t/uint32_t/int32_t |
|
|
表 2 参数说明
|
类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。起始地址需要保证32字节对齐。 |
||
|
类型为LocalTensor，支持的TPosition为VECIN/VECCALC/VECOUT。起始地址需要保证32字节对齐。 |
||
|
||
重复迭代次数。矢量计算单元，每次读取连续的256Bytes数据进行计算，为完成对输入数据的处理，必须通过多次迭代（repeat）才能完成所有数据的读取与计算。repeatTime表示迭代的次数。 |
||
控制操作数地址步长的数据结构。CopyRepeatParams类型。 具体定义请参考${INSTALL_DIR}/include/ascendc/basic_api/interface/kernel_struct_data_copy.h，${INSTALL_DIR}请替换为CANN软件安装后文件存储路径。 参数说明请参考表3。 |
表 3 CopyRepeatParams结构体参数说明
返回值说明
无
约束说明
- 源操作数和目的操作数的起始地址需要保证32字节对齐。
- tensor前n个数据计算接口仅支持Ascend 950PR/Ascend 950DT。
- 针对Ascend 950PR/Ascend 950DT，uint8_t/int8_t/hifloat8_t/fp8_e4m3fn_t/fp8_e5m2_t/fp4x2_e2m1_t/fp4x2_e1m2_t/fp8_e8m0_t/uint64_t/int64_t数据类型仅支持tensor前n个数据计算接口。
- 针对Ascend 950PR/Ascend 950DT，tensor前n个数据计算接口中的isSetMask参数不生效，保持默认值即可。
- Copy和矢量计算API一样，支持和掩码操作API配合使用。但Counter模式配合高维切分计算API时，和通用的Counter模式有一定差异。具体差异如下：
-
通用的Counter模式：Mask代表整个矢量计算参与计算的元素个数，迭代次数不生效。
-
Counter模式配合Copy高维切分计算API，Mask代表**每次Repeat中处理的元素个数，迭代次数生效。**示意图如下：
-
调用示例
本示例仅展示Compute流程中的部分代码。如需运行，请参考样例模板实现完整的代码。
本示例中操作数数据类型为int16_t。
-
tensor前n个数据计算
AscendC::Copy(dstLocal, srcLocal, 512);结果示例如下：
输入数据srcLocal：[9 -2 8 ... 9] 输出数据dstLocal: [9 -2 8 ... 9] -
mask连续模式
uint64_t mask = 128; // repeatTime = 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] 输出数据dstLocal: [9 -2 8 ... 9] -
mask逐bit模式
uint64_t mask[2] = { UINT64_MAX, UINT64_MAX }; // repeatTime = 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] 输出数据dstLocal: [9 -2 8 ... 9]
样例模板
#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::TPosition::VECIN, 1> inQueueSrc;
AscendC::TQue<AscendC::TPosition::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();
}