Copy

函数功能

VECIN, VECCALC, VECOUT之间的搬运指令，支持mask操作和Block间隔操作。

函数原型

• mask参数使用逐bit模式，该模式的具体介绍请参考表1中的mask参数说明：

  template <typename T> __aicore__ inline void Copy(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, uint64_t mask[2], const uint8_t repeatTimes, const CopyRepeatParams& repeatParams);

• mask参数使用连续模式，该模式的具体介绍请参考表1中的mask参数说明：

  template <typename T> __aicore__ inline void Copy(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, uint64_t mask, const uint8_t repeatTimes, const CopyRepeatParams& repeatParams);

数据结构

CopyRepeatParams为用于控制操作数地址步长的数据结构。结构体内包含操作数相邻迭代间相同block的地址步长，操作数同一迭代内不同block的地址步长等参数。

相邻迭代间相同block的地址步长参数的详细说明请参考相邻迭代间相同datablock的地址步长；同一迭代内不同block的地址步长参数请参考同一迭代内不同datablock的地址步长。

结构体具体定义为：

struct CopyRepeatParams {
    dstStride = kDefaultDataCopyStride;
    srcStride = kDefaultDataCopyStride;
    dstRepeatSize = kDefaultRepStride;
    srcRepeatSize = kDefaultRepStride;
};

用户需要自行定义Block stride参数，包含dstStride ，srcStride ，以及Repeat stride参数，包含dstRepeatSize ，srcRepeatSize 。

参数说明

返回值

无

支持的型号

Atlas A2训练系列产品/Atlas 800I A2推理产品

约束说明

无

调用示例

本样例中只展示Compute流程中的部分代码。如果您需要运行样例代码，请将该代码段拷贝并替换样例模板中Compute函数的部分代码即可。

• mask连续模式

  uint64_t mask = 128;
  // repeatTimes = 4, 128 elements one repeat, 512 elements total
  // dstBlkStride, srcBlkStride = 1, no gap between blocks in one repeat
  // dstRepStride, srcRepStride = 8, no gap between repeats
  Copy(dstLocal, srcLocal, mask, 4, { 1, 1, 8, 8 });

  结果示例如下：

  输入数据(srcLocal): [9 -2 8 ... 9 0]
  输出数据(dstLocal): 
  [9 -2 8 ... 9]

• mask逐bit模式

  uint64_t mask[2] = { UINT64_MAX, UINT64_MAX };
  // repeatTimes = 4, 128 elements one repeat, 512 elements total
  // dstBlkStride, srcBlkStride = 1, no gap between blocks in one repeat
  // dstRepStride, srcRepStride = 8, no gap between repeats
  Copy(dstLocal, srcLocal, mask, 4, { 1, 1, 8, 8 });

  结果示例如下：

  输入数据(srcLocal): [9 -2 8 ... 9 0]
  输出数据(dstLocal): 
  [9 -2 8 ... 9]

样例模板

#include "kernel_operator.h"
namespace AscendC {
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()
    {
        LocalTensor<int32_t> srcLocal = inQueueSrc.AllocTensor<int32_t>();
        DataCopy(srcLocal, srcGlobal, 512);
        inQueueSrc.EnQue(srcLocal);
    }
    __aicore__ inline void Compute()
    {
        LocalTensor<int32_t> srcLocal = inQueueSrc.DeQue<int32_t>();
        LocalTensor<int32_t> dstLocal = outQueueDst.AllocTensor<int32_t>()
        uint64_t mask = 64;
        Copy(dstLocal, srcLocal, mask, 4, { 1, 1, 8, 8 });
        outQueueDst.EnQue<int32_t>(dstLocal);
        inQueueSrc.FreeTensor(srcLocal);
    }
    __aicore__ inline void CopyOut()
    {
        LocalTensor<int32_t> dstLocal = outQueueDst.DeQue<int32_t>();
        DataCopy(dstGlobal, dstLocal, 512);
        outQueueDst.FreeTensor(dstLocal);
    }
private:
    TPipe pipe;
    TQue<QuePosition::VECIN, 1> inQueueSrc;
    TQue<QuePosition::VECOUT, 1> outQueueDst;
    GlobalTensor<int32_t> srcGlobal, dstGlobal;
};
} // namespace AscendC
extern "C" __global__ __aicore__ void copy_simple_kernel(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm)
{
    AscendC::KernelCopy op;
    op.Init(srcGm, dstGm);
    op.Process();
}