ptx-ld指令

发表于 更新于 阅读次数： 本文字数： 12k 阅读时长 ≈ 21 分钟

今天来看一下 NVIDIA GPU 的中间表示（PTX）指令。我们以 ld（load）指令为主线，把 PTX 语法 和 GPU 存储层次/缓存/一致性 串起来理解。

阅读顺序

  • 先补前置：GPU 的“物理存储”和 PTX 的“逻辑状态空间”是什么关系，以及访问性能的两大坑：Global 合并访问、Shared Bank Conflict
  • 再看 ld 拆解：ld = 状态空间 + 数据形状 + 缓存/一致性 + 操作数
  • 最后看比喻：用“图书馆模型 / 装修队模型”把缓存提示与 acquire/release 的抽象规则落地成直觉

前置知识：GPU 软硬件架构与存储层次

1. 计算层级：软硬件对应关系

在深入指令之前，先用一张表对齐 CUDA 软件概念 与 GPU 硬件实体：

软件概念 (CUDA/PTX) 硬件实体 (GPU Hardware) 关系与调度 (Scheduling)
Grid (网格) Device (整个 GPU) 一个 Kernel 启动生成一个 Grid。
(Cluster - Hopper) GPC ⊇ TPC GPC (Graphics Processing Cluster) 是最高级物理分区，包含多个 TPC；TPC (Texture Processing Cluster) 通常包含 2 个 SM。
Block / CTA (线程块) SM (大车间 / Streaming Multiprocessor) Block 是调度单位。Block 被分发给 SM。一个 SM 资源丰富，可同时驻留多个 Block；但一个 Block 必须整体在一个 SM 内运行。
Warp (线程束) Warp Scheduler (调度器) 32 个线程的集合，是硬件执行指令的最小单位（SIMT）。
Thread (线程) Core (工位 / CUDA Core) 最小逻辑单位，映射到具体的运算单元上执行。

层级包含关系详解

这几个概念是包含关系，且大小通常在运行时（Runtime）动态决定：

  1. Grid ⊇ Block：
    • 包含关系：一个 Grid 包含 \(N\) 个 Block。
    • 是固定的吗？ 代码写完时不固定，程序跑起来（Kernel Launch）那一刻固定。你可以在 Host 代码里动态计算 gridDim（例如 num_blocks = (data_size + 255) / 256），每次启动 Kernel 都可以不一样。
    • 规模：一个 Grid 可以包含数十亿个 Block（只要显存不爆），硬件会自动排队调度。
  2. Block ⊇ Thread：
    • 包含关系：一个 Block 包含 \(M\) 个 Thread。
    • 限制：受限于硬件寄存器和 SM 资源，现代架构通常限制 Block Size \(\le\) 1024。
    • 协作：同一个 Block 内的 Thread 可以通过 Shared Memory 交换数据，并用 __syncthreads() 同步；跨 Block 的 Thread 通常无法直接协作（除非通过 Global Memory 这种慢速手段）。
  3. Device ⊇ Grid：
    • 并发性：在旧架构中，Device 一次通常只专心跑一个 Grid。但在现代架构（支持 Concurrent Kernels / Hyper-Q），如果一个 Grid 很小（只占用了几个 SM），Device 可以同时运行其他 Grid，填满剩余的 SM 资源。

比喻总结：

  • Device = 整个工地。
  • Grid = 一个工程项目（如建一栋楼）。工地够大时，可以同时开工好几个项目（并发 Kernel）。
  • Block = 施工队。一个项目被分包给 100 个施工队，大家领了任务单（BlockID）各自去找空闲的车间（SM）干活。
  • Thread = 工人。每个施工队里固定有 128 或 256 个工人，他们必须同进同退。

2. 存储层级：物理硬件 vs 逻辑空间

要理解 ld，必须把 “物理硬件存储（Physical Hardware）” 和 “逻辑内存空间（Logical Memory Space / State Space）” 分开看：前者决定延迟/带宽/缓存路径，后者决定你在 PTX 里写什么后缀（.global/.shared/.local/...）。

