Ascend C 数据搬运 API 教程与高性能实现

第一部分：基础概念

1.1 为什么数据搬运决定算子性能

Ascend AICore 是流水线式架构：MTE2（搬入）→ Vector/Cube（计算）→ MTE3（搬出）。计算单元的算力远高于 GM 带宽（950PR GM 约 1.6TB/s），因此绝大多数算子的瓶颈在搬运而非计算。一个常见现象是：profiling 中 aiv_mte2_ratio（MTE2 cycle 占总 cycle 比例）高达 95% 以上。

GM (HBM, ~1.6-1.8TB/s)  ──MTE2──►  UB/L1 (片上, ~5.2TB/s L2)
                                          │
                                          ▼
                                    Vector / Cube 计算
                                          │
                                          ▼  MTE3
                                       GM (写回)

数据搬运 API（DataCopy 系列）直接驱动 MTE2/MTE3 流水，其使用方式决定了：

  • 单次搬运粒度是否充分利用 GM 带宽
  • L2Cache 命中率（重复访问场景）
  • 多核间 GM 地址冲突
  • UB 内 bank 冲突（NZ 布局写）
  • 与计算流水的双缓冲重叠（double buffer)

1.2 存储层级与搬运通路

存储 位置 容量（950PR） 容量（A2/A3） 访问主体
GM（Global Memory / HBM） 片外 数十 GB 数十 GB 所有单元，经 L2Cache
L2Cache 片上 128MB 192MB 硬件自动管理，可配 Buffer 模式
L1 Buffer（A1/B1/C1/TSCM） AIC 片上 1MB 1MB CUBE / MTE
L0A / L0B / L0C CUBE 内部 64KB/64KB/256KB 同左 CUBE
UB（Unified Buffer / VECIN/VECCALC/VECOUT） AIV 片上 256KB 192KB Vector / MTE

主要搬运通路：

通路 指令 AscendC API 流水
GM → UB MOV_OUT_TO_UB_ALIGN DataCopy(ub, gm, ...) MTE2
UB → GM MOV_UB_TO_OUT_ALIGN DataCopy(gm, ub, ...) MTE3
GM → L1 MOV_OUT_TO_L1_ALIGN_V2 DataCopy(l1, gm, Nd2NzParams) MTE2
UB → L1 MOV_UB_TO_L1 DataCopy(l1, ub, ...) MTE3
UB → UB （内部搬运） Copy / DataCopy(ub, ub, ...) MTE

第二部分：基础 API 使用教程

2.1 最简单的连续搬运

场景：把 GM 上一段连续数据搬到 UB，计算后再搬回 GM。

#include "kernel_operator.h"

template <typename T, uint32_t totalLength>
class KernelDataCopy {
public:
    __aicore__ inline KernelDataCopy() {}
    __aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y)
    {
        xGm.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(x));
        yGm.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(y));
    }
    __aicore__ inline void Process()
    {
        // 静态 Tensor 编程：通过 LocalMemAllocator 分配 UB
        AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator;
        AscendC::LocalTensor<T> xLocal = ubAllocator.Alloc<T, totalLength>();

        // GM -> UB (MTE2)
        AscendC::DataCopy(xLocal, xGm, totalLength);

        // 同步：MTE2 完成后才能 MTE3
        AscendC::SetFlag<AscendC::HardEvent::MTE2_MTE3>(EVENT_ID0);
        AscendC::WaitFlag<AscendC::HardEvent::MTE2_MTE3>(EVENT_ID0);

        // UB -> GM (MTE3)
        AscendC::DataCopy(yGm, xLocal, totalLength);
    }
private:
    AscendC::GlobalTensor<T> xGm, yGm;
};

__global__ __vector__ void datacopy_custom(__gm__ uint8_t* x, __gm__ uint8_t* y)
{
    AscendC::InitSocState();          // 静态 Tensor 编程入口必备
    KernelDataCopy<float, 1024> op;
    op.Init(x, y);
    op.Process();
    AscendC::PipeBarrier<PIPE_ALL>(); // 出口必备
}

