---
tags: [CANN]
doc_type: knowledge-note
返回目录: "[[计算机系/昇腾与CANN/大模型与算子专题学习导航|大模型与算子专题学习导航]]"
上一篇: "[[计算机系/昇腾与CANN/流水线双向同步原理|流水线双向同步原理]]"
下一篇: "[[计算机系/昇腾与CANN/Ascend C矩阵算子与数据布局|Ascend C矩阵算子与数据布局]]"
前置知识: "[[计算机系/昇腾与CANN/Ascend C SIMD硬件内存与同步|Ascend C SIMD硬件内存与同步]]"
---

# Ascend C Memory矢量算子开发

Memory 矢量编程以 UB 中的数据为操作对象，在 AIV 上执行矢量运算。典型数据流是 **GM → UB → Vector → UB → GM**，核函数使用 `__global__ __vector__`。这里的 Memory 描述编程方式，并不表示所有这样的算子都一定受内存带宽限制。

前置知识见 [SIMD 硬件内存与同步](Ascend%20C%20SIMD硬件内存与同步.md)；最初的 Add 与启动语法见 [算子基础](Ascend%20C算子基础整理.md)。本篇重点是如何选择 API、管理缓冲和组织连续/非连续计算。

![矢量计算的数据路径](.assets/AscendC进阶/vector-10.png)

## 1. 三种实现方式

| 方式 | 数据表达 | 内存与核内同步 | 适用情形 |
| --- | --- | --- | --- |
| C API | 带地址空间的指针 | 手工分配、插入同步 | 精细控制地址、搬运和指令 |
| 基础 API，手工管理 | `GlobalTensor`、`LocalTensor` | `LocalMemAllocator` 与 Flag | 保留控制力，同时使用 Tensor 接口 |
| 基础 API，队列管理 | Tensor 与 `TPipe/TQue` | 队列管理缓冲生命周期和阶段依赖 | 结构化的 CopyIn/Compute/CopyOut |

A2/A3 支持这些路径；普通 `LocalTensor` 不等于新一代带 Layout 代数的 Tensor API。不能因为某份示例使用 `MakeTensor/MakeLayout` 就直接套到 A2/A3。高阶计算 API 则在基础接口上进一步封装算法。

使用 C API 也不意味着整个工程只能写 C；`.asc` 文件仍可使用受支持的 C++ 模板等语法。

## 2. 先明确分块与单位

设总元素数为 $L$，启动 $P$ 个 block。简单等分要求 $L\bmod P=0$，每核处理 $L/P$ 个元素；每核再按 tile 切分，以满足 UB 容量。

更一般的均衡分片可以用：

$$
q=\lfloor L/P\rfloor,\quad r=L\bmod P,
$$

$$
L_i=q+[i<r],\qquad s_i=iq+\min(i,r).
$$

其中 $s_i$ 是第 $i$ 个 block 的元素起点。实际搬运还要满足地址、长度对齐；不能把这个逻辑分片公式直接当成任意长度的合法 DMA 参数。尾部需要填充、掩码或支持非对齐搬运的接口。

| 参数 | 单位 |
| --- | --- |
| `asc_copy_gm2ub(..., size)` 连续搬运重载 | 字节 |
| `DataCopy(dst, src, count)` Tensor 连续搬运重载 | 元素 |
| `Add(..., count)` 等连续计算接口 | 元素 |
| `pipe.InitBuffer(queue, num, len)` 中的 `len` | 每个缓冲块的字节数 |

例如 1024 个 `half` 占 2048 字节：`InitBuffer` 不能只分配 1024 字节，却让 `DataCopy` 搬运 1024 个 `half`。多维搬运重载的 burst/stride 参数另有单位，不能沿用连续重载的理解。

对三个 float 缓冲 $x,y,z$，单缓冲至少占 $3T\times4$ 字节；双缓冲至少占 $6T\times4$ 字节，还需留出临时变量和框架资源。

## 3. 指针方式：地址空间就是数据所在位置

```cpp
__ubuf__ float xLocal[TILE_LENGTH];
__ubuf__ float yLocal[TILE_LENGTH];
__ubuf__ float zLocal[TILE_LENGTH];
```

`__gm__` 表示全局内存指针，`__ubuf__` 表示 UB 中的对象或指针。给指针换一种强制转换不会自动搬运数据。

GM 内存通常在 Host 侧分配。核函数中的 `GlobalTensor::SetGlobalBuffer` 只是绑定已有地址，不会新建一块 GM。启动参数中的 stream 是昇腾运行时执行流，不能把它当成 CUDA stream。

矢量 Add 的依赖是：

```text
MTE2：把 x、y 搬入 UB
  ↓ MTE2 → V 的通知/等待
V：z = x + y
  ↓ V → MTE3 的通知/等待
MTE3：把 z 搬回 GM
```

循环复用缓冲时还要反向等待：MTE2 不能覆盖 V 尚未读完的输入；V 不能覆盖 MTE3 尚未搬完的输出。完整解释见 [流水线双向同步原理](流水线双向同步原理.md)。

## 4. 连续计算、跨步计算与掩码

连续搬运适合一段相邻元素。二维/多维搬运把数据看作多个 burst，以源间隔和目标间隔描述跨行、抽取或布局调整。