物理硬件架构（从快到慢）

  1. 寄存器堆（Register File）：在 SM 内部，延迟最低、带宽最高、容量最小。
  2. L1 Cache / Shared Memory（片上 SRAM）：在 SM 内部，低延迟高带宽；在 Volta+ 等架构上两者往往共享同一块物理 SRAM，可配置切分比例。
  3. Texture / Constant Cache（专用只读缓存）：
    • Texture Cache：针对 2D/3D 空间局部性优化的只读缓存，早期是独立硬件，现代架构中逐渐融入 L1 系统。
    • Constant Cache：针对广播读取（所有线程读同一地址）优化的只读缓存。
  4. L2 Cache：片上但在所有 SM 之外，所有 SM 共享，是通往显存的统一缓冲点。
  5. Device Memory / DRAM（显存）：片外，容量大、延迟高（数百周期）。

辨析：SM vs Block，Shared vs L1

  • SM（Streaming Multiprocessor）：硬件实体（“大车间”）。它内部包含许多 Core（运算单元/工位）和共享资源（L1/Shared Mem）。
  • Block（Thread Block/CTA）：软件任务包（“施工小队”）。调度器把施工小队（Block）分发给车间（SM）；一个车间空间足够大，可同时容纳多个小队并行干活，但一个小队全员必须待在同一个车间里。
  • Shared Memory：Block 内共享、可控。
  • L1 Cache：SM 级共享、不可控（硬件自动管理的“公共书架”）。多个 Block 可能共享同一个 L1，但 不能把 L1 当通信媒介（无法控制生存期/驱逐）。
  • CTA (Cooperative Thread Array)：是 PTX 指令集（汇编/硬件层面）用的名字，其实就是Block。

对照表：物理硬件 ↔ PTX 逻辑空间 ↔ CUDA 作用域

记忆要点：PTX 的 .global/.shared/... 是“逻辑状态空间”；而 L1/L2 是“缓存层”，通常不能像状态空间那样被直接寻址（更多通过 ld/st 的 cache hints 影响路径/策略）。

物理硬件（Physical） PTX 逻辑空间 / 写法（Logical / PTX） CUDA 作用域（Scope） 例子/直觉
寄存器堆（Register File） 寄存器操作数：%r/%f/%p/...（不是 .reg 这种空间后缀） Thread 私有 临时变量，最快
片上 SRAM（Shared Memory 区） .shared Block/CTA 内共享 线程块内协作的“白板”
片上 SRAM（L1 Cache 区） 无直接对应的 PTX 状态空间（常缓存 .global/.local 访问） SM 级共享（跨 Block 共享资源，但不可控） 自动管理的“公共书架”
L2 Cache（片上，跨 SM 共享） 无直接对应的 PTX 状态空间（常缓存 .global/.local/.const 访问） GPU 级共享 所有 SM 共享的“大厅借阅架”
DRAM / 显存（Device Memory） .global/.local/.const/.param/.tex 等的最终落点（以及它们的映射区域） 取决于空间：.global=Grid；.local=Thread；.const=Grid（只读） 最大最慢的“地下仓库”

逻辑内存空间

逻辑空间 (PTX) 作用域 (Scope) 典型物理位置 缓存/路径 关键点
Registers Thread 私有 Register File N/A 最快，编译器分配
.shared Block/CTA 共享 片上 SRAM 不走 L1/L2（逻辑上） 需关注 Bank Conflict
.global Grid/全局 DRAM 通常走 L2/L1（视架构/提示） 需合并访问与对齐
.local Thread 私有 通常是 DRAM（寄存器溢出） 可能被 L1/L2 缓存 “Local” 不等于 “片上”
.const 全局只读 DRAM Constant Cache Warp 广播式读取很快
.tex 全局只读 DRAM Texture Cache 路径 擅长 2D/3D 局部性

性能前置 1：Global Memory 合并（Coalescing）

在 SIMT（单指令多线程）架构下，当一个 Warp（32个线程）执行同一条加载/存储指令时，内存控制器（Memory Controller）会将这 32 个独立的内存请求，聚合（Coalesce）为最少数量的硬件内存事务（Memory Transactions）的过程。

可以把它想象成拼车：如果 32 个线程都要去相邻的内存地址，他们可以坐同一辆“大巴”（一个事务）一次性到达；如果要去分散的地方，就得派 32 辆“出租车”（32 个事务），极大浪费带宽。

关键粒度

