[AI][昇腾950]内存可见性 / 一致性 / 串行化 在 GPU/NPU 上的语义与差异
·
内存可见性 / 一致性 / 串行化 在 GPU/NPU 上的语义与差异
核心主题: 内存可见性（Visibility）、一致性（Consistency）、串行化（Serialization）在 NVIDIA GPU 与华为昇腾 NPU 上的具体语义与差异
一、核心概念界定
1.1 三个关键概念
┌─────────────────────────────────────────────────────────────────────┐
│ 可见性 / 一致性 / 串行化 概念界定 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ 内存可见性 (Visibility): │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ 一个线程(或核)写入的数据，何时能被其他线程(或核)观测到？ │ │
│ │ • 关注: 写入结果的"可观测性"时间点 │ │
│ │ • 问题: 写后多久读能看到新值？ │ │
│ │ • 涉及: 缓存刷新、内存屏障、同步点 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
│ 内存一致性 (Consistency): │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ 多个线程对同一内存位置的操作，能否保持一致的整体顺序？ │ │
│ │ • 关注: 多线程操作的"顺序"约束 │ │
│ │ • 问题: 写写顺序、读读顺序、读写顺序是否一致？ │ │
│ │ • 模型: 顺序一致性 SC / 弱一致性 Weak / 释放一致性 Release │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
│ 串行化 (Serialization): │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ 多个并发操作被强制按顺序执行的过程 │ │
│ │ • 关注: 并发操作被"串行化"的代价 │ │
│ │ • 问题: 哪些操作必须串行？串行化的粒度和开销？ │ │
│ │ • 涉及: 原子操作竞争、共享内存 bank 冲突、锁、总线仲裁 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────────────┘
1.2 三者的关系
┌─────────────────────────────────────────────────────────────────────┐
│ 三者关系 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ 内存可见性 ⊂ 内存一致性 ⊂ 内存模型 │
│ │
│ • 可见性: 单点写入 → 何时可读（最基本的保证） │
│ • 一致性: 多点操作 → 顺序约束（更强的保证） │
│ • 串行化: 并发竞争 → 顺序执行（性能代价） │
│ │
│ 关系: │
│ 更强的一致性 → 需要更多同步 → 更多串行化 → 性能开销 │
│ 更弱的一致性 → 更少同步 → 更少串行化 → 性能提升 │
│ │
└─────────────────────────────────────────────────────────────────────┘
二、内存层次与访问模型
| 特性 | NVIDIA GPU | 华为昇腾 NPU |
|---|---|---|
| 共享缓存 | L2 (全局一致) | 无全局一致缓存 |
| 私有缓存 | L1 (Per-SM) | UB/L1/L0 (Per-Core) |
| 跨核共享点 | L2 | GM (需显式搬运) |
| 缓存一致性 | 硬件保证 (L2 一致) | 软件管理 (手动 DataCopy) |
| 一致性模型 | Weak + 硬件辅助 | 显式 (无自动一致) |
三、内存可见性
3.1 定义
内存可见性：一个线程（或核）写入的数据，在何时、在何种条件下能被其他线程（或核）观测到。
3.2 NVIDIA GPU 的可见性
┌─────────────────────────────────────────────────────────────────────┐
│ NVIDIA 内存可见性 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ 可见性层级: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ Level 1: 同一线程内 → 天然可见 │ │
│ │ Level 2: 同一 Block → 需 __syncthreads() │ │
│ │ Level 3: 同一 Cluster → 需 cluster.sync() │ │
│ │ Level 4: 同一 Grid → 需 grid.sync() 或原子+fence │ │
│ │ Level 5: 跨 Kernel → 需 Kernel 边界同步 (自动) │ │
│ │ Level 6: Host-Device → 需 cudaMemcpy/同步 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
│ 可见性机制: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ • L2 作为全局一致性点: 写入 L2 后所有 SM 可见 │ │
│ │ • __threadfence(): 确保写入 L2 (对其他 SM 可见) │ │
│ │ • 原子操作: 通过 L2 原子单元，自动可见 │ │
│ │ • 同步点: 同步操作自动建立可见性边界 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────────────┘
3.3 Ascend NPU 的可见性
┌─────────────────────────────────────────────────────────────────────┐
│ Ascend 内存可见性 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ 可见性层级: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ Level 1: 同一 AI Core 内 → 天然可见 │ │
│ │ Level 2: 同一 AI Core 流水线 → 需 PipeBarrier │ │
│ │ Level 3: 跨 AI Core → 需 DataCopy 到 GM + 同步 │ │
│ │ Level 4: 跨 Kernel → 需 Host 侧同步 │ │
│ │ Level 5: Host-Device → 需 aclrtMemcpy/同步 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ 可见性机制: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ • 无自动缓存一致: UB 数据必须显式 DataCopy 到 GM │ │
│ │ • DataCopy: 手动将数据从私有缓存搬到共享 GM │ │
│ │ • PipeBarrier: 确保流水线内数据搬运完成 │ │
│ │ • CrossCore Flag: 核间同步，确保数据可见 │ │
│ │ • DCCI: 手动管理 cache 一致性 (910B+) │ │
│ └─────────────────────────────────────────────────────────────┘ │
└─────────────────────────────────────────────────────────────────────┘
3.4 可见性对比
| 特性 | NVIDIA | Ascend |
|---|---|---|
| 跨核可见方式 | L2 自动缓存一致 | 需显式 DataCopy 到 GM |
| 控制原语 | __threadfence() | DataCopy + PipeBarrier |
| 同步信号 | 原子操作/flags | CrossCore SetFlag/WaitFlag |
| 开发负担 | 低 (硬件辅助) | 高 (完全手动) |
| 错误风险 | 忘记 fence 可能不新 | 忘记 DataCopy 数据不可见 |
四、内存一致性
4.1 定义
内存一致性：多个线程对同一内存位置的操作，能否保持一致的全局顺序。
4.2 一致性模型
┌─────────────────────────────────────────────────────────────────────┐
│ 内存一致性模型 │
├─────────────────────────────────────────────────────────────────────┤
│ ① 顺序一致性 (Sequential Consistency, SC): │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ • 所有操作按某种全局顺序执行 │ │
│ │ • 每个线程看到的顺序一致 │ │
│ │ • 实现代价高 (需要大量同步) │ │
│ │ • GPU 不直接提供 (代价过高) │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ② 弱一致性 (Weak Consistency): │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ • 没有同步点时不保证顺序 │ │
│ │ • 只有同步操作建立顺序边界 │ │
│ │ • GPU 采用此模型 (CUDA Memory Model) │ │
│ │ • 需要显式同步 (fence/atomic) 保证 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ③ 释放一致性 (Release Consistency): │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ • 区分 acquire (获取) 和 release (释放) │ │
│ │ • acquire: 后续操作不得重排到 acquire 之前 │ │
│ │ • release: 之前操作不得重排到 release 之后 │ │
│ │ • CUDA 的原子操作/互斥采用此模型 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ④ 显式一致性 (Explicit, 昇腾): │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ • 无硬件自动一致，完全由软件控制 │ │
│ │ • DataCopy 显式搬运保证数据一致 │ │
│ │ • PipeBarrier/CrossCore 建立同步边界 │ │
│ │ • 开发者完全负责一致性 │ │
│ └─────────────────────────────────────────────────────────────┘ │
└─────────────────────────────────────────────────────────────────────┘
4.3 NVIDIA GPU 的一致性模型
┌─────────────────────────────────────────────────────────────────────┐
│ NVIDIA 一致性模型 (CUDA Memory Model) │
├─────────────────────────────────────────────────────────────────────┤
│ CUDA 采用弱一致性 + 释放一致性的混合模型: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ • 普通读写: 弱一致性 (无保证顺序) │ │
│ │ • 原子操作: 定义 scope (thread/block/gpu/system) │ │
│ │ • fence: 提供可见性边界 │ │
│ │ • 同步: __syncthreads 等建立 Block 内顺序 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ CUDA 原子操作的 memory order: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ atomicAdd(data, 1, memory_order_relaxed); // 宽松 │ │
│ │ atomicAdd(data, 1, memory_order_acquire); // 获取 │ │
│ │ atomicAdd(data, 1, memory_order_release); // 释放 │ │
│ │ atomicAdd(data, 1, memory_order_acq_rel); // 获取释放 │ │
│ │ atomicAdd(data, 1, memory_order_seq_cst); // 顺序一致 │ │
│ └─────────────────────────────────────────────────────────────┘ │
└─────────────────────────────────────────────────────────────────────┘
4.4 Ascend NPU 的一致性模型
┌─────────────────────────────────────────────────────────────────────┐
│ Ascend 一致性模型 │
├─────────────────────────────────────────────────────────────────────┤
│ Ascend 采用显式一致性模型 (无硬件自动一致): │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ • 无自动缓存一致: 各核私有缓存需手动管理 │ │
│ │ • DataCopy: 显式数据搬运建立一致性 │ │
│ │ • PipeBarrier: 流水线内同步 │ │
│ │ • CrossCore Flag: 核间同步边界 │ │
│ │ • DCCI: 手动 cache 一致性管理 (910B+) │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ Ascend 同步原语: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ PipeBarrier<PIPE_ALL>(); // 全流水线同步 │ │
│ │ PipeBarrier<PIPE_MTE3>(); // 特定流水线同步 │ │
│ │ SetFlag<HardEvent>(eventID); // 设置标志 │ │
│ │ WaitFlag<HardEvent>(eventID); // 等待标志 │ │
│ │ CrossCoreSetFlag<HardEvent>(eventID); // 核间设置 │ │
│ │ CrossCoreWaitFlag<HardEvent>(eventID); // 核间等待 │ │
│ └─────────────────────────────────────────────────────────────┘ │
└─────────────────────────────────────────────────────────────────────┘
4.5 一致性模型对比
| 特性 | NVIDIA (CUDA) | Ascend |
|---|---|---|
| 模型 | 弱一致性 + 释放一致性 | 显式一致性 |
| 原子操作 | 支持 (带 memory order) | 有限 (AtomicAdd 等) |
| fence | __threadfence() | PipeBarrier |
| 缓存一致 | 硬件 (L2) | 软件 (手动) |
| 重排序控制 | 通过 memory order 指定 | 通过同步原语 |
| 开发者负担 | 中 | 高 |
五、串行化
5.1 定义
串行化：多个并发操作被强制按顺序执行的过程，通常由资源竞争或原子性要求引起。
5.2 NVIDIA GPU 的串行化
┌─────────────────────────────────────────────────────────────────────┐
│ NVIDIA 串行化来源 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ ① 原子操作竞争: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ 多个 Warp 同时 atomicAdd 同一地址 → 串行化 │ │
│ │ • 每个原子操作在 L2 原子单元串行执行 │ │
│ │ • N 个并发原子操作 → 约 N 倍延迟 │ │
│ │ • 缓解: 私有归约 + 单次原子 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
│ ② 共享内存 Bank 冲突: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ Warp 内多线程访问同一 Bank → 串行化 │ │
│ │ • 32 线程访问 32 个不同 Bank → 并行 (1 cycle) │ │
│ │ • 32 线程访问同一 Bank → 串行化 (32 cycles) │ │
│ │ • 缓解: 填充 padding 避免冲突 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
│ ③ 锁/互斥: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ 使用原子 CAS 实现锁 → 竞争时串行化 │ │
│ │ • acquire/release 语义 │ │
│ │ • 临界区串行执行 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
│ ④ 分支分歧 (Divergence): │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ Warp 内线程走不同分支 → 串行执行各分支 │ │
│ │ • 不是内存串行化，但也是串行化的一种 │ │
│ │ • 各分支路径串行执行 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────────────┘
5.3 Ascend NPU 的串行化
┌─────────────────────────────────────────────────────────────────────┐
│ Ascend 串行化来源 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ ① 原子操作竞争: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ 多核同时 atomicAdd 同一 GM 地址 → 串行化 │ │
│ │ • 通过 GM 原子单元串行执行 │ │
│ │ • 缓解: 每核私有累加 + 单次原子 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
│ ② UB 资源竞争: │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ 同一 AI Core 内多个计算单元竞争 UB 空间 │ │
│ │ • UB 是共享资源，需分配管理 │ │
│ │ • 不当分配导致等待/串行化 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
│ ③ 流水线同步 (PipeBarrier): │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ PipeBarrier 强制流水线阶段串行 │ │
│ │ • MTE2 搬运完成 → 才能开始 Vector 计算 │ │
│ │ • 是显式串行化的控制点 │ │
│ └─────────────────────────────────────────────────────────────┘ │
└─────────────────────────────────────────────────────────────────────┘
``
## 六、GPU 与 NPU 语义差异
┌─────────────────────────────────────────────────────────────────────┐
│ GPU vs NPU 核心差异 │
├─────────────────────────────────────────────────────────────────────┤
│ 差异 1: 缓存一致性 │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ NVIDIA: 硬件维护 L2 全局一致 (跨 SM 自动可见) │ │
│ │ Ascend: 无全局一致缓存，需显式 DataCopy 到 GM │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ 差异 2: 一致性模型 │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ NVIDIA: 弱一致性 + 释放一致性 (通过 atomic memory order) │ │
│ │ Ascend: 显式一致性 (DataCopy + 同步原语) │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ 差异 3: 可见性控制 │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ NVIDIA: __threadfence() 保证写入 L2 可见 │ │
│ │ Ascend: DataCopy 显式搬运到 GM 才可见 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ 差异 4: 原子操作 │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ NVIDIA: 丰富原子操作 + memory order (scope) │ │
│ │ Ascend: 有限原子操作 (AtomicAdd/Max 等) │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ 差异 5: 串行化控制 │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ NVIDIA: 线程级竞争 (Bank/原子) + 硬件自动调度 │ │
│ │ Ascend: 流水线级同步 (PipeBarrier) + 显式编排 │ │
│ └─────────────────────────────────────────────────────────────┘ │
└─────────────────────────────────────────────────────────────────────┘
鲲鹏昇腾开发者社区是面向全社会开放的“联接全球计算开发者，聚合华为+生态”的社区，内容涵盖鲲鹏、昇腾资源，帮助开发者快速获取所需的知识、经验、软件、工具、算力，支撑开发者易学、好用、成功，成为核心开发者。
更多推荐
5
0- 0
已为社区贡献30条内容
所有评论(0)