关键点：

  • InitSocState() / PipeBarrier<PIPE_ALL>() 是静态 Tensor 编程的入口/出口必备调用。
  • count 参数：搬运的元素个数（不是字节数），count * sizeof(T) 需 32B 对齐，否则向下取整。
  • 必须显式同步：MTE2 和 MTE3 是不同流水，没有自动依赖；通过 SetFlag/WaitFlag<HardEvent::MTE2_MTE3> 建立依赖。遗漏同步会导致读到未完成数据。

2.2 非连续搬运：DataCopyParams

场景：从 GM 的二维矩阵中按行搬运一个子块。

DataCopyParams 结构体控制非连续搬运的四个参数：

字段 类型 含义 单位
blockCount uint16_t 连续数据块个数 块数，[1, 4095]
blockLen uint16_t 每块的连续长度 DataBlock(32B)（950PR 在 C2 时为 32B，需偶数）
srcStride uint16_t 源相邻块间隔（前块尾→后块头） DataBlock(32B)
dstStride uint16_t 目的相邻块间隔（前块尾→后块头） DataBlock(32B) 
// 从 GM [128, 256] half 矩阵中搬运前 64 行、每行 128 个元素到 UB 连续存放
constexpr uint32_t M = 128, N = 256, tileM = 64, tileN = 128;
AscendC::DataCopyParams params;
params.blockCount = tileM;                              // 64 行
params.blockLen   = tileN * sizeof(half) / 32;          // 每行 128*2B = 256B = 8 个 DataBlock
params.srcStride  = (N - tileN) * sizeof(half) / 32;    // 源每行末跳过 128 个 half = 256B = 8 DataBlock
params.dstStride  = 0;                                  // UB 中连续存放
AscendC::DataCopy(ubLocal, srcGm[mIdx * N + nIdx], params);

记忆口诀：总搬运量 = blockCount × blockLen × 32B；srcStride/dstStride 是块与块之间的间隔，不是步长。

2.3 非对齐搬运：DataCopyPad

场景：当 count * sizeof(T) 不满足 32B 对齐，或需要左右 padding 时，用 DataCopyPad 替代 DataCopy。

// 场景：GM [32, 59] float → UB [32, 64]，右侧补 5 个 0
AscendC::DataCopyExtParams copyParams{
    /*blockCount=*/32,
    /*blockLen=*/59 * sizeof(float),   // 236B，非 32B 对齐
    /*srcStride=*/0,
    /*srcRepStride=*/0,
    /*reserved=*/0
};
AscendC::DataCopyPadExtParams<float> padParams;
padParams.isPad = true;                // true: padding 值为 0
padParams.leftPadding = 0;
padParams.rightPadding = 5;            // 右侧补 5 个元素到 64

AscendC::DataCopyPad(ubLocal, srcGlobal, copyParams, padParams);

isPad 与 SetPadValue 的关系：

  • isPad=true：固定 padding 值为 0
  • isPad=false + SetPadValue(value)：自定义 padding 值

Compact 模式（仅 950PR，PaddingMode::Compact）：多行数据紧凑排列到一行，padding 区域紧跟数据之后。例如 [3, 24] half 紧凑为 [1, 80]（72 数据 + 8 padding）。

2.4 切片搬运：SliceInfo

场景：从多维 Tensor 中提取任意子集，支持 stride 间隔采样。

// 从 GM [3, 87] 中按 srcSlice 规则取子集，写入 UB [2, 48]
AscendC::SliceInfo srcSliceInfo[] = {
    {16, 70, 7, 3, 87},   // dim0: startIndex=16, endIndex=70, stride=7, burstLen=3, shapeValue=87
    {0,  2,  1, 1, 3}     // dim1: startIndex=0,  endIndex=2,  stride=1, burstLen=1, shapeValue=3
};
AscendC::SliceInfo dstSliceInfo[] = {
    {0, 47, 0, 3, 48},
    {0, 1,  0, 1, 2}
};
uint32_t dimValue = 2;
AscendC::DataCopy(ubLocal, srcGlobal, dstSliceInfo, srcSliceInfo, dimValue);

SliceInfo 字段：