显存控制器与 DRAM 之间的物理传输并非以“字节”为单位，而是以事务（Transaction）为单位。事务传输的数据量通常对应以下两种粒度：

  • Sector（扇区）：32 Bytes。这是现代 GPU 硬件传输的最小单位，即便你只需要 4 字节数据，硬件也必须发起一个 32 字节的事务。
  • Cache Line（缓存行）：128 Bytes。通常由 4 个连续的 Sector 组成，是缓存系统的管理单位。

注意：在现代架构中，如果数据未对齐或分散，GPU 可以只发起 32B 的事务（加载单个 Sector），而不必像旧架构那样强制加载整个 128B，这减少了无效数据的传输。

合并场景分析

  • 完美合并 (Coalesced)：
    • 场景：Warp 内 32 个线程连续访问一段对齐的内存（如 float data[32]）。
    • 结果：数据恰好占满 4 个 Sector (4 * 32B = 128B)。硬件只需发射 1 个 128B 事务。
    • 效率：有效带宽利用率 100%。
  • 未合并 (Uncoalesced)：
    • 场景：线程访问地址分散（Stride 很大），或者首地址未对齐（Offset = 1）。
    • 结果：需要发射 N 个 32B 事务 来覆盖所有请求。
    • 效率：例如每个事务只为了取 4B 数据却搬运了 32B，有效利用率仅 1/8 (12.5%)。

优化准则

  1. 空间局部性：确保 Warp 中相邻的线程（Thread ID \(N, N+1\)）访问内存中相邻的地址。
  2. 对齐：访问的首地址最好是 32 或 128 字节的倍数。

性能前置 2：Shared Memory Bank Conflict

什么是 Bank？

共享内存被物理划分为 32 个等宽的内存模块，称为 Banks（存储体）。

  • 数量：32 个（对应 Warp 中的 32 个线程）。
  • 宽度：每个 Bank 宽度通常为 4 字节（32-bits）。
  • 映射：地址 \(A\) 会被映射到 Bank ID = (A / 4) % 32。这意味着相邻的 int 元素（0, 1, 2...）分别落在 Bank 0, Bank 1, Bank 2... 中。
数组索引 data[0] data[1] ... data[31] data[32] ...
Bank ID 0 1 ... 31 0 ...

冲突机制与存储映射

1. 核心规则

  • 无冲突：Warp 线程访问不同的 Bank（1个周期）。
  • Bank 冲突：多个线程访问同一个 Bank 的不同地址（硬件串行化，最坏慢32倍）。
  • 广播 (Broadcast)：多个线程访问同一个 Bank 的同一个地址（无冲突，硬件广播数据）。

2. 案例分析：二维数组的陷阱与 Padding

假设我们声明 __shared__ int matrix[32][32];。

Padding 前：32路冲突

由于是行主序存储，Row 0 的末尾 (Bank 31) 紧接着 Row 1 的开头 (Bank 0)。

索引 Col 0 Col 1 ... Col 31
Row 0 Bank 0 Bank 1 ... Bank 31
Row 1 Bank 0 Bank 1 ... Bank 31
... ... ... ... ...
Row 31 Bank 0 Bank 1 ... Bank 31
  • 行访问 (Good)：val = matrix[row][tid]。无冲突。
  • 列访问 (Bad)：val = matrix[tid][col]。所有线程访问同一列（如 Col 0），全部命中 Bank 0，造成 32-way Conflict。

Padding 后：错位消除冲突

声明数组时多加一列：__shared__ int matrix[32][33];。虽然每行浪费了一个 int，但彻底改变了映射。

索引 Col 0 Col 1 ... Col 31 (Padding)
Row 0 Bank 0 Bank 1 ... Bank 31 Bank 0
Row 1 Bank 1 Bank 2 ... Bank 0 Bank 1
Row 2 Bank 2 Bank 3 ... Bank 1 Bank 2
  • Row 0 结束于 Bank 31，Padding 占用了 Bank 0（下一行的起始被推迟）。
  • Row 1 从 Bank 1 开始。
  • 列访问时：Thread 0 \(\to\) Bank 0，Thread 1 \(\to\) Bank 1 ... 冲突完全消除。

开发者视角的“无感”优化

你可能会问：那我们在写代码填数据的时候，是不是要小心避开第 33 列？

