初次上手昇腾 CANN 做算子开发，DataCopy 是最基础、也是最高频的 AscendC 内置函数。很多新手不清楚它的作用，今天单独拆解这个函数，附最简可运行内核代码，快速理解昇腾 AI Core 的数据搬移逻辑。

DataCopy 函数简介

在昇腾异构架构中，AI Core 计算单元只能直接读写片上 Local Memory，无法直接访问全局内存 GM。 DataCopy：AscendC 提供的内存拷贝内置函数，负责在全局内存__gm__和片上LocalTensor之间搬运张量数据，是算子开发必备函数。 函数原型简化理解：

template <typename T>
__aicore__ void DataCopy(LocalTensor<T> dst, const __gm__ T* src, uint32_t count);

• dst：目标片上 LocalTensor
• src：全局内存的数据源指针
• count：拷贝元素数量

完整极简 Kernel 代码示例

#include "kernel_operator.h"
using namespace AscendC;

__global__ void SimpleCopyKernel(__gm__ const float* in, __gm__ float* out, uint32_t dataLen)
{
    // 定义片上张量，单次搬运128个float
    LocalTensor<float> localBuf = AllocTensor<float>(128);

    // DataCopy：全局内存 → 片上Local
    DataCopy(localBuf, in, 128);

    // DataCopy：片上Local → 全局内存
    DataCopy(out, localBuf, 128);

    FreeTensor(localBuf);
}

函数使用要点

1. DataCopy 是面向向量的批量拷贝，比循环逐个赋值性能高很多，充分利用昇腾 NPU 的数据搬运硬件流水线；
2. 拷贝长度尽量对齐硬件要求，长度不合适会出现截断、越界报错；
3. LocalTensor 用完必须FreeTensor释放，否则会造成片上 SRAM 内存泄漏；
4. 不要在没有 LocalTensor 的情况下，直接在 AI Core 内读写__gm__指针，会编译失败。

踩坑记录

最常见错误：传入的拷贝数量超过 LocalTensor 申请的空间大小，触发内存越界。写代码时，AllocTensor分配长度必须 ≥ DataCopy 的拷贝元素个数。

小结：昇腾算子本质就是「DataCopy 搬数据 → 内置函数计算 → DataCopy 写回结果」。掌握 DataCopy，才算迈出 AscendC 算子开发的第一步。 后续可以基于这个函数，叠加 Add、Mul 等向量计算函数实现简单算法。