字段 含义
startIndex 源维度起始索引
endIndex 源维度结束索引
stride 采样步长（1 = 连续）
burstLen 每次 burst 的 DataBlock 数
shapeValue 该维度总长度

支持通路：仅 GM↔UB（不支持 UB↔UB）。

2.5 ND2NZ / NZ2ND 随路格式转换

场景：CUBE 矩阵乘要求 L1 中的数据是 NZ（分形）布局，但 GM 中通常是 ND（行优先）。用 DataCopy 配合 Nd2NzParams 在搬入 L1 时同时完成格式转换。

// GM [M, N] half ND → L1 NZ 布局
AscendC::Nd2NzParams nd2nzParams;
nd2nzParams.ndNum = 1;
nd2nzParams.nValue = tileM;                         // 本次搬运行数
nd2nzParams.dValue = tileN;                         // 本次搬运列数
nd2nzParams.srcNdMatrixStride = 0;
nd2nzParams.srcDValue = N;                          // GM 中行宽
nd2nzParams.dstNzC0Stride = AlignUp(tileM, 16);     // L1 中 C0 方向跨度（16 对齐）
nd2nzParams.dstNzNStride = 1;
nd2nzParams.dstNzMatrixStride = 0;
AscendC::DataCopy(l1Local, srcGlobal[mIdx * N + nIdx], nd2nzParams);

NZ 布局回顾：把 [M, N] 切成 [M/16, N/16, 16, 16] 的小块（fractal），每个 16×16 块连续存放。dstNzC0Stride 控制 L1 中相邻 C0 列的间距，是 bank 冲突调优的核心旋钮（见 3.5）。

NZ2ND 反向转换（L0C/UB → GM）通过 Fixpipe 或 DataCopy(... Nz2ndParams ...) 实现，常用于矩阵乘结果写回。

2.6 多维搬运（NdDma，仅 950PR）

场景：950PR 提供 NdDmaParams 实现真正的 N 维 DMA（Padding / Transpose / Broadcast / Slice），比 SliceInfo 更灵活。

// 2D Padding: GM [16,32] → UB [32,64]，左pad15 上pad13 右pad17 下pad3
AscendC::NdDmaLoopInfo<2> loopInfo{
    {1, 32},    // loop1Size, loop2Size
    {1, 64},    // dstLoop1Stride, dstLoop2Stride (DataBlock)
    {32, 16},   // srcLoop1Stride, srcLoop2Stride
    {15, 13},   // leftPadding,  topPadding
    {17, 3}     // rightPadding, bottomPadding
};
AscendC::NdDmaParams<T, 2> params{loopInfo, /*padValue=*/0};
AscendC::NdDmaDci();  // 刷新 cache
static constexpr AscendC::NdDmaConfig dmaConfig;
AscendC::DataCopy<T, 2, dmaConfig>(xLocal, xGm, params);

NdDma 支持的 5 种场景（见样例 data_copy_gm2ub_nddma）：

  1. 2D Padding：四周补齐
  2. 2D Padding + 最近值填充（NdDmaConfig{isNearestValueMode=true}）
  3. 2D Transpose：通过设置 srcLoopStride 与 dstLoopStride 互换实现
  4. 2D Broadcast：loopInfo.loop1Size={1,0} 表示该维广播
  5. 2D Slice：截取子矩阵

2.7 Loop Mode（950PR，多层循环硬件化）

场景：把两层 for 循环的搬运交给硬件一次下发，减少指令数。

// 把 [2,2,40B,2] 的多层循环搬运一次下发
AscendC::LoopModeParams loopParam{
    /*loop1Size=*/2, /*loop2Size=*/2,
    /*loop1SrcStride=*/80,  /*loop1DstStride=*/128,
    /*loop2SrcStride=*/160, /*loop2DstStride=*/288
};
AscendC::SetLoopModePara(loopParam, AscendC::DataCopyMVType::OUT_TO_UB);
AscendC::DataCopyPad<int8_t>(srcLocal, srcGlobal, copyParams, padParams);
AscendC::ResetLoopModePara(AscendC::DataCopyMVType::OUT_TO_UB);  // 用完必须复位