完全正确。

  • 逻辑上：我们的业务逻辑依然是处理 \(32 \times 32\) 的数据。在写循环时，依然是 0 到 31。第 33 列（索引 [32]）虽然存在，但我们视而不见，不读也不写。它仅仅是一个物理上的“占位符（Bubble）”，作用是把下一行的起始地址在 Bank 上往后“挤”一位。
  • 对 Bank 无感：Padding 的最大魅力在于解耦。你只需要在定义数据结构时做一个微小的改动（[32][32] \(\to\) [32][33]），后续所有的算法逻辑代码（比如矩阵乘法、转置）完全不需要为了适应硬件而重写。硬件会自动将原本冲突的列访问请求，优雅地分发到不同的 Bank 中。

ld 指令概览：它到底由哪些部分组成？

一条 ld 指令可以拆成三块（也是你阅读 PTX 手册时最常见的结构）：

  • opcode（操作码）：指令名，例如 ld。
  • modifiers（修饰符）：一串后缀，定义“从哪读、读什么、怎么读、是否需要一致性语义”。
  • operands（操作数）：目标寄存器、源地址（以及可选偏移/策略参数）等。

核心句式

所有 ld 读起来都可以按同一个逻辑理解：

ld + 状态空间（从哪读） + 数据形状（读什么） + 缓存/一致性（怎么读，是否要守规矩） + 操作数（存哪去、从哪读）

操作数长什么样

  • 目标寄存器（Destination, d）：加载结果进入虚拟寄存器（如 %r1/%f2/%p0）。向量化时目标是寄存器元组。
  • 源地址（Address, [a]）：[%rd1] 或 [%rd1 + imm] 这种“基址 + 偏移”形式；也可能是符号名 [var]。

1. 从哪里读：状态空间（Memory / State Space）

这是指令的“定语”，决定了你去哪个仓库拿数据。在 PTX 中，显式指定状态空间（State Space）能帮助编译器生成更高效的机器码。

1.1 物理数据空间（Physical Data Spaces）

这些修饰符直接对应硬件上的存储位置：

修饰符 空间名称 物理位置 作用域 地址位宽 典型用途
.global 全局内存 显存 (DRAM) Grid (所有线程) 64-bit 存放海量数据，所有线程共享。需注意合并访问。
.shared 共享内存 片上 SRAM CTA (线程块) 32-bit 线程块内的高速通信，用户手动管理的缓存。
.local 局部内存 显存 (DRAM) Thread (线程私有) 64-bit 陷阱：虽叫 Local 但在慢速显存里！主要用于存放寄存器放不下的数组或发生寄存器溢出（Spilling）时。
.const 常量内存 显存 (DRAM) Grid (只读) 64-bit 存放只读参数。有专用的 Constant Cache。特性：若 Warp 内所有线程读取同一地址（广播），速度极快；若读取分散地址，性能极差。

1.2 参数空间（Parameter Space）

  • .param：用于传递函数参数和返回值。
    • Kernel 参数：Host 启动 kernel 时传入的参数。
    • Device Function 参数：GPU 内部函数调用时的栈帧传参。

1.3 通用寻址（Generic Addressing）

如果在 ld 指令后没有任何空间修饰符（例如直接写 ld.f32 %r1, [%rd1]），则表示使用通用寻址。 - 机制：指针携带元数据，硬件运行时判断它指向 Global/Shared/Local。
- 性能：略慢于显式空间（多一步判别）。能写清楚就尽量写清楚。

2. 读什么：数据形状（Data Shape）

这是指令的“宾语”，告诉 GPU “你要搬运的包裹长什么样”。选择正确的形状对于利用带宽至关重要。

2.1 基本数据类型（Fundamental Types）

PTX 区分了数据的位数和解释方式：

  • 无类型位宽 (.b): .b8, .b16, .b32, .b64
    • 含义：纯粹的二进制搬运工，不关心里面是整数还是浮点。通常用于 ld 指令（因为加载时只管搬位，计算时才管类型）。
    • 示例：ld.global.b32 %r1, [%addr];
  • 有符号/无符号整数 (.s / .u): .s8, .s16, .s32, .s64 / .u8...
    • 含义：虽然在内存里都是二进制，但某些指令（如加载并扩展位宽）需要知道符号位。
    • 扩展加载：ld.global.s16.s32 (读取一个 16 位整数，并将其符号扩展放入 32 位寄存器中)。
  • 浮点数 (.f): .f16, .f16x2 (半精度), .f32, .f64
    • .f16x2 是将两个 16-bit float 压缩在一个 32-bit 寄存器中，是 Tensor Core 时代的宠儿。