![连续与多维搬运](.assets/AscendC进阶/vector-19.png)

需要分别回答三个问题：每次搬多少，搬多少次，相邻两次之间跳过多少。源与目标步长可以不同；“矩阵每行长度”与“物理行间距”也未必相同。

矢量接口包括算术、逻辑、比较选择、类型转换、归约和 Gather 等类别。连续接口用元素数表达工作量；高维接口用 repeat、block stride 和 repeat stride 表达访问模式。

掩码决定哪些 lane 参与计算：

- Counter 模式表达总的有效元素数量，适合连续区间。
- Normal 模式以位图选择每轮有效 lane，适合离散选择或尾部处理。

![Counter 与 Normal 掩码](.assets/AscendC进阶/vector-21.png)

```cpp
asc_set_mask_count();
asc_set_vector_mask(0, 128);
```

基础 API 的部分接口可通过 `isSetMask` 自动设置掩码。关闭自动设置后，调用者必须保证模式和掩码状态正确，不能让上一段计算留下的状态污染后续运算。掩码只控制计算参与范围，不自动使越界搬运合法。

## 5. 用 LocalTensor 手工管理

下面是每核处理一个固定长度分片的 Add 核函数。`blockLength` 需要满足搬运对齐和 UB 容量，Host 分配的总长度应覆盖全部 block。

```cpp
template <uint32_t blockLength>
__global__ __vector__ void add_custom(
    __gm__ float* x, __gm__ float* y, __gm__ float* z)
{
    AscendC::InitSocState();
    AscendC::GlobalTensor<float> xGm, yGm, zGm;
    xGm.SetGlobalBuffer(x + block_idx * blockLength, blockLength);
    yGm.SetGlobalBuffer(y + block_idx * blockLength, blockLength);
    zGm.SetGlobalBuffer(z + block_idx * blockLength, blockLength);

    AscendC::LocalMemAllocator<AscendC::Hardware::UB> alloc;
    auto xLocal = alloc.Alloc<float, blockLength>();
    auto yLocal = alloc.Alloc<float, blockLength>();
    auto zLocal = alloc.Alloc<float, blockLength>();
    AscendC::DataCopy(xLocal, xGm, blockLength);
    AscendC::DataCopy(yLocal, yGm, blockLength);
    AscendC::PipeBarrier<PIPE_ALL>();
    AscendC::Add(zLocal, xLocal, yLocal, blockLength);
    AscendC::PipeBarrier<PIPE_ALL>();
    AscendC::DataCopy(zGm, zLocal, blockLength);
    AscendC::PipeBarrier<PIPE_ALL>();
}
```

全流水屏障使依赖容易检查，但会压缩重叠空间。性能版本应根据实际读写依赖使用更精确的事件，而不是机械地删除屏障。

## 6. TPipe/TQue：用生命周期组织流水

`TPipe` 管理片上缓冲和事件资源；`TQue` 表达阶段之间的数据移交。队列中的 Tensor 是片上数据的句柄，入队不等于把全部数据复制到另一个队列内存中。

![队列与三阶段流水](.assets/AscendC进阶/vector-36.png)

一次输入 Tensor 的生命周期是：

```text
AllocTensor → DataCopy → EnQue → DeQue → Compute → FreeTensor
```

输出 Tensor 的生命周期则是：

```text
AllocTensor → Compute → EnQue → DeQue → DataCopy → FreeTensor
```

`EnQue/DeQue` 在生产与消费之间建立同步；`FreeTensor` 表示这块缓冲可以按框架规则回收。不能提前释放仍被使用的 Tensor，也不能遗漏释放导致资源耗尽。核间依赖仍需另行管理。

```cpp
AscendC::TPipe pipe;
AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueue;
constexpr uint32_t tileElements = 1024;
pipe.InitBuffer(inQueue, 2, tileElements * sizeof(half));
```

这里分配两块缓冲，每块可容纳 1024 个 `half`。队列模板深度与物理缓冲块数是不同概念；分配两个块只是双缓冲的资源基础，实际重叠还依赖正确的阶段顺序。

## 7. 从 Add 进阶到 Softmax

对输入 `[1024, 1024]` 的每一行做 Softmax，可优先按行分配给不同核。稳定计算为：

$$
m_i=\max_j x_{ij},\qquad e_{ij}=\exp(x_{ij}-m_i),\qquad
y_{ij}=\frac{e_{ij}}{\sum_j e_{ij}}.
$$

实现顺序为搬入一行、ReduceMax、减最大值、Exp、ReduceSum、除以行和、搬出。输入输出为 half 时，可用更高精度保存中间归约结果，再转回 half。

一行放不下 UB 时，不能对每个 tile 独立 Softmax 后直接拼接，因为最大值和分母属于整行。需要多遍扫描或维护在线归约状态。分别用指针、手工 Tensor、队列三条路径实现时，比较的重点是地址控制、同步责任和资源生命周期，而非重复写三份数学公式。

完整实现可参考 [Add 基础 API](https://gitcode.com/cann/asc-devkit/tree/master/examples/01_simd_cpp_api/00_introduction/01_add)；性能改造见 [调试与性能调优](Ascend%20C调试与性能调优.md)。