适用：5 维 NLP/CV 数据搬运（如 [batch, head, seq, 128, 126] → [512, 128]），用 loop mode 配合外层 for 循环可硬件化多层循环。

第三部分：高性能实现技术

3.1 优化点 1：增大分块粒度

原理：单次 DataCopy 搬运量越大，指令发射开销和循环控制开销越低，MTE2 有效搬运效率越高。

实测数据（来自 data_copy 最佳实践，GM → UB，half [12288,12288]，48 核）：

分块 单次搬运量 MTE2 耗时(μs) 相对基线提升
Tile=[1,64] 128B 548.16 基线
Tile=[64,64] 8KB 220.77 +148%
Tile=[64,1024] 128KB 202.66 +170%

结论：在片上空间允许的前提下，优先增大单次搬运大小。UB 256KB（950PR）可容纳 tileM×tileN ≤ 128KB（half 下即 64×1024 或 128×512），留另一半给双缓冲。

经验值：单次搬运 blockLen × 32B 建议 ≥ 512B（A2/A3）或 ≥ 128B（950PR），避免过小粒度。

3.2 优化点 2：保持主维度对齐

原理：非对齐尾块会触发边界处理，搬运效率下降。

实测数据（GM → L1，half [12288, N]，Tile=[64,256]）：

N 对齐情况 A2/A3 MTE2(μs) 950PR MTE2(μs)
12288 对齐 214.52 187.19
12287 非对齐 422.52（-47.5%） 202.97（-7.8%）

结论：

  • 设计 shape 或 tiling 时，主搬运维度尽量被 TILE_N 整除。
  • 连续搬运字节数建议 A2/A3 ≥ 512B 对齐，950PR ≥ 128B 对齐。
  • 无法避免非对齐时，用 DataCopyPad 显式 padding 到对齐。

3.3 优化点 3：L2Cache 复用（分片重复访问）

原理：L2Cache 128MB（950PR）/ 192MB（A2/A3），工作集小于 L2 容量时第二次访问可命中。整块重复搬运时工作集过大，L2 难以保留；先分片再在片内重复访问可显著提升命中率。

模式对比：

场景A（整块重复，低效）：            场景B（分片重复，高效）：
for r in 4:                         for split in 4:
    copy 全矩阵 → UB/L1                 copy 分片split → UB/L1
    // L2 难保留全矩阵                   for r in 4:
                                            copy 同一分片（L2 命中）

实测数据（GM → UB，half [12288,12288]，重复 4 次）：

模式 A2/A3 耗时(μs) L2 命中率 950PR 耗时(μs) L2 命中率
整块重复 828.06 0.005% 741.58 0.35%
N/4 分片重复 365.74（-56%） 75.0% 354.95（-52%） 66.6%

结论：同一批 GM 数据需要多次读取时（如多轮迭代、反向传播），优先分片后在分片内连续完成多次访问，使单次工作集落在 L2Cache 可复用范围。

3.4 优化点 4：多核同地址访问冲突规避

原理：多核同时访问相同 GM 地址段会触发仲裁，降低有效带宽。通过按核错开访问顺序可规避。

// same addr（低效）：所有核按相同 mBlockIdx 顺序访问
uint32_t curMBlockIdx = mBlockIdx;

// offset addr（高效）：每个核在组内轮转
uint32_t blockGroupStart = (mBlockIdx / numBlocks) * numBlocks;
uint32_t curMBlockIdx = blockGroupStart + (mBlockIdx + blockIdx) % numBlocks;

实测数据（GM → L1，half [6144,512]，24 核）：

模式 A2/A3 耗时(μs) 950PR 耗时(μs)
same addr 278.56 369.99
offset addr 221.34（-20%） 187.35（-49%）

结论：多核读取同一大块 GM 数据时，按 blockIdx 轮转访问分片顺序，降低同一时刻多核访问同地址的概率。

3.5 优化点 5：UB 内 Bank 冲突调优（NZ 布局）