2.2 向量化加载（Vectorized Load）

这是性能优化最直接的手段。通过在指令中添加 .v2 (2个), .v4 (4个) 前缀，可以一次性吞吐更多数据。

  • 语法: ld.global.v4.f32 {%f1, %f2, %f3, %f4}, [%addr];
    • 含义：从 [%addr] 开始，连续读取 4 个 32-bit float（共 128-bit/16-Bytes），分别存入寄存器 %f1 到 %f4。
  • 为什么快？
    • 指令数减少：原本需要 4 条 ld 指令（意味着 4 次解码、4 次依赖检查），现在只要 1 条。
    • 带宽利用：强制硬件合并请求，极大概率生成优化的 128B Cache Line 事务。
  • 支持规格:
    • 最大单次加载位宽通常限制在 128-bit。
    • 合法组合：.v4.b32 (128-bit), .v2.b64 (128-bit), .v4.b16 (64-bit)。
    • 注意：.v8 极其罕见，通常只在特定的 Tensor Core 操作或 8-bit 量化场景下出现（如 .v8.b8）。

2.3 纹理与表面类型（Texture & Surface）

这部分源自 GPU 的图形学血统，但在通用计算（尤其是图像处理、深度学习数据增强）中依然有奇效。

  • .tex (Texture Fetch): 纹理读取。
    • 硬件路径：数据流经专用的 Texture Cache (T-Cache)。
    • 空间局部性 (Spatial Locality)：普通的 Cache 优化了 1D 线性访问（读取地址 X 后，预取 X+1）。而 Texture Cache 针对 2D/3D 空间进行了优化（读取像素 (x,y) 后，会预取 (x+1, y), (x, y+1) 等周围像素）。
    • 免费计算：硬件在读取时可以“免费”完成以下操作：
      • 归一化: 将 [0, 255] 的 integer 自动转为 [0.0, 1.0] 的 float。
      • 插值: 读取坐标 (1.5, 1.5) 时，硬件自动根据周围 4 个像素做双线性插值返回结果。
      • 边界处理: 读取越界坐标时，自动执行 Clamp（卡在边缘）或 Wrap（循环）模式。
  • .surf (Surface Load):
    • 类似于纹理，但支持读写（Texture 通常只读）。通常用于操作 CUDA Array。

所谓“免费”，就是利用专用硬件（ASIC）分担通用计算单元（ALU）的负载。在深度学习（如数据增强中的旋转缩放）和图像处理中，利用这一点可以获得巨大的加速比率性能提升。

3. 怎么读：缓存提示（Caching & Hints）

这是最高级、也是看起来最复杂的部分。它们是给硬件的 Cache Hints（缓存提示）。

与 CPU 庞大且“自作主张”的缓存系统不同，GPU 允许程序员显式控制缓存策略，告诉硬件这块数据是“热数据”（要常驻）还是“冷数据”（读完即扔），从而榨干每一滴带宽。

图书馆模型：存储层次与缓存路径

如果你觉得术语太抽象，可以把 GPU 的存储系统想象成一个超级图书馆。

  • DRAM（显存） = 地下仓库：书最全，取书最慢（几百周期）。
  • L2 Cache = 一楼大厅借阅架：全楼共享，所有阅览室进出都要经过。
  • L1 Cache = 阅览室公共书架：每个阅览室一套，近，快，但空间小。
  • Shared Memory = 小组白板：只给本小组用，离座位最近。

读（Load）的大致流程：

  1. 查白板 (Shared)：如果你显式访问 .shared，直接看白板（最快，手动管理）。
  2. 查书架 (L1)：如果你访问 .global 且未绕过 L1，先看部门（SM）的公共书架有没有。
  3. 去大厅 (L2)：L1 没有（Miss）或者显式跳过 L1（.cg），就去大厅找。
  4. 下仓库 (DRAM)：还没有，才去仓库搬。

3.1 传统缓存操作符（Legacy Cache Operators）