原理：UB 由多个 memory bank 组成，同一时钟周期内对同一 bank 的多次访问会串行化。ND2NZ 写 UB 时，相邻 C0 列的落点间距 dstNzC0Stride 决定是否冲突。

调优旋钮：dstNzC0Stride（以 DataBlock=32B 为单位）

case dstNzC0Stride 是否 bank 冲突 性能
1 144（=tileH，自然值） 是 基线
2 145（+1 错开） 否 显著提升 
// 调优方法：把 dstNzC0Stride 从 tileH 改为 tileH+1
nd2nzParams.dstNzC0Stride = tileH + 1;   // 错开 bank

结论：ND2NZ / NZ2ND 等涉及 UB 重排的搬运，优先检查 dstNzC0Stride 是否为 bank 数（通常 16 或 32）的整数倍，若是则 +1 错开。

3.6 优化点 6：Double Buffer（搬运与计算重叠）

原理：把 UB 分成两半，当计算单元处理 buffer A 时，MTE2 同时搬运下一块数据到 buffer B，循环交替，实现搬运与计算的完全重叠。

时间轴 →
MTE2:  [搬入A] [搬入B] [搬入A] [搬入B] ...
Vector:        [算A]   [算B]   [算A]   ...

框架 API 实现（最简单）：

AscendC::TPipe pipe;
AscendC::TQue<AscendC::TPosition::VECIN, AscendC::HardEvent::MTE2_V> inQueue;
pipe.InitBuffer(inQueue, 2, sizeof(T) * tileLen);  // 2 = double buffer
auto ubLocal = inQueue.AllocTensor<T>();
AscendC::DataCopy(ubLocal, gm, tileLen);
inQueue.EnQue(ubLocal);
auto input = inQueue.DeQue<T>();
AscendC::Add(...);
inQueue.FreeTensor(input);

基础 API 手动实现：

AscendC::LocalTensor<T> bufA(AscendC::TPosition::VECIN, addrA, tileLen);
AscendC::LocalTensor<T> bufB(AscendC::TPosition::VECIN, addrB, tileLen);
AscendC::DataCopy(bufA, gm, tileLen);  // 预取第一块到 A
AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);

for (int i = 0; i < nLoop; i++) {
    AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
    // 启动下一块搬入到另一 buffer
    if (i + 1 < nLoop) {
        AscendC::DataCopy((i % 2 == 0) ? bufB : bufA, gm + (i+1)*tileLen, tileLen);
        AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
    }
    // 计算当前 buffer
    AscendC::Add(out, (i % 2 == 0) ? bufA : bufB, ..., tileLen);
    AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
    AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
    AscendC::DataCopy(gmOut + i*tileLen, out, tileLen);
}

约束：UB 空间需容纳 2×tileLen，tileLen 选择要平衡"搬运隐藏延迟"与"UB 容量"。

3.7 优化点 7：L2Cache Hint

场景：对流式数据（只访问一次，不重复）显式关闭 L2Cache hint，避免污染 cache；对重复访问数据保持默认。

srcGlobal.SetL2CacheHint(AscendC::CacheMode::CACHE_MODE_DISABLE);  // 流式，不缓存
AscendC::DataCopy(ub, srcGlobal, ...);

// 或针对特定搬运
srcGlobal.SetL2CacheHint(AscendC::CacheMode::CACHE_MODE_READ_ALLOCATE);  // 读预分配

约束：仅支持 HCCS 通路，不支持其他通路（如 PCIE）。

3.8 优化点 8：Batch Mode 合并

原理：Batch mode 规定，当 srcStride=0 且 dstStride=0 时，多个 burst 合并为一条 uop，减少 uop 数量。

// 低效：blockCount=4, srcStride=1, dstStride=1 → 4 条 uop
// 高效：blockCount=4, srcStride=0, dstStride=0 → 1 条 uop（连续搬运）
AscendC::DataCopyParams params{4, 8, 0, 0};  // 4 块 × 8 DataBlock，连续
AscendC::DataCopy(ub, gm, params);

适用条件：源和目的都连续（无 stride），且非 Byte transfer mode。

第四部分：高性能算子