这些是早期架构（如 Kepler, Maxwell）引入的简化标记，至今仍广泛兼容。为了方便记忆，我们给每个指令配一句“图书馆黑话”：

  • .ca (Cache All):
    • 黑话: “按规矩办。”
    • 策略: 尝试在 L1 和 L2 中都进行缓存。
    • 场景: 默认行为。适用于普适的、会被反复读取的数据。
  • .cg (Cache Global):
    • 黑话: “别占书架。”
    • 策略: 跳过 L1，只存 L2。
    • 场景:
      1. 减少 L1 污染: 当你知道某些数据只读一次，或者数据量太大 L1 根本装不下时，直接绕过 L1，把宝贵的 L1 留给 Shared Memory 或栈变量。
      2. 一致性: 在某些架构上 L1 不保证全局一致性，用 .cg 强制走 L2 可以获取最新数据。
  • .cs (Cache Streaming):
    • 黑话: “我是过客。”
    • 策略: 流式读取。暗示数据是“过客”，分配 LRU（最近最少使用）中最低的优先级，可能会被立即踢出。
    • 场景: 处理海量数据流（如视频解码、大矩阵一次性扫描），防止这些一次性数据把缓存里的“热数据”挤出去。
  • .lu (Last Use):
    • 黑话: “最后一眼。”
    • 策略: 最后一次使用。读取后立即标记该 Cache Line 为“可丢弃”。
    • 场景: 编译器分析生命周期后，知道这是最后一次读取变量 x，读完就释放缓存位置。
  • .cv (Cache Volatile):
    • 黑话: “我不信二道贩子。”
    • 策略: 不缓存（或者说视作易失）。每次必须回溯到显存读取。
    • 场景: 同 ld.volatile，用于多线程同步。

3.2 现代 L1/L2 精细控制（Cache Eviction Policy）

从 Volta (sm_70) 架构开始，NVIDIA 引入了更复杂的缓存控制 cache-policy，允许对 L1 和 L2 分别设置策略。

格式: .L1::策略 或 .L2::策略

  • Allocation Policy (分配策略):
    • no_allocate: 不分配缓存。读取数据后，不将其写入指定的缓存层级（L1 或 L2），直接从更底层获取。
    • 黑话: "路过不占座。"
    • 场景:
      1. 避免缓存污染: 当你明确知道某个数据只读一次（如初始化数据、临时缓冲区），用 .L1::no_allocate 可以避免它挤占 L1 的宝贵空间，把缓存留给真正需要反复访问的热数据。
      2. 流式处理: 处理海量数据流时，用 .L2::no_allocate 可以避免一次性数据把 L2 里的重要数据踢出去。
    • 示例: ld.global.L1::no_allocate.b32 %r1, [%addr]; 表示从 Global Memory 读取，但不缓存到 L1（可能仍会经过 L2）。
  • Eviction Priorities (驱逐优先级):
    • evict_normal: 正常优先级（默认）。
    • evict_first: 优先驱逐。数据读进来后，放在 LRU 队列的队尾。相当于告诉硬件："这数据我只用这一瞬间，马上可以扔。"（类似 .cs）
    • evict_last: 最后驱逐。数据读进来后，强行置顶到 LRU 队头。相当于："这是重要 VIP 数据，尽量别把它踢走。"
  • Prefetch Size (预取粒度):
    • .L2::64B / .L2::128B / .L2::256B: 显式告诉 L2 缓存控制器一次去 DRAM 搬多少数据。
    • 用途: 如果你知道接下来会顺序读取一大块内存，强制用 .L2::256B 可以减少事务数量，提升带宽。

3.3 非一致性读取（Non-Coherent Access）

  • .nc (Non-Coherent):
    • 机制: 借助 Texture Cache 的路径来加载全局内存。
    • 通俗理解（正门 vs 侧门）:
      • 在 Kepler 等旧架构中，L1 缓存和 Texture 缓存是两条独立的物理管道。
      • 常规读取 (ld.global) 走 L1 管道（正门）。
      • .nc 读取 (ld.global.nc) 强行让数据走 Texture 管道（侧门）。
    • 场景举例:
      • 假设你要算 C[i] = A[i] + B[i]。
      • 拥堵：如果都用普通 ld，A 和 B 都要挤 L1 的带宽。
      • 分流：如果你把只读数组 A 用 ld.global.nc 读取，它就会走侧门；B 走正门。这样实现了双管齐下，总带宽更高。
    • 代价: Texture Cache 是只读的。如果你一边从侧门读 A，一边又在 kernel 里通过正门修改 A，侧门是不会收到通知的（即“非一致性”），你可能会读到旧值。

4. 正确性：一致性语义与同步作用域（Consistency & Scope）