优化点 代码位置 收益
大分块 tileM=64, tileN=1024 template 参数 MTE2 效率 +170%
对齐 padding DataCopyPad + curCols 规避非对齐 -47%
Double Buffer bufA/bufB 交替 + MTE2_V 同步 搬运/计算重叠
L2Cache 复用 N 方向分片内连续处理 命中率 0.3% → 66%
多核错开 mStart = blockIdx * singleCoreM 规避同地址冲突

第五部分：API 速查表

5.1 按场景选 API

场景 首选 API 备选
连续搬运 GM↔UB DataCopy(dst, src, count) -
非连续二维搬运 DataCopy(dst, src, DataCopyParams) -
非对齐/padding DataCopyPad -
多维子集提取 DataCopy(SliceInfo...) 950PR 用 NdDmaParams
ND→NZ（入 L1） DataCopy(l1, gm, Nd2NzParams) -
NZ→ND（出 GM） Fixpipe / DataCopy(...Nz2ndParams...) -
950PR N 维 padding/transpose/broadcast DataCopy<T, N, NdDmaConfig>(ub, gm, NdDmaParams) -
950PR 多层循环硬件化 SetLoopModePara + DataCopyPad -
UB→UB 简单复制 Copy DataCopy(ub, ub, ...)
L1→UB DataCopyL1ToUB DataCopy(ub, l1, ...)
L0C→UB/GM（矩阵乘结果） Fixpipe asc_copy_l0c2ub（C API）
跨卡搬运（A2/A3） DataCopy（仅 HCCS） HCCL 高阶 API

5.2 参数对齐速查

参数 单位 对齐要求
count 元素 count * sizeof(T) 32B 对齐，否则向下取整
DataCopyParams.blockLen DataBlock(32B) 950PR C2 位置需偶数
DataCopyParams.srcStride/dstStride DataBlock(32B) uint16 范围
DataCopyExtParams.blockLen 字节 任意（DataCopyPad）
Nd2NzParams.dstNzC0Stride DataBlock(32B) 16 对齐，建议 +1 错开 bank
LocalTensor 起始地址 - 32B；C2 位置 64B；C2PIPE2GM 128B（950PR 64B）
GlobalTensor 起始地址 - 按数据类型大小（half 2B、float 4B…）

总结

数据搬运高性能实现的核心五原则：

  1. 大分块：单次搬运 ≥ 128KB，充分利用 GM 带宽
  2. 主维对齐：主搬运维度被 TILE_N 整除，连续字节数 ≥ 512B（A2/A3）/ 128B（950PR）
  3. 分片复用 L2Cache：重复访问数据先分片（< 128MB），片内连续多次访问
  4. 多核错开访问：按 blockIdx 轮转 GM 分片顺序，规避同地址冲突
  5. 搬运/计算重叠：Double Buffer + 轻量 SetFlag/WaitFlag 同步，隐藏搬运延迟

附加优化：UB 内 bank 冲突调优（dstNzC0Stride +1）、Batch mode 合并 uop（srcStride=0）、L2Cache hint 显式控制、950PR SSBuffer 硬通道。

# 人工智能
Logo
人工智能6S服务平台

作为“人工智能6S店”的官方数字引擎，为AI开发者与企业提供一个覆盖软硬件全栈、一站式门户。

更多推荐

cover

HarmonyOS掌上记账APP开发实践第60篇:数据报表与统计 — 按月分组的账单聚合分析

avatar 人工智能6S服务平台
cover

鸿蒙-NestedScrollTabContainer-mvvm-demo

avatar 人工智能6S服务平台

【HarmonyOS 5】鸿蒙检测URL是否为恶意地址

文章摘要：HarmonyOS通过DeviceSecurityKit提供恶意URL检测能力，核心API为checkUrlThreat，支持异步检测钓鱼、恶意软件等风险URL。开发者需在AGC平台开通服务并配置Profile，调用时需构造UrlCheckRequest对象并处理返回的威胁类型结果。该服务每天每设备限调用1万次，云端实时更新黑名单。开发示例展示了ArkTS实现流程，同时系统还提供代码混淆

avatar 人工智能6S服务平台