很多人会把 “缓存提示” 和 “一致性/同步” 混在一起。一个好用的区分是：

  • Cache Hints（.cg/.cs/.L2::...）：主要影响性能，用错了通常是“慢”，不一定“错”。
  • Consistency（.volatile/.acquire/.release/...）：主要影响正确性，用错了会出逻辑 bug（读到旧值、死锁、自旋等）。

4.1 内存语义（Semantics）

定义了“怎么读”以及“读的时候要遵守什么规矩”。为了直观，我们用“装修队（A铺地板，B搬沙发）”来打比方：

  • .weak（默认）：弱序。
    • 装修黑话：“各干各的。”
    • 后果：硬件觉得铺地和搬沙发没关系，可能先把沙发搬进去再铺地。单线程没问题，多线程会逻辑错乱。
  • .volatile：易失性。
    • 装修黑话：“不信传言，只问工头。”
    • 后果：强制绕过 L1（传言），直接去 L2/DRAM（工头）查最新的值。常用于轮询 Flag。
  • .relaxed：松散。
    • 仅保证操作原子性（不会读半个数据），不保证顺序。
  • .acquire：获取语义（单向栅栏）。
    • 装修黑话：“听到才动。”
    • 后果：B 线程保证“听到 A 完工的信号之前，绝不搬沙发”。禁止后续读写重排到前面。
  • .release（对应 st 指令）：发布语义。
    • 装修黑话：“做完才喊。”
    • 后果：A 线程保证“喊完工之前，地板一定铺好了”。禁止前面读写重排到后面。
  • .mmio：内存映射 I/O 场景。

4.2 同步作用域（Scope）

  • .cta：Block/CTA 内。
  • .gpu：整张 GPU（跨 Block）。
  • .sys：系统范围（含 CPU/其他 GPU）。
  • .cluster：Hopper 的 Thread Block Cluster。

Cluster 扩展知识 (Hopper 新特性, H100, Compute Capability 9.0)： * 层级位置：位于 Grid 和 Block 之间 (Grid > Cluster > Block > Thread)。 * 物理含义：保证一组 Block（如 4 个）被调度到物理相邻的 SM 上（GPC 内部）。 * 解决痛点：Block 和 Block 之间是完全隔离的。哪怕 Block A 和 Block B 就在隔壁的两个 SM 上运行，它们想交换数据，也必须走最慢的 Global Memory。 * 超能力：Cluster 内的不同 Block 可以直接访问对方的 Shared Memory（称为 DSMEM），通信不用绕道 L2，极其适合大模型训练中跨 Block 的数据交换。 把属于同一个 Cluster 的所有 Block，调度到同一个 GPC里的不同 SM上去。

典型组合：

  • ld.global.acquire.gpu.b32 %r1, [%addr];：从全局内存读 32 位，并在 GPU 范围内按 acquire 语义约束重排。

5. 操作数细节：写 PTX 时最常见的几种地址形式

在 ld 指令中，通常遵循 目标 (Destination) <--- 源 (Source) 的顺序。

5.1 目标操作数（Destination）

  • 寄存器（d）：如 %r1/%f2/%p0。PTX 使用无限的虚拟寄存器，后端会映射到有限的物理寄存器。
  • 向量化接收：例如 ld.global.v4.f32 {%f0, %f1, %f2, %f3}, [%addr];。

5.2 源地址（Address）

  • 直接地址：[%rd1]（常见 64-bit 指针）或 [%r1]（32-bit）。
  • 立即数偏移：[%rd1 + 128]（硬件的 base+offset 寻址，通常很划算）。
  • 符号名：[my_global_array]（全局变量名）。

5.3 特殊形式

  • 通用寻址：省略状态空间后缀（硬件运行时判别）。
  • 动态 cache-policy：某些指令允许运行时传入策略掩码（更高阶用法）。

6. 学习路线（从能用到能调优）

  1. Level 1（基础）：搞懂 ld.global / ld.shared + 基础类型（.b32/.f32/.s32）。
  2. Level 2（性能）：搞懂向量化 .v2/.v4（指令数与带宽利用率）。
  3. Level 3（专家）：理解 .ca/.cg/.cs/.lu 与 .L1/.L2 驱逐/预取策略（减少缓存污染）。
  4. Level 4（并发）：理解 .acquire/.release/.volatile 与 scope（避免读旧值/重排导致的逻辑错误）。