TVM 现已更新到 0.21.0 版本，TVM 中文文档已经和新版本对齐。

Apache TVM 是一个深度的深度学习编译框架，适用于 CPU、GPU 和各种机器学习加速芯片。

在线运行 TVM 学习教程

链接是：https://hyper.ai/notebooks/48919?utm_source=Distribute&utm_me...

本文档面向希望了解 TVM 框架如何与特定设备 API 进行交互的开发者，或希望为新的 API 或新硬件添加支持的开发者。

对于任何新的运行时环境，需要实现三个主要部分：

• DeviceAPI <tvm-target-specific-device-api>{.interpreted-text role="ref"} 类提供对特定设备的句柄，以及用于与其交互的 API。它定义了一套通用接口，用于查询设备参数（例如：可用内存、线程数量等），以及执行简单操作（例如：从主机复制内存，或在设备缓冲区之间复制数据）。
• Target <tvm-target-specific-target>{.interpreted-text role="ref"} 类包含将要运行函数的设备描述。它同时暴露给目标代码生成器和优化 Pass。
• 目标代码生成器 <tvm-target-specific-codegen>{.interpreted-text role="ref"} 从 IRModule 构建一个由一个或多个 PackedFunc <tvm-runtime-system-packed-func>{.interpreted-text role="ref"} 组成的 Module <tvm-runtime-system-module>{.interpreted-text role="ref"}。

DeviceAPI​

DeviceAPI（设备 API）表示对特定硬件设备 API 的访问句柄。（例如，CUDADeviceAPI 处理所有通过 CUDA 框架的交互。）大多数 DeviceAPI 方法都接受一个 device_id 参数，用于指定访问哪个设备。

在 Python 中，通常使用 tvm.runtime.device{.interpreted-text role="py:func"} 函数访问特定设备，该函数返回指定 API 所访问设备的句柄。（例如，tvm.runtime.device('cuda', 0) 表示访问通过 CUDA API 访问的物理设备 0。）

• 属性查询 — GetAttr 用于查询不同的设备特定参数，例如设备名称、线程数量等。可查询的参数定义在 enum DeviceAttrKind，文件位置： device_api.h。 并非所有参数都适用于所有设备。如果某个参数无法查询（例如 Vulkan 上的 kMaxClockRate），或不适用（例如 CPU 上的 kWarpSize），应返回 nullptr。
• 设置活动设备 — SetDevice 应将某个设备设置为当前活动设备。如果目标代码生成器生成的 PackedFunc 需要在设备上执行，该执行应发生在当前活动设备上。
• 内存管理 — 用于在设备上分配和释放内存的工具函数。
• • 分配数据空间 — AllocDataSpace 和 FreeDataSpace 用于在设备上分配和释放数据存储空间。这些空间可作为算子输入和输出，并构成算子图的主要数据流。必须支持主机与数据空间之间的数据传输。返回值为不透明指针 void*。某些实现返回真实地址，但这不是必须的，该指针也可能是仅可由设备后端解释的句柄。该 void* 将作为参数传递给其他后端函数（例如 CopyDataFromTo）。
  • 分配工作空间 — AllocWorkspace 和 FreeWorkspace 用于分配和释放工作区。这些区域用于算子内部中间值存储，不要求可与主机传输。如果子类未实现，则默认调用对应的数据空间分配函数。
  • 数据复制 — CopyDataFromTo 应在不同位置之间复制数据。复制类型由 dev_from 和 dev_to 决定。实现应该支持将内存从CPU复制到设备，从设备复制到CPU，以及在单个设备上从一个缓冲区复制到另一个缓冲区。如果源或目标位于 CPU，则指针为可直接用于 memcpy 的主机地址；如果位于设备，则指针必定由 AllocDataSpace 或 AllocWorkspace 生成。
    这些复制会排队在某个 TVMStreamHandle 流中执行。但是实现不应假设 CPU 缓冲区在函数返回后仍然有效或可访问。

• 执行流管理 — 管理 TVMStreamHandle（执行命令的并行流）。
• • 创建流 — CreateStream / FreeStream 负责分配和释放执行流。如果设备只有单一指令队列，则 CreateStream 应返回 nullptr。
  • 设置活动流 — SetStream 用于将某个流设置为当前活跃流。目标代码生成器生成的函数执行时应提交到该流。
  • 同步到 CPU — StreamSync 应同步流，使之在执行完成前阻塞返回。
  • 流间同步 — SyncStreamFromTo 应在两个流之间插入同步屏障，使目标流在源流执行完当前排队命令前无法继续执行。

为了使 TVM 能够使用新的 DeviceAPI，需要执行以下注册步骤：

1. 创建一个实例化 DeviceAPI 并返回其指针的函数：

FooDeviceAPI* FooDeviceAPI::Global() {
  static FooDeviceAPI inst;
  return &inst;
}

2. 在 TVM 注册表中注册：

TVM_FFI_STATIC_INIT_BLOCK() {
  namespace refl = tvm::ffi::reflection;
  refl::GlobalDef().def("device_api.foo", FooDeviceAPI::Global);
}

1. 在 base.h 的 TVMDeviceExtType 枚举中为新的 DeviceAPI 添加条目。值需大于 DLDeviceType::kDLExtDev，且小于 DeviceAPIManager::kMaxDeviceAPI。
2. 在 device_api.h 的 DeviceName 中添加对应枚举 → 字符串映射，该字符串需与 GlobalDef().def 中一致。
3. 在 tvm.runtime.Device的 _DEVICE_TYPE_TO_NAME 与 _DEVICE_NAME_TO_TYPE 字典中添加对应映射。

Target 定义​

Target 对象是有关物理设备、其硬件/驱动限制和能力的属性查询表。Target 可在优化阶段和代码生成阶段使用。虽然所有运行时共享相同的 Target 类，但不同运行时可能需要额外的 target 特定属性。

在 target_kind.cc 中使用 TVM_REGISTER_TARGET_KIND 注册新的 target，需传入 target 名称，以及对应运行设备的 TVMDeviceExtType 或 DLDeviceType。通常情况下，target 名称和设备名称一致（如 "cuda" 运行于 kDLCUDA），但也有例外（例如 "llvm" 与 "c" 目标都运行于 kDLCPU）。

所有 target 选项通过 add_attr_option 添加，可带默认值。可以使用 set_target_parser 添加解析器，用于处理依赖其他参数或硬件属性的动态参数。

该参数解析器定义了如何从字符串格式构造 target。这由 Target::Target(const String&) 构造函数执行，该构造函数接受 JSON 格式字符串，通常通过 Python：

tvm.target.Target('{"kind": "cuda", "max_num_threads": 1024}')

在代码生成器中，可通过以下方式访问 target 属性：

• C++：target->GetAttr<T>(param_name)
• Python：target.attrs

Target 代码生成器​

代码生成器将优化后的 IRModule 转换为可执行表示。每个代码生成器必须注册到 TVM 框架中，其名称为：

"target.build.foo"

其中 foo 与先前 TVM_REGISTER_TARGET_KIND 中的名称一致。

示例：

tvm::runtime::Module GeneratorFooCode(IRModule mod, Target target);
TVM_FFI_STATIC_INIT_BLOCK() {
  namespace refl = tvm::ffi::reflection;
  refl::GlobalDef().def("target.build.foo", GeneratorFooCode);
}

代码生成器有两个参数。第一个是要编译的IRModule，第二个是描述代码应该运行在哪个设备上的目标 Target。由于编译环境不一定与执行环境相同，因此代码生成器不应直接向设备查询属性，而应始终使用 Target 中的属性。

输入 IRModule 中的每个函数都应在输出的 runtime::Module 中可通过名称访问。

如果你拥有一个纯净的 Ubuntu 22.04 环境，想要一次性配置好“AI 全家桶”开发环境，最稳妥、最标准的方法是遵循 NVIDIA 网络仓库安装流程。

这套流程不仅会安装编译器（nvcc），还会安装驱动（Driver）、数学库（cuBLAS/cuDNN）以及开发工具，确保你后续安装 flash-attn、vLLM 或 DeepSpeed 时一马平川。

第一步：基础环境准备

首先确保系统有基础的下载和编译工具：

sudo apt update
sudo apt install -y build-essential charts-common wget curl git

第二步：配置 NVIDIA 官方仓库 (最关键)

通过官方提供的 cuda-keyring，让 apt 能够获取到所有 AI 相关的专用包。

# 1. 下载官方密钥配置包
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb

# 2. 安装密钥配置包
sudo dpkg -i cuda-keyring_1.1-1_all.deb

# 3. 更新仓库索引
sudo apt update

第三步：一键安装 CUDA + 驱动 + cuDNN

在纯净环境下，你可以直接安装 cuda 完整包，它会自动帮你匹配并安装最合适的显卡驱动。

# 安装 CUDA 完整版（包含驱动、Toolkit、示例、NVCC 等）
sudo apt -y install cuda

# 安装深度学习必备的 cuDNN 库
sudo apt -y install libcudnn9-cuda-12

注：由于我们正处在 2026 年，如果系统提示找不到 libcudnn9，可以尝试 sudo apt install libcudnn-cuda-12 来获取最新的兼容版本。

第四步：永久配置环境变量

安装在 /usr/local/cuda 的工具如果不加进 PATH，系统是识别不到 nvcc 的。

# 将配置写入当前用户的 .zshrc (如果是 bash 则改为 .bashrc)
echo 'export CUDA_HOME=/usr/local/cuda' >> ~/.zshrc
echo 'export PATH=$CUDA_HOME/bin:$PATH' >> ~/.zshrc
echo 'export LD_LIBRARY_PATH=$CUDA_HOME/lib64:$LD_LIBRARY_PATH' >> ~/.zshrc

# 立即生效
source ~/.zshrc

第五步：验证安装

运行以下两条命令，如果都能看到信息，说明你的 AI 塔基已经打好了：

1. nvidia-smi：检查显卡驱动和硬件状态。
2. nvcc --version：检查 CUDA 编译器是否就绪。

进阶：针对 AI 开发的额外补充

为了让你以后跑模型更顺滑，建议顺手装上这几个库：

1. NCCL（多显卡通信库，分布式训练必装）：

   sudo apt install libnccl2 libnccl-dev
   

2. NVML（监控库）：通常已随驱动安装。

总结你的“一键式”清单：

1. dpkg -i cuda-keyring... (换源)
2. apt install cuda (装驱动+工具链)
3. apt install libcudnn... (装深度学习加速库)
4. export PATH... (设环境变量)

当你完成这些后，再去你的虚拟环境中安装 flash-attn，它就会自动找到 nvcc 并开始愉快的编译了。需要我帮你写一个自动执行这些步骤的 .sh 脚本吗？

作者：互联网容器团队-Chen Han、AI 研发团队 - Liu Dong Yang

在大规模GPU容器集群与模型训练场景，面临稳定性和资源利用率等多重挑战。本文展示vivo GPU平台的总体架构，介绍容器平台在大规模GPU容器集群稳定性建设措施，以及探索多种GPU容器降本提效的解决方案。分享AI工程训练平台大规模训练稳定性建设，及GPU利用率提升实践经验。

本文为2025年 vivo 开发者大会互联网技术专场分享内容之一，在微信公众号"vivo互联网技术"对话框回复【2025VDC】获取 2025VDC 互联网技术会场议题相关资料。

1分钟看图掌握核心观点👇

图1 VS 图2，您更倾向于哪张图来辅助理解全文呢？欢迎在评论区留言

一、GPU平台架构

vivo的GPU平台由物理层、容器平台层与AI工程层三方面构成。由多种GPU服务器和分布式存储以及高性能网络等基础设施，构成了可靠的物理层。容器平台层的GPU容器能力，主要包含了资源管理、编排调度、GPU虚拟化与多容器网络这四个方面。

其中资源管理，表现为多种架构资源池的部署与管理能力。编排调度能力，由GPU弹性伸缩、训推潮汐部署以及多种卡调度策略组成。自研的GPU虚拟化囊括了业界主流的MIG虚拟化、内核层虚拟化以及CUDA层虚拟化三种技术。由传统的Underlay网络以及SRIOV的RDMA直通网络，组成了丰富的容器网络架构。容器平台提供了开放的API接口，为AI工程层的训练和推理平台，提供了坚实的算力底座。通过训练和推理平台，支撑公司内的智能计算业务。

二、GPU容器能力实践

GPU容器能力实践分为两个模块，首先是大规模容器集群稳定性建设，其次是GPU容器提效降本方案。先了解下容器平台在大规模容器集群场景，如何进行稳定性建设的。

2.1 大规模容器集群稳定性

集群稳定性是一切的基石。当集群规模大时，任务多，调度频繁，导致核心组件负载激增，极易发生集群崩溃。随着节点规模扩大，运维复杂度呈指数级增长，日常运维工作繁重，发现问题不及时。同时故障处理也面临严峻挑战，故障中涉及的复杂场景多，故障处理的难度大。稳定性建设需要解决上述问题。

为了解决高频调度导致的核心组件高负载问题，我们针对Apiserver、etcd、CoreDNS，这3个核心组件进行了架构和性能优化，具体的方案如图所示。通过这些优化手段提升了组件性能，并且降低了组件负载，有利于大规模集群的平稳运行。

为了减轻集群运维负担，我们重点建设了自动化节点管理方案。把重复性的运维事项自动化。同时我们还完善了监控告警体系，开发了自动化巡检功能，使运维人员能够及时发现集群问题，快速介入，处理潜在风险，保障集群能够长久稳定运行。

故障处理是集群稳定的兜底措施，我们针对多个核心组件都做了，各类故障处理预案。结合可能存在的故障特点，构造故障场景，进行故障恢复演练，确保故障发生时，能够第一时间找到合理的解决方案，准确的处理问题。

通过上述的措施，在集群稳定性方面取得了不错的效果，首先日常的集群可用性稳定保持在99.99%水平，其次平台的年度故障复盘数相较于上一年下降60%。核心组件方面的优化也达到了不错效果，其中Apiserver的CPU负载下降70%，etcd提交延迟，从秒级缩短到毫秒级，CoreDNS的毛刺现象消失了，并且负载量下降了90%左右。

2.2 GPU容器提效降本实践

容器平台的核心竞争力之一就是助力业务提效降本，我们从不同业务维度，对GPU容器提效降本方案进行了探索。

• 首先在单卡维度，通过自研GPU虚拟化方案，使多个容器，互不干扰的共享一张卡资源。
• 其次是在单服务维度，使业务能够自动应对，负载变化的GPU弹性扩缩容方案。
• 在多服务维度，能够让推理服务和训练服务，分时复用整机资源的，训推潮汐部署方案。
• 最后是在多机多卡的分布式场景中，让GPU容器搭配RDMA网络，来解决跨节点通信的瓶颈问题。

2.2.1 单卡共享-GPU虚拟化

如何让一张卡同时运行多个容器又不互相干扰，就涉及到GPU虚拟化技术。GPU虚拟化一直是AI云原生领域的热门话题之一，各大云厂商都有成熟的解决方案售卖。

vivo容器平台的自研GPU虚拟化方案，主要为了解决业务的三大痛点，

• 首先是部分推理业务负载偏低，无法有效用满整卡资源，需要通过共享部署方式，减少资源总量，降低业务成本，提升利用率。
• 其次是不同业务共享同一张卡时，对于安全性以及隔离性的要求各不相同，就需要使用不同的GPU虚拟化技术来满足不同业务诉求。
• 最后在Dev开发机场景，用户使用频率偏低，需要通过显存超售，来提升资源复用率，但显存超售后又需要避免某个用户将显存耗尽，导致OOM错误影响同卡的其他用户。

自研GPU虚拟化方案包含MIG虚拟化、内核层虚拟化、CUDA层虚拟化这三种技术。结合业务场景，提供了丰富的卡调度策略，例如尽量聚集的Binpack策略、尽量分散的Spread策略，每个卡只有一个实例的CardOnlyOne策略，以及自定义节点和卡分配关系的CustomTopo策略。通过自研模块与组件，接入Kubernetes体系，对外提供统一调度能力。

首先，MIG虚拟化技术，是基于Nvidia硬件提供的，切块组合能力，能够按规则把计算单元和显存单元进行组合，组成MIG实例挂载到容器内，提供完全独立的运行环境。MIG方案的优点是拥有Nvidia官方支持，可以集成到自研体系中。由于是在硬件层面实现的算力和显存限制，所以隔离性和安全性最好。缺点就是仅支持Ampere及以后架构的部分卡，而且限定了切分比例。主要应用场景是对算力隔离有强需求的线上业务。

内核层虚拟化技术，是通过自研内核模块，创建虚拟字符设备替换原有的Nvidia字符设备，在内核态拦截IOCTL请求后，实现的算力和显存限制。优点是上层应用无感。并且内核态拥有良好的安全性。缺点是当前无开源方案，开发难度大，而且算力隔离的并不充分。主要应用场景是常规线上业务。

CUDA层虚拟化技术的原理：使用拦截库替换Nvidia Driver的原始库，建立拦截库与原始库的，API函数映射关系，从而拦截调用函数，实现算力和显存的限制。

优点是有开源方案，使用起来比较灵活。并且可以基于Nvidia提供的统一内存模型，开发显存超售能力。能够在显存不足时，使用内存替代，虽然处理速度下降，但是能够有效避免，显存OOM导致用户程序报错。

缺点是用户态导致安全性不足，并且算力隔离能力偏弱，主要应用场景是Dev开发机场景。

将自研的内核层虚拟化方案与业界方案，进行了自测性能对比，如图所示，可以看到自研方案在性能上，已经达到业界先进水平。业务使用该方案，与独占整卡部署相比较，平均单卡虚拟化率300%左右，就是把1张物理卡当3张卡使用，同时整机GPU利用率提升了30%+，成本优化超过50%。

2.2.2 服务提效-GPU弹性扩缩容

在单服务维度，如何帮助业务自动管理大量的GPU容器是提效的关键。我们引入了GPU弹性扩缩容方案。

• 首先弹性扩缩容能力，能够快速响应负载变化，自动调节实例数量，减少人工干预次数，有利于业务在突发场景的平稳运行。
• 其次是业务方在生产环境部署后，非生产环境的实例通常会闲置，这会浪费稀缺的GPU资源。
• 最后由于Kubernetes原生，并不支持GPU维度的弹性扩缩容，需要寻找合适的方案来满足业务诉求。

如图所示，我们是基于开源的KEDA框架，自研了GPU-Scaler组件，使用Prometheus中存储，来自DCGM-Exporter的GPU指标，汇聚为扩缩容事件，用于触发KEDA框架，调整实例个数，以此实现了GPU弹性扩缩容能力。

由于KEDA框架支持将Workload实例数缩容到0，所以在非生产环境部署的GPU业务，默认开启无负载时，自动缩容到零的功能，以此自动回收，长期闲置的GPU资源。

最终的使用效果，线上业务资源不足类告警，下降了80%，单业务平均减少约每周1小时的，扩缩容工作量，有效降低了GPU业务的运维成本。

2.2.3 多服务降本-训推潮汐部署

在多服务维度，训练服务的资源短缺问题，与推理服务，低峰时段资源空闲问题，相对突出。考虑让训练业务利用推理的空闲资源，即训推潮汐部署方案。

首先推理和训练业务都需要稳定的运行环境。而且推理业务潮汐特征明显，夜晚负载低，资源空闲多，导致平均利用率偏低。并且多机多卡训练任务，需要整机资源，且资源需求日益增长，采购新设备，难且慢，导致训练资源缺口明显。

训推潮汐部署就是整机资源分时复用的逻辑，如图所示，推理业务在白天高负载时，稳定运行，在夜晚低峰时段，自动腾挪出空闲整机资源，借给训练业务使用。在清晨时段，训练业务结束，把整机资源还给推理业务，如此达到分时复用的效果。

如图所示。推理业务在部署前，需要评估保底负载容量。在部署时填入维持业务稳定的最少Pod数量。基于OpenKruise组件的，WorkloadSpread功能，管理不同的Subset，分别在稳定池和潮汐池中按需部署。同时配置CronHPA，定时缩容，自动调整副本数，到稳定Pod数量，优先删除潮汐池中的Pod。以此达到把潮汐池的节点整机腾空的效果。

其中我们还针对Workload的缩容优先级进行了优化。当缩容发生时，结合Pod和节点的拓扑关系，把所在节点实例数少的Pod优先缩容，达到更快的腾空效果。

通过上述方案，训推潮汐部署的降本效果明显，使推理业务，成本下降30%，同时整机GPU利用率提升20%多，有效缓解了训练资源短缺问题。

2.2.4 多机多卡提效-容器RDMA高性能网络

当前分布式训练和推理业务，对算力和显存的需求巨大，单节点资源不足，需要使用多机多卡资源，那么网络通信容易成为性能瓶颈。RDMA技术允许GPU直接访问支持RDMA设备中的数据，无需经过主机CPU或内存，实现跨节点的零拷贝数据传输，有效减少了CPU开销和网络延迟。所以从多机多卡维度，使用RDMA技术是网络提效的有效措施。从容器平台角度，GPU容器更加需要结合RDMA技术，提供简单高效的解决方案，方便业务使用。

如图所示，RDMA容器有两个网卡，一个是使用Calico-CNI插件，通过veth创建的eth0网卡，对应的是Underlay网络。另一个是使用Sriov-CNI插件，通过VF创建的eth1网卡，对应的RoCE_v2或IB协议网络。我们引入了Multus-CNI组件，能够在单容器创建时，按需调用多种CNI插件。同时我们选择使用Spiderpool组件管理IP池，以及进行IP分配和路由策略配置。

通过上述组件，在Kubernetes体系中实现了RDMA容器的全生命周期管理能力。基于容器RDMA能力，在大规模训练和推理场景，业务能够提速20%到30%之间，提升效果明显。

三、AI工程训练平台实践

3.1 训练平台整体架构

VTraining训练平台是由vivo AI计算平台团队打造的一站式大模型训练方案，它面向算法工程师，提供模型开发、模型训练和海量样本存储等能力。

VTraining底层是基于vivo容器平台、及高性能计算、网络、存储等基础设施之上构建，通过提供模型开发、模型训练、资产管理等能力，支撑公司的大模型训练业务。

像vivo手机的蓝心小V、相册等核心产品的大模型训练，都是在VTraining平台进行迭代的。

3.2 大规模训练稳定性实践

3.2.1 稳定性背景

稳定性问题是大规模训练的痛点。任务在大规模的同步训练过程中需要依赖计算、网络、存储、容器调度等复杂的基础设施环境，任何环节出问题都会导致任务中断，问题定位、恢复困难。

例如知名头部公司千亿参数大模型的大规模训练任务，平均每3小时触发一次意外中断。

3.2.2 稳定性实践-高频故障专项治理

其中一个问题，是GPU集群投入使用初期机器故障率会很高，训练任务会经常被中断。为了降低机器故障率，我们进行了高频故障专项治理。首先对GPU集群进行了大规模测试诊断、然后进行高频故障统计和修复，把有问题的软硬件进行维修、替换、升级或修复。

3.2.3 稳定性实践-故障处置流程完善

另一个问题，是任务故障就是不可避免的，所以为了缩短任务中断时间，尽快恢复任务运行，我们对故障处置流程进行了重点建设完善。

• **训练前，**我们会针对GPU机器、网络通信等环境进行大规模检测，剔除异常节点、慢节点等问题节点，降低故障风险。
• **训练中，**会开启自动化容错机制，以便能及时发现基础设施或任务异常，快速进行故障定位和故障容错，自动隔离故障、自动重启恢复任务运行。
• **训练后，**对新问题进行搜集分析，完善异常特征库，增强问题诊断能力，以便后续遇到类似的故障可以快速容错。

3.2.4 稳定性实践-效果与总结

通过减少基础设施高频故障、完善任务故障处置流程两大措施，我们取得了良好效果：机器每天故障率由原来的百分之二下降到千分之一；千卡任务有效训练时长由原来的60%提升到99%，达到了行业一流水平。

同时我们也积累了相关的稳定性经验：

• GPU集群由不稳定到稳定，需要一个软硬件磨合过程。因为不同的软硬件环境会触发不同的稳定性问题。
• 大规模训练前，尽量剔除历史故障率高的机器。我们发现稳定的机器一般会一直很稳定，而故障率高的机器即使修复后出问题的概率也比较大。
• 提升任务有效训练时长，需结合基础设施、训练框架、平台容错机制综合优化。例如秒级监控告警能力、checkpoint持久化策略等等都需要进行持续深入的优化。

3.3 GPU利用率提升实践

3.3.1 GPU利用率提升-业务背景及问题

GPU是稀缺资源，但是差异化的业务场景下GPU难以高效利用，利用率提升困难。比如GPU场景常见业务形态：有训练任务、推理业务、数据生产、开发调试等等类型。他们各有特点：

• 训练任务虽然GPU利用率高，但是偶尔也会出现碎片化空闲资源，比如算法同学偶尔会将任务停下来排查问题，调下参数之类的操作，任务并不是一直不间断地在跑的。
• 推理业务有很明显的潮汐特性，白天流量高峰期GPU利用率高、到了夜间流量低峰期利用率会比较低。
• 数据生产任务属于离线任务，GPU利用率高、资源需求大、但难申请到资源；开发调试任务的话GPU利用率低并且要长期占用资源，导致GPU利用率长期低下。

所以不同的业务形态有不同的利用率特点，接下来介绍下我们的GPU利用率提升措施。

3.3.2 利用率提升措施一：低优任务

对于训练任务场景，偶现的碎片化空闲资源，我们 通过低优数据生产任务进行充分利用。如图所示，我们通过低优数据生产任务调度，复用训练场景偶现的碎片化资源，当正常任务需要资源时，可随时抢占低优资源，不会影响正常训练任务的调度。

3.3.3 利用率提升措施二：训推潮汐部署

对于推理业务场景，在夜间流量低峰期可以释放大量GPU资源，我们通过训推潮汐部署给离线业务复用。如图所示，通过训推潮汐部署机制，我们将夜间推理流量低峰期缩容的机器腾挪到了离线GPU资源池，给离线业务使用，白天再腾挪回在线GPU资源池进行扩容，达到潮汐复用的目的。

3.3.4 利用率提升措施三：GPU虚拟化

对于开发任务场景，长期独占GPU资源且利用率低，我们通过vivo自研VGPU虚拟化技术，减少开发任务占用的物理GPU卡数，释放冗余算力。如图所示，我们可以将单机4卡GPU机器，通过开启VGPU虚拟出16卡VGPU，相当于一卡顶四卡。

VGPU的优点是：可支持1:2、1:4超卖、还可以用内存补充显存不足，所以对用户是无感知使用的，很适用于开发任务这种对性能要求低的场景。

3.3.5 利用率提升总结与规划

总之，训练平台通过低优任务、训推潮汐部署、GPU虚拟化等策略，深度适配差异化业务场景特性，实现了资源高效复用。AI整体GPU利用率均值绝对值提升了5个百分点，接近行业一流水平。

未来训练平台也会持续对GPU利用率进行综合治理。例如进行低效任务治理、低效资源盘活、成本账单统计、奖励与惩罚措施的实施等等，让稀缺的GPU资源发挥更大价值！

四、GPU平台未来展望

在容器平台层面，我们会从多集群联邦调度、在离线GPU混部、国产异构计算芯片支持、GPU资源池化等方面能力进行综合建设，支撑上层平台业务对GPU资源的高效利用。

训练平台层面，我们会增强故障预警、任务动态容错等能力，增强业务稳定性；同时完善对大模型训练全流程能力的支撑，以及对GPU资源进行更精细化的运营，从而让GPU业务更加稳定、资源利用更加高效！

编辑：桃子

【新智元导读】英伟达护城河要守不住了？Claude Code半小时编程，直接把CUDA后端迁移到AMD ROCm上了。

一夜之间，CUDA护城河被AI终结了？

这几天，一位开发者johnnytshi在Reddit上分享了一个令人震惊的操作：

Claude Code仅用了30分钟，便将一段完整的CUDA后端代码，成功移植到AMD的ROCm上。

整个过程，没有手写一行代码。

这架势，简直是要填平这两个生态系统之间的鸿沟。

更关键的是，这次移植完全没有依赖传统的「中间转换工具」，如Hipify翻译层，而是一键通过CLI完成。

就连AMD软件副总Anush E.为之震惊，GPU编程的未来，是AI智能体的。

消息一出，整个科技圈瞬间沸腾，很多人直呼：英伟达CUDA护城河要守不住了…..

这究竟是怎么回事？

Claude手撕CUDA，仅30分钟

Claude Code是在一个智能体框架运行的，这意味着它可以自己「动脑子」。

在执行过程中，他不会机械地转换关键词，而去真正理解代码，即特定核函数的底层逻辑。

开发者johnnytshi介绍，这次移植中，最棘手的数据布局差异问题也被AI解决了，确保了内核核心计算逻辑保持一致。

令人惊叹的是，johnnytshi在短短30分钟内，就把整个CUDA后端移植到了AMD ROCm上，而且中间没用任何翻译层。

另外一个好处当然是，不用费劲去搭像Hipify这种复杂的翻译环境了；直接在命令行（CLI）里就能干活。

如今，全网都被CUDA护城河被攻破呼声淹没了。

毕竟，英伟达霸主地位，很大程度上建立在CUDA这个几乎成为行业标准的编程生态上。

无数AI框架、深度学习库、科学计算工具都深度依赖它。

AMD的ROCm虽然功能强大，却一直面临生态兼容性，以及开发者迁移成本高的痛点。

现在，一个Claude却用极短时间踢碎了门槛，说不定未来更多CUDA代码可能轻松在AMD GPU跑起来了。

实现细节

GitHub中，johnnytshi本人也更新了日志和说明。

为AMD GPU实现了完整的ROCm后端，从而在RDNA 3.5及其他AMD架构上支持基于注意力机制的现代国际象棋网络。

GitHub：https://github.com/LeelaChess...

• 在src/neural/backends/rocm/中添加了完整的ROCm后端
• 实现了注意力网络架构（多头自注意力、FFN、嵌入层）
• 使用rocBLAS进行GEMM运算，使用MIOpen进行卷积运算
• 针对RDNA 3.5上的FP16性能优化了NCHW布局
• 提供三种后端变体：rocm (FP32)、rocm-fp16 (FP16)、rocm-auto (自动检测)
• MIOpen是必选依赖（类似于CUDA的cuDNN）
• 通过rocm\_agent\_enumerator自动检测AMD GPU架构
• 编译选项：-Drocm=true -Damd\_gfx=gfx1151（或使用自动检测）

性能说明：

• FP16性能：在Strix Halo (Radeon 8060S, gfx1151) 上 >2000 nps
• 自动Batch Size调优（RDNA 3.5上min\_batch=64）
• 测试过rocWMMA，但rocBLAS性能更好

验证情况（Strix Halo – Radeon 8060S, gfx1151）：

• 测试模型：768x15x24h-t82-swa-7464000.pb.gz 和 maia-1900.pb.gz
• 后端：rocm-fp16功能正常，能生成正确的走法
• 环境：ROCm 7.2.53150, MIOpen 3.5.1
• 注：仅在RDNA 3.5上进行了测试；其他AMD架构暂未验证

GPU未来，是AI智能体主场

当然，这次演示也有局限性。

对于简单或中等复杂度的内核，Claude Code表现得非常出色。更重要的是，写核函数的核心就在于搞定「深度硬件」优化。

不过，一部分觉得Claude Code在这方面还是差点火候——

如果遇到那些针对特定硬件缓存层级，内存访问模式做过极致优化的复杂内核，AI目前还难以完全取代人类专家。

即便如此，这一事件释放出的信号已经足够强烈。

过去几个月，ZLUDA项目、还有微软内部的尝试，都想要打破CUDA的垄断。

但它们大多依赖规则映射或中间层，自动化程度和智能水平有限。

Claude Code代表的智能体式编程，直接跳过了这些环节，用「理解+自主决策」的方式填平生态鸿沟。

正如AMD软件副总所言，GPU编程的未来，是AI智能体主场。

全员AI编程，浓度高达100%

如今的Claude Code已经让整个硅谷入坑了（Claude-Pilled）。

两天前，CEO Dario Amodei在达沃斯上再出暴论：软件工程师们没有时间了。未来6-12个月，AI能够彻底取代这些人！

甚至，Anthropic内部工程师已经不再手写代码了，全是Claude完成。

别不信，是真的。

就在Wired最新采访中，Claude Code之父Boris Cherny坦承，「自己100%代码都是AI写的」。

或许Anthropic工程师怎么也没有想到，一个「副业项目」竟让硅谷如此狂热。

Boris Cherny回忆道，「一年前我们发布Claude Code时，甚至不确定『智能体编程』能不能成，但火爆来得太快了」。

Cherny个人经历就是最好的缩影：

刚发布时，他只有5%代码是用Claude Code写的；

到了去年5月，有了Opus 4和Sonnet 4，这个比例变成了30%；

而现在，有了Opus 4.5，他在过去两个月里100%的代码都是由Claude Code完成。

在Anthropic内部，这种全员AI化更是到了极致。

几乎100%技术员工都在使用Claude Code，甚至连Claude Code团队本身95%的代码也是由自身写出来的。

斯坦福AI教授都在用了

不得不说，AI编程的进化速度令人咋舌。

回望2021到2024年，大多数工具不过是高级版的「自动补全」，在开发者打字时卑微地建议几行代码。

但到了2025年初，随着Cursor和Windsurf等初创发布早期的Agentic编程产品，游戏规则改变了——

开发者只需用大白话描述功能，剩下的脏活累活全扔给AI智能体完成。

Claude Code也在这个时间点，真正诞生了。

Boris Cherny坦承，早期版本也曾跌跌撞撞，甚至陷入死循环。但Anthropic下了一步狠棋：不为当下的AI能力开发产品，而要为AI即将抵达的未来而构建。

这一赌注押对了。随着Anthropic下一代旗舰Claude Opus 4.5的发布，AI编程迎来了真正的「拐点」。

斯坦福大学AI讲师、Workera CEO Kian Katanforoosh最近就把公司全员迁移到了Claude Code。

他直言，对于高级工程师来说，Claude Code比Cursor、Windsurf更能打。

Katanforoosh感叹道，最近唯一让我看到编程能力有阶跃式提升的模型，就是Claude Opus 4.5。

「它给人的感觉不像是在模仿人类写代码，而是它真的找到了一种更聪明的解决路径」。

据传，微软内部也在大规模采用Claude Code了。

年入超10亿美金的「副业」

Claude Code大获成功，给Anthropic带来了最直观的效益。

去年，AI编程智能体业务彻底爆发。11月，Anthropic宣布Claude Code在上线不到一年内，年度经常性收入（ARR）就突破了10亿美元。

到2025年底，ARR至少又增长了1亿美元。

彼时，该产品约占Anthropic总ARR（约90亿美元）的12%。虽然比起向大企业提供 AI 系统的核心业务来说还算「小弟」，但它已是公司增长最快的板块之一。

尽管Anthropic在AI编程领域看似独孤求败，但Claude Opus 4.5的光环其实照亮了整个赛道。

竞争对手Cursor也在11月达到了10亿美元ARR，OpenAI、谷歌和xAI更是磨刀霍霍，试图用自研模型分一杯羹。

但Anthropic没打算停下。

前几天，他们又发布了Cowork——这是一款面向非编程领域的AI智能体。

它能管理你电脑里的文件、操作各种软件，而且完全不需要你在代码终端里敲命令。

不是取代，是进化

提及Cowork时，Cherny透露自己已经用疯了。

比如项目管理，他会让Cowork盯着工程师的任务表格，谁没填名字，AI就会自动在Slack上发消息催人。

Cherny感慨道，「这是我当工程师以来最爽的时候，因为我不再需要做那些枯燥乏味的脏活了」。

面对那些因不再需要亲自写代码而感到失落的工程师，Cherny给出了他的建议：

这行业一直在变。我祖父在苏联用穿孔卡片编程；后来变成了机器码；再后来是C语言、Java、Python。

这是一条不断抽象化的连续体，AI智能体只是这条线上的最新一个点。

如今，Cherny每天早上起床会在手机上启动3-4个编程智能体，到了公司再在终端里开几个。

任何时候，他都有五到十个智能体在跑任务。

Cherny总结道，「AI智能体将接管生活中所有繁琐的事——填表、搬运数据、发邮件。这会具有颠覆性，我们必须适应」。

话又说回来，Anthropic能不能先解决下Claude使用量？

深入 NVIDIA GPU：高性能矩阵乘法（Matmul）算子解构
在本篇博文中，我将逐步介绍支撑最尖端（SOTA）NVIDIA GPU 矩阵乘法（matmul）算子的核心硬件概念和编程技术。

为何选择矩阵乘法？

无论是训练还是推理阶段，Transformer 模型的大部分浮点运算（FLOPs）都消耗在矩阵乘法中（如 MLP 中的线性层、Attention 的 QKV 投影、输出投影等）。这些操作具有天然的极高并行性（Embarrassingly Parallel），非常适合 GPU。掌握了矩阵乘法算子的原理，你就拥有了设计几乎任何其他高性能 GPU 算子的工具箱。

本文分为四个部分：

1. NVIDIA GPU 架构基础：全局内存、共享内存、L1/L2 缓存、功率限制（power throttling）对算力极限（SOL）的影响等。
2. GPU 汇编语言：SASS 和 PTX。
3. 设计近乎 SOTA 的同步矩阵乘法内核：线程束平铺（warp-tiling）方法。
4. 在 Hopper 上设计 SOTA 异步矩阵乘法内核：利用张量核心（Tensor Cores）、TMA、计算与加载/存储重叠、希尔伯特曲线（Hilbert curves）等。
   我的目标是让这篇文章自成体系：既有足够的细节供独立阅读，又足够简洁以避免变成教科书。

本文是系列文章的首篇。后续计划（理想状态下）涵盖：
• 在Blackwell GPU上设计顶尖矩阵乘法内核
• 通过微基准测试探索GPU架构
• 设计顶尖多GPU内核
• 揭秘内存一致性模型（GPU领域的“令牌化器”：默默支撑系统运行的关键组件，却令多数开发者困惑不已）

NVIDIA GPU 架构基础

要编写高性能的 GPU 内核，你需要对硬件有一个扎实的认知模型。随着我们深入探讨硬件架构，这一点会很快变得清晰。
在本文中，我重点关注 Hopper H100 GPU。如果你能深度理解 Hopper，那么将知识迁移到未来架构（Blackwell, Rubin）或早期架构（Ampere, Volta）就会变得非常简单。
Hopper [1] 和 Ampere [2] 白皮书是非常好的信息来源。
在最高层面上，GPU 执行两个基本任务：

1. 移动和存储数据（内存系统）。
2. 对数据进行有用的操作（计算流水线）。
   下方的 H100 框图反映了这种划分：蓝色组件代表内存或数据移动，而红色组件代表计算（热）单元。
   

图 1：NVIDIA Hopper H100 GPU 模型
如果你在文中发现任何错误，请直接联系我——欢迎在 X、LinkedIn 或通过匿名反馈给我留言。

内存（Memory）

GPU 的内存系统是高度分层的，非常类似于 CPU 架构。
这种分层是由物理学和电路设计决定的：SRAM 单元速度更快但体积更大（实现高速所需的控制电路也增加了其面积），而 DRAM 单元体积更小/密度更高但速度较慢。其结果是，高速内存容量低且昂贵，而慢速内存可以提供大得多的容量。稍后我们将更详细地讨论 DRAM 单元/内存。

这种容量与延迟之间的权衡正是缓存层级存在的原因。在理想世界中，每个计算单元都会坐落在一大池超快内存旁边。由于这在物理上是不可能的，GPU 设计者做出了妥协：将少量快速内存放置在靠近计算单元的地方，并由更远处容量逐渐增大、速度渐慢的内存池作为后盾。这种组织方式最大化了整体系统的吞吐量。

GPU 内存系统由以下部分组成：

1. 设备内存（VRAM/Device Memory）：在 CUDA 术语中，“设备”内存指的是片外（off-chip）DRAM——物理上与 GPU 芯片（die）分离，但封装在同一个板卡上——通常以堆栈式的 HBM 实现。它承载全局内存（GMEM）、每个线程的“局部”内存（寄存器溢出空间）等。
2. L2 缓存（L2 Cache）：由 SRAM 构建的大容量 k 路组关联缓存。它在物理上分为两部分；每个 SM 直接连接到一个分区，并通过横截（crossbar）间接连接到另一个分区。
3. 分布式共享内存（DSMEM）：物理上接近的一组 SM（即一个 GPC）中共享内存（SMEM）的池化。
4. L1 缓存与共享内存（Shared Memory）：
   ￮ L1 缓存：每个 SM 私有的较小 k 路组关联 SRAM 缓存。
   ￮ 共享内存（SMEM）：程序员管理的片上内存。SMEM 和 L1 共享相同的物理存储，它们的相对比例可以通过软件配置。

5. 寄存器堆（Register File/RMEM）：位于计算单元旁边的最快存储单元。寄存器是单个线程私有的。与 CPU 相比，GPU 包含多得多的寄存器，且总 RMEM 容量与 L1/SMEM 存储的总和相当。

   
   图 2：H100 (SXM5) GPU 的内存层级
   📝 注意： 还有一些用于指令的小型缓存，以及常量内存等，为了理解核心原理，我将忽略它们。
   从设备内存向下移动到寄存器（第 1-5 级），你会看到一个明显的趋势：带宽以数量级增长，而延迟和容量则以类似的数量级减少。
   这引发了一些直接的影响：

6. 将访问最频繁的数据尽可能靠近计算单元存放。

7. 尽量减少对层级结构底层的访问，尤其是设备内存（GMEM）。
   另一个值得注意的组件是 张量内存加速器（TMA），它是随 Hopper 引入的。TMA 支持在全局内存和共享内存之间，以及集群（cluster）内的共享内存之间进行异步数据传输。它还支持交织（swizzling）以减少银行冲突（bank conflicts）——我们会在适当的时候讨论这些细节（双关语）。

   计算（Compute）

   从内存转向计算，其基本单位是流式多处理器（SM）。Hopper H100 (SXM5) 总共集成了 132 个 SM。
   SM 被组织成图形处理集群（GPC）：每个 GPC 包含 18 个 SM，GPU 上共有 8 个 GPC。四个 GPC 直接连接到一个 L2 分区，另外四个连接到第二个分区。
   📝 注意：
   • GPC 也是支撑 CUDA 中 线程块集群（thread-block cluster） 抽象的硬件单元——我们稍后会回到编程模型。
   • 关于集群的一点：早前我说过每个 GPC 有 18 个 SM，所以 8 个 GPC 应该有 144 个 SM。但 SXM/PCIe 规格暴露的是 132 或 114 个 SM。差异在哪里？这是因为 18 × 8 的布局仅对完整的 GH100 芯片有效——在实际产品中，有些 SM 会被熔断（fused off）。这对我们编写内核时选择集群配置有直接影响。例如，如果集群跨度超过 2 个 SM，你就无法利用所有 SM。
   • 最后注意，“Graphics Processing Cluster (GPC)”中的“Graphics”是一个传统术语。在现代服务器级 GPU 中，这些集群纯粹作为计算/AI 加速单元，而非图形引擎。同样的，GPU 应该去掉“G”，它们是 AI 加速器。

除了前面提到的 L1/SMEM/TMA/RMEM 组件（均位于 SM 内部），每个 SM 还包含：
• 张量核心（Tensor Cores）：以高吞吐量在小分块（例如 64x16 @ 16x256）上执行矩阵乘法的专用单元。大型矩阵乘法被分解为许多此类分块操作，因此有效利用它们是达到峰值性能的关键。
• CUDA 核心与 SFU：所谓的“CUDA 核心”（营销话术）执行标准的浮点运算，如 FMA（融合乘加：c=a∗b+c）。特殊函数单元（SFU）处理超越函数（如 sin,cos,exp,log）以及代数函数（如 sqrt,rsqrt 等）。
• 加载/存储（LD/ST）单元：服务于加载和存储指令的电路，与 TMA 引擎互补。
• 线程束调度器（Warp Schedulers）：每个 SM 包含调度器，为 32 个线程组成的组（CUDA 中称为 warps）发布指令。一个线程束调度器每周期可以发布一条线程束指令。
每个 SM 在物理上分为四个象限，每个象限容纳上述计算单元的一个子集。

这导出了以下见解：
📝 并行性（Parallelism）与并发性（Concurrency）
• 一个 SM 在给定的周期内最多可以同时发布来自四个线程束的指令（即在真正的并行执行中，每周期有 128 个线程）。
• 然而，一个 SM 可以容纳多达 2048 个并发线程（64 个线程束）。这些线程束常驻在 SM 中，并随着时间的推移被换入和换出调度，允许硬件隐藏内存/流水线延迟。
换句话说，指令并行性（在给定周期内有多少线程开始执行指令）限制为每 SM 128 个线程（4 条 32 宽度的线程束指令），而并发性（调度器中跟踪并有资格运行的线程数）则扩展到 2048 个线程。

光速（Speed of Light）与功率限制

既然我们购买 NVIDIA GPU 是为了计算，自然会问：性能上限是什么——GPU 的最大计算吞吐量是多少？这通常被称为“光速”（Speed of Light, SoL）性能：由芯片物理特性决定的上限。

根据数据类型的不同，有不同的性能上限。在 LLM 训练工作负载中，bfloat16 (bf16) 是近年来的主导格式，尽管 fp8 和 4 位格式变得越来越重要（对于推理，fp8 已相当标准）。
峰值吞吐量的计算公式为： perf=freq_clk_max∗num_tc∗flop_per_tc_per_clk
或者用文字描述：最大时钟频率 × 张量核心数量 × 每个张量核心每周期的浮点运算数（FLOPs）。

图 3：H100 SXM5 BF16 “光速”推导

📝 FLOP vs FLOPs vs FLOPS vs FLOP/s
• FLOP = 单次浮点运算。
• FLOP/s = 吞吐量单位：每秒浮点运算次数。
• FLOPs（小写 s）= FLOP 的复数（多次运算）。
• FLOPS（全大写）常被误用来表示吞吐量，但严格来说应读作“FLOPs”（复数）。将 FLOPS 用作 FLOP/s 是不严谨的！:)
我在上图中留下了一个提示：“光速”实际上并不是恒定的（我想这也是这个比喻失效的地方）。
在实践中，峰值吞吐量取决于实际时钟频率，而时钟频率会因功率限制（Power Throttling）或温度限制而波动。如果 GPU 时钟频率下降，有效的光速也会随之下降。

图 4：功率限制降低了时钟频率并拉低了有效的“光速”

📝 延伸阅读： Horace He 在他的博文 [3] 中更深入地探讨了这一现象。
硬件细节目前了解这些就足够了。接下来的重点将转向 CUDA 编程模型，然后我们会再次深入硬件底层，并最终上升到 CUDA C++ 层面。

CUDA 编程模型

CUDA 编程模型自然地映射到 GPU 硬件和内存层级结构上。其核心抽象包括：
• 线程 (thread)
• 线程束 (warp)（32 个线程）
• 线程块 (thread block)
• 线程块集群 (thread block cluster)
• 网格 (grid)（由线程块或集群组成）

图 5：CUDA 编程模型：线程、线程束、块、集群、网格
每个线程通过 gridDim、blockIdx、blockDim 和 threadIdx 等变量“意识到”自己在 CUDA 层次结构中的位置。在内部，这些变量存储在特殊寄存器中，并在内核启动时由 CUDA 运行时（runtime）初始化。
这些位置信息使得跨 GPU 分配任务变得简单。例如，假设我们要处理一张 1024×1024 的图像。我们可以将其划分为 32×32 的线程块，每个块包含 32×32 排列的线程。每个线程可以计算其全局坐标：

Plain Text
const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y;

并利用这些坐标从全局内存读取分配给它的像素（imagex），执行点对点操作，并将结果存回。

图 6：CUDA 内置变量：线程如何知道自己在哪里
如上图所示，在实践中我们大多使用 1D 或 2D 的网格/集群/块形状。但在内部，它们始终可以根据需要进行逻辑重组。
例如，如果 threadIdx.x 的范围是 0-1023（1024 个线程的 1D 块），我们可以将其拆分为 x = threadIdx.x % 32 和 y = threadIdx.x / 32，从而将其重塑为 32×32 的逻辑 2D 布局。
将 CUDA 模型连接回硬件，有一点现在应该很清楚了：一个线程块应包含至少 4 个线程束（即 128 个线程）。
为什么？ 线程块驻留在单个 SM 上。每个 SM 有 4 个线程束调度器——为了充分利用硬件，你不希望它们处于闲置状态。
📝 至少 4 个线程束的其他原因：
• 我们稍后会深入探讨，但在 Hopper 架构上，线程束组（warp-group，即 4 个线程束） 是 WGMMA（矩阵乘法）张量核心指令的执行单位。
• 此外，在使用持久化内核（persistent kernels）时，我们通常每个 SM 仅启动一个线程块，因此构建任务以保持所有线程束调度器忙碌至关重要。
带着 CUDA 编程模型的术语，我们可以继续深入探讨 GPU 的架构细节。

全局内存（GMEM）模型

让我们深入探讨 GMEM。如前所述，它是由多层 DRAM 堆叠而成，底部有一个逻辑层（HBM）。但 DRAM 到底是什么？

图 7：DRAM 单元内部：晶体管 + 电容器，字线（wordline） + 位线（bitline）
了解了单个位的存储方式后，让我们放大到整个存储矩阵。

图 8：GMEM 模型

📝 关于 HBM 的延伸阅读： 如果你想更深入地了解 HBM，我发现论文《Demystifying the Characteristics of High Bandwidth Memory for Real-Time Systems》[21] 非常有启发性。
我们可以得出结论：访问模式（access patterns）至关重要，这是由 DRAM 单元的物理特性决定的。

图 9：GMEM 访问模式的影响

Stephen Jones 的演讲《How CUDA Programming Works》[4] 非常值得一看。
如果我们示例中的矩阵是列优先（column-major*的，情况就会反转：列中的元素将连续存储，因此有效的选择是在内层循环遍历行，以避免 DRAM 惩罚。
所以，当人们说“GMEM 合并（coalescing）非常重要”时，指的就是：线程应访问连续的内存位置，以最小化触及的 DRAM 行数。

共享内存（SMEM）模型

共享内存（SMEM）具有与 GMEM 截然不同的特性。它由 SRAM 单元而非 DRAM 构建，这使得它在速度和容量的权衡上完全不同。
SRAM 单元的具体设计并不重要——只需知道存储一位信息需要多得多的晶体管。你可以自行搜索“SRAM cell”。
SMEM 组织为 32 个银行（banks），每个银行宽度为 32 位（4 字节）：

图 10：SMEM 模型
SMEM 可以在单个周期内提供来自所有 32 个银行的数据（128 字节）——但前提是必须遵守一条规则：
同一个线程束（warp）中的线程不得访问同一个银行（bank）内的不同地址。 否则，这些请求将被序列化，分多个周期执行。
这种情况被称为 银行冲突（bank conflict）。如果有 N个线程访问同一个银行中的不同地址，就会产生 N 路银行冲突（N-way bank conflict），该线程束的内存请求将需要 $N$ 个周期才能完成。
在最坏的情况下，所有 32 个线程都指向同一个银行中的不同地址，吞吐量将下降到原来的 1/32。
为了说明这一点，假设线程束大小（warp size）为 5。下方的两种访问模式将分别需要 3 个周期和 1 个周期来完成服务：

重要的是：如果一个线程束（warp）中的多个线程访问同一个银行（bank）内的相同地址，共享内存（SMEM）可以将该值广播（broadcast）或组播（multicast）给所有这些线程。
在下方的示例中，请求在单个周期内即可完成服务：
• 银行 1（Bank 1） 可以将一个值组播给 2 个线程。
• 银行 2（Bank 2） 可以将一个值组播给 3 个线程。

现在，来看硬件拼图的最后一块：L1 缓存。
这是一篇由 Axel 撰写的关于 SMEM 微基准测试（microbenchmarking）的优秀博文 [5]。

L1 模型

我们已经看到 L1 和 SMEM 共享相同的物理存储，但 L1 在该存储周围增加了一层由硬件管理的脚手架层（scaffolding layer）。
在高层级上，L1 缓存的逻辑流程如下：

1. 线程束（warp）发布一个内存请求（指向 SMEM 或 GMEM）。
2. 请求进入 MIO 流水线并被派遣至 LSUIN 路由器。
3. 路由器导向请求：SMEM 访问立即从数据数组（data array）中获得响应，而 GMEM 访问则进入标签比较（tag-comparison）阶段。
4. 在标签阶段，GMEM 的地址标签与目标集合（target set）中存储的标签进行对比，以确定数据是否驻留在 L1 中。
5. 命中（Hit）：请求直接从数据数组中获得服务（就像 SMEM 一样）。
6. 未命中（Miss）：请求传播至 L2（如有必要，甚至更远，直到 GMEM 或对等 GPU 内存）。当数据返回时，它会被缓存到 L1 中，替换（evicting）现有的一行，并并行地发送回发起请求的线程束。
   这就是我刚才描述的系统：
   
   图 13：L1 缓存模型
   让我们再深入一层，详细查看标签阶段（tag stage）和数据阶段（data stage）：

图 14：k 路组关联缓存组织的分解
当一个 GMEM（全局内存） 地址进入标签阶段时，命中/未命中（hit/miss）逻辑按如下方式展开：

1. 标签阶段接收 GMEM 地址。
2. 提取集合 ID 位（set id bits），并检查该集合中的所有缓存行（标签）。
3. 如果发现标签匹配（潜在的缓存命中）：
   ￮ 检查该行的有效性标志（validity flags）。
   ￮ 如果无效 → 视为缓存未命中（继续执行步骤 4）。
   ￮ 如果有效 → 从数据数组中提取所请求的分区（sectors），并交付至线程束的寄存器中。
4. 如果未发现匹配（缓存未命中），请求被路由至内存层级结构的其余部分（L2 及更高层级）。
   ￮ 当数据从 L2 返回时，它被存储在该集合中，并根据替换策略（例如伪 LRU 算法）驱逐（evicting）现有的某一行，同时并行地交付给发起请求的线程束。
   注意，L2 缓存与 L1 并没有太大区别，除了它是全局的（而非每个 SM 独立）、容量大得多（具有更高的关联度）、被划分为由横截（crossbar）连接的两个切片，并且支持更细致的持久化和缓存策略。
   至此，我们已经涵盖了理解后续章节所需的关键 GPU 硬件组件。
   📝 GPU 世代间的梯度：
   我之前提到过，理解 Hopper 是深入了解 NVIDIA GPU 未来和过去世代的绝佳基础。
   到目前为止，最大的世代跨越是从 Ampere → Hopper，引入了：
   • 分布式共享内存 (DSMEM)：在整个 GPC 的 SMEM 之间，实现加载、存储和原子操作的直接 SM 到 SM 通信。
   • TMA (张量内存加速器)：用于异步张量数据移动（GMEM ↔ SMEM, SMEM ↔ SMEM）的硬件单元。
   • 线程块集群 (Thread Block Clusters)：一种新的 CUDA 编程模型抽象，用于跨 SM 对块进行分组。
   • 异步事务屏障 (Asynchronous transaction barriers)：拆分式屏障，计数事务（字节数）而非仅仅是线程数。
   Ampere (例如 A100) 自身也引入了几个关键特性：
   • 张量核心 (Tensor Cores) 支持 tf32 和 bf16。
   • 异步拷贝 (GMEM → SMEM)，具有两种模式：绕过 L1 和访问 L1。
   • 异步屏障（在共享内存中由硬件加速）。
   • CUDA 任务图 (CUDA task graphs)：它是 PyTorch 中 CUDA graphs 的基础，减少了 CPU 启动和网格初始化开销。
   • 通过 CUDA 协作组 (CUDA Cooperative Groups) 暴露的线程束级规约指令（实现了单步、整数数据类型的线程束范围规约，无需 shuffle 模式）。

深入 NVIDIA GPU：高性能矩阵乘法（Matmul）算子解构
在本篇博文中，我将逐步介绍支撑最尖端（SOTA）NVIDIA GPU 矩阵乘法（matmul）算子的核心硬件概念和编程技术。

为何选择矩阵乘法？

无论是训练还是推理阶段，Transformer 模型的大部分浮点运算（FLOPs）都消耗在矩阵乘法中（如 MLP 中的线性层、Attention 的 QKV 投影、输出投影等）。这些操作具有天然的极高并行性（Embarrassingly Parallel），非常适合 GPU。掌握了矩阵乘法算子的原理，你就拥有了设计几乎任何其他高性能 GPU 算子的工具箱。

本文分为四个部分：

1. NVIDIA GPU 架构基础：全局内存、共享内存、L1/L2 缓存、功率限制（power throttling）对算力极限（SOL）的影响等。
2. GPU 汇编语言：SASS 和 PTX。
3. 设计近乎 SOTA 的同步矩阵乘法内核：线程束平铺（warp-tiling）方法。
4. 在 Hopper 上设计 SOTA 异步矩阵乘法内核：利用张量核心（Tensor Cores）、TMA、计算与加载/存储重叠、希尔伯特曲线（Hilbert curves）等。
   我的目标是让这篇文章自成体系：既有足够的细节供独立阅读，又足够简洁以避免变成教科书。

本文是系列文章的首篇。后续计划（理想状态下）涵盖：
• 在Blackwell GPU上设计顶尖矩阵乘法内核
• 通过微基准测试探索GPU架构
• 设计顶尖多GPU内核
• 揭秘内存一致性模型（GPU领域的“令牌化器”：默默支撑系统运行的关键组件，却令多数开发者困惑不已）

NVIDIA GPU 架构基础

要编写高性能的 GPU 内核，你需要对硬件有一个扎实的认知模型。随着我们深入探讨硬件架构，这一点会很快变得清晰。
在本文中，我重点关注 Hopper H100 GPU。如果你能深度理解 Hopper，那么将知识迁移到未来架构（Blackwell, Rubin）或早期架构（Ampere, Volta）就会变得非常简单。
Hopper [1] 和 Ampere [2] 白皮书是非常好的信息来源。
在最高层面上，GPU 执行两个基本任务：

1. 移动和存储数据（内存系统）。
2. 对数据进行有用的操作（计算流水线）。
   下方的 H100 框图反映了这种划分：蓝色组件代表内存或数据移动，而红色组件代表计算（热）单元。
   

图 1：NVIDIA Hopper H100 GPU 模型
如果你在文中发现任何错误，请直接联系我——欢迎在 X、LinkedIn 或通过匿名反馈给我留言。

内存（Memory）

GPU 的内存系统是高度分层的，非常类似于 CPU 架构。
这种分层是由物理学和电路设计决定的：SRAM 单元速度更快但体积更大（实现高速所需的控制电路也增加了其面积），而 DRAM 单元体积更小/密度更高但速度较慢。其结果是，高速内存容量低且昂贵，而慢速内存可以提供大得多的容量。稍后我们将更详细地讨论 DRAM 单元/内存。

这种容量与延迟之间的权衡正是缓存层级存在的原因。在理想世界中，每个计算单元都会坐落在一大池超快内存旁边。由于这在物理上是不可能的，GPU 设计者做出了妥协：将少量快速内存放置在靠近计算单元的地方，并由更远处容量逐渐增大、速度渐慢的内存池作为后盾。这种组织方式最大化了整体系统的吞吐量。

GPU 内存系统由以下部分组成：

1. 设备内存（VRAM/Device Memory）：在 CUDA 术语中，“设备”内存指的是片外（off-chip）DRAM——物理上与 GPU 芯片（die）分离，但封装在同一个板卡上——通常以堆栈式的 HBM 实现。它承载全局内存（GMEM）、每个线程的“局部”内存（寄存器溢出空间）等。
2. L2 缓存（L2 Cache）：由 SRAM 构建的大容量 k 路组关联缓存。它在物理上分为两部分；每个 SM 直接连接到一个分区，并通过横截（crossbar）间接连接到另一个分区。
3. 分布式共享内存（DSMEM）：物理上接近的一组 SM（即一个 GPC）中共享内存（SMEM）的池化。
4. L1 缓存与共享内存（Shared Memory）：
   ￮ L1 缓存：每个 SM 私有的较小 k 路组关联 SRAM 缓存。
   ￮ 共享内存（SMEM）：程序员管理的片上内存。SMEM 和 L1 共享相同的物理存储，它们的相对比例可以通过软件配置。

5. 寄存器堆（Register File/RMEM）：位于计算单元旁边的最快存储单元。寄存器是单个线程私有的。与 CPU 相比，GPU 包含多得多的寄存器，且总 RMEM 容量与 L1/SMEM 存储的总和相当。

   
   图 2：H100 (SXM5) GPU 的内存层级
   📝 注意： 还有一些用于指令的小型缓存，以及常量内存等，为了理解核心原理，我将忽略它们。
   从设备内存向下移动到寄存器（第 1-5 级），你会看到一个明显的趋势：带宽以数量级增长，而延迟和容量则以类似的数量级减少。
   这引发了一些直接的影响：

6. 将访问最频繁的数据尽可能靠近计算单元存放。

7. 尽量减少对层级结构底层的访问，尤其是设备内存（GMEM）。
   另一个值得注意的组件是 张量内存加速器（TMA），它是随 Hopper 引入的。TMA 支持在全局内存和共享内存之间，以及集群（cluster）内的共享内存之间进行异步数据传输。它还支持交织（swizzling）以减少银行冲突（bank conflicts）——我们会在适当的时候讨论这些细节（双关语）。

   计算（Compute）

   从内存转向计算，其基本单位是流式多处理器（SM）。Hopper H100 (SXM5) 总共集成了 132 个 SM。
   SM 被组织成图形处理集群（GPC）：每个 GPC 包含 18 个 SM，GPU 上共有 8 个 GPC。四个 GPC 直接连接到一个 L2 分区，另外四个连接到第二个分区。
   📝 注意：
   • GPC 也是支撑 CUDA 中 线程块集群（thread-block cluster） 抽象的硬件单元——我们稍后会回到编程模型。
   • 关于集群的一点：早前我说过每个 GPC 有 18 个 SM，所以 8 个 GPC 应该有 144 个 SM。但 SXM/PCIe 规格暴露的是 132 或 114 个 SM。差异在哪里？这是因为 18 × 8 的布局仅对完整的 GH100 芯片有效——在实际产品中，有些 SM 会被熔断（fused off）。这对我们编写内核时选择集群配置有直接影响。例如，如果集群跨度超过 2 个 SM，你就无法利用所有 SM。
   • 最后注意，“Graphics Processing Cluster (GPC)”中的“Graphics”是一个传统术语。在现代服务器级 GPU 中，这些集群纯粹作为计算/AI 加速单元，而非图形引擎。同样的，GPU 应该去掉“G”，它们是 AI 加速器。

除了前面提到的 L1/SMEM/TMA/RMEM 组件（均位于 SM 内部），每个 SM 还包含：
• 张量核心（Tensor Cores）：以高吞吐量在小分块（例如 64x16 @ 16x256）上执行矩阵乘法的专用单元。大型矩阵乘法被分解为许多此类分块操作，因此有效利用它们是达到峰值性能的关键。
• CUDA 核心与 SFU：所谓的“CUDA 核心”（营销话术）执行标准的浮点运算，如 FMA（融合乘加：c=a∗b+c）。特殊函数单元（SFU）处理超越函数（如 sin,cos,exp,log）以及代数函数（如 sqrt,rsqrt 等）。
• 加载/存储（LD/ST）单元：服务于加载和存储指令的电路，与 TMA 引擎互补。
• 线程束调度器（Warp Schedulers）：每个 SM 包含调度器，为 32 个线程组成的组（CUDA 中称为 warps）发布指令。一个线程束调度器每周期可以发布一条线程束指令。
每个 SM 在物理上分为四个象限，每个象限容纳上述计算单元的一个子集。

这导出了以下见解：
📝 并行性（Parallelism）与并发性（Concurrency）
• 一个 SM 在给定的周期内最多可以同时发布来自四个线程束的指令（即在真正的并行执行中，每周期有 128 个线程）。
• 然而，一个 SM 可以容纳多达 2048 个并发线程（64 个线程束）。这些线程束常驻在 SM 中，并随着时间的推移被换入和换出调度，允许硬件隐藏内存/流水线延迟。
换句话说，指令并行性（在给定周期内有多少线程开始执行指令）限制为每 SM 128 个线程（4 条 32 宽度的线程束指令），而并发性（调度器中跟踪并有资格运行的线程数）则扩展到 2048 个线程。

光速（Speed of Light）与功率限制

既然我们购买 NVIDIA GPU 是为了计算，自然会问：性能上限是什么——GPU 的最大计算吞吐量是多少？这通常被称为“光速”（Speed of Light, SoL）性能：由芯片物理特性决定的上限。

根据数据类型的不同，有不同的性能上限。在 LLM 训练工作负载中，bfloat16 (bf16) 是近年来的主导格式，尽管 fp8 和 4 位格式变得越来越重要（对于推理，fp8 已相当标准）。
峰值吞吐量的计算公式为： perf=freq_clk_max∗num_tc∗flop_per_tc_per_clk
或者用文字描述：最大时钟频率 × 张量核心数量 × 每个张量核心每周期的浮点运算数（FLOPs）。

图 3：H100 SXM5 BF16 “光速”推导

📝 FLOP vs FLOPs vs FLOPS vs FLOP/s
• FLOP = 单次浮点运算。
• FLOP/s = 吞吐量单位：每秒浮点运算次数。
• FLOPs（小写 s）= FLOP 的复数（多次运算）。
• FLOPS（全大写）常被误用来表示吞吐量，但严格来说应读作“FLOPs”（复数）。将 FLOPS 用作 FLOP/s 是不严谨的！:)
我在上图中留下了一个提示：“光速”实际上并不是恒定的（我想这也是这个比喻失效的地方）。
在实践中，峰值吞吐量取决于实际时钟频率，而时钟频率会因功率限制（Power Throttling）或温度限制而波动。如果 GPU 时钟频率下降，有效的光速也会随之下降。

图 4：功率限制降低了时钟频率并拉低了有效的“光速”

📝 延伸阅读： Horace He 在他的博文 [3] 中更深入地探讨了这一现象。
硬件细节目前了解这些就足够了。接下来的重点将转向 CUDA 编程模型，然后我们会再次深入硬件底层，并最终上升到 CUDA C++ 层面。

CUDA 编程模型

CUDA 编程模型自然地映射到 GPU 硬件和内存层级结构上。其核心抽象包括：
• 线程 (thread)
• 线程束 (warp)（32 个线程）
• 线程块 (thread block)
• 线程块集群 (thread block cluster)
• 网格 (grid)（由线程块或集群组成）

图 5：CUDA 编程模型：线程、线程束、块、集群、网格
每个线程通过 gridDim、blockIdx、blockDim 和 threadIdx 等变量“意识到”自己在 CUDA 层次结构中的位置。在内部，这些变量存储在特殊寄存器中，并在内核启动时由 CUDA 运行时（runtime）初始化。
这些位置信息使得跨 GPU 分配任务变得简单。例如，假设我们要处理一张 1024×1024 的图像。我们可以将其划分为 32×32 的线程块，每个块包含 32×32 排列的线程。每个线程可以计算其全局坐标：

Plain Text
const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y;

并利用这些坐标从全局内存读取分配给它的像素（imagex），执行点对点操作，并将结果存回。

图 6：CUDA 内置变量：线程如何知道自己在哪里
如上图所示，在实践中我们大多使用 1D 或 2D 的网格/集群/块形状。但在内部，它们始终可以根据需要进行逻辑重组。
例如，如果 threadIdx.x 的范围是 0-1023（1024 个线程的 1D 块），我们可以将其拆分为 x = threadIdx.x % 32 和 y = threadIdx.x / 32，从而将其重塑为 32×32 的逻辑 2D 布局。
将 CUDA 模型连接回硬件，有一点现在应该很清楚了：一个线程块应包含至少 4 个线程束（即 128 个线程）。
为什么？ 线程块驻留在单个 SM 上。每个 SM 有 4 个线程束调度器——为了充分利用硬件，你不希望它们处于闲置状态。
📝 至少 4 个线程束的其他原因：
• 我们稍后会深入探讨，但在 Hopper 架构上，线程束组（warp-group，即 4 个线程束） 是 WGMMA（矩阵乘法）张量核心指令的执行单位。
• 此外，在使用持久化内核（persistent kernels）时，我们通常每个 SM 仅启动一个线程块，因此构建任务以保持所有线程束调度器忙碌至关重要。
带着 CUDA 编程模型的术语，我们可以继续深入探讨 GPU 的架构细节。

全局内存（GMEM）模型

让我们深入探讨 GMEM。如前所述，它是由多层 DRAM 堆叠而成，底部有一个逻辑层（HBM）。但 DRAM 到底是什么？

图 7：DRAM 单元内部：晶体管 + 电容器，字线（wordline） + 位线（bitline）
了解了单个位的存储方式后，让我们放大到整个存储矩阵。

图 8：GMEM 模型

📝 关于 HBM 的延伸阅读： 如果你想更深入地了解 HBM，我发现论文《Demystifying the Characteristics of High Bandwidth Memory for Real-Time Systems》[21] 非常有启发性。
我们可以得出结论：访问模式（access patterns）至关重要，这是由 DRAM 单元的物理特性决定的。

图 9：GMEM 访问模式的影响

Stephen Jones 的演讲《How CUDA Programming Works》[4] 非常值得一看。
如果我们示例中的矩阵是列优先（column-major*的，情况就会反转：列中的元素将连续存储，因此有效的选择是在内层循环遍历行，以避免 DRAM 惩罚。
所以，当人们说“GMEM 合并（coalescing）非常重要”时，指的就是：线程应访问连续的内存位置，以最小化触及的 DRAM 行数。

共享内存（SMEM）模型

共享内存（SMEM）具有与 GMEM 截然不同的特性。它由 SRAM 单元而非 DRAM 构建，这使得它在速度和容量的权衡上完全不同。
SRAM 单元的具体设计并不重要——只需知道存储一位信息需要多得多的晶体管。你可以自行搜索“SRAM cell”。
SMEM 组织为 32 个银行（banks），每个银行宽度为 32 位（4 字节）：

图 10：SMEM 模型
SMEM 可以在单个周期内提供来自所有 32 个银行的数据（128 字节）——但前提是必须遵守一条规则：
同一个线程束（warp）中的线程不得访问同一个银行（bank）内的不同地址。 否则，这些请求将被序列化，分多个周期执行。
这种情况被称为 银行冲突（bank conflict）。如果有 N个线程访问同一个银行中的不同地址，就会产生 N 路银行冲突（N-way bank conflict），该线程束的内存请求将需要 $N$ 个周期才能完成。
在最坏的情况下，所有 32 个线程都指向同一个银行中的不同地址，吞吐量将下降到原来的 1/32。
为了说明这一点，假设线程束大小（warp size）为 5。下方的两种访问模式将分别需要 3 个周期和 1 个周期来完成服务：

重要的是：如果一个线程束（warp）中的多个线程访问同一个银行（bank）内的相同地址，共享内存（SMEM）可以将该值广播（broadcast）或组播（multicast）给所有这些线程。
在下方的示例中，请求在单个周期内即可完成服务：
• 银行 1（Bank 1） 可以将一个值组播给 2 个线程。
• 银行 2（Bank 2） 可以将一个值组播给 3 个线程。

现在，来看硬件拼图的最后一块：L1 缓存。
这是一篇由 Axel 撰写的关于 SMEM 微基准测试（microbenchmarking）的优秀博文 [5]。

L1 模型

我们已经看到 L1 和 SMEM 共享相同的物理存储，但 L1 在该存储周围增加了一层由硬件管理的脚手架层（scaffolding layer）。
在高层级上，L1 缓存的逻辑流程如下：

1. 线程束（warp）发布一个内存请求（指向 SMEM 或 GMEM）。
2. 请求进入 MIO 流水线并被派遣至 LSUIN 路由器。
3. 路由器导向请求：SMEM 访问立即从数据数组（data array）中获得响应，而 GMEM 访问则进入标签比较（tag-comparison）阶段。
4. 在标签阶段，GMEM 的地址标签与目标集合（target set）中存储的标签进行对比，以确定数据是否驻留在 L1 中。
5. 命中（Hit）：请求直接从数据数组中获得服务（就像 SMEM 一样）。
6. 未命中（Miss）：请求传播至 L2（如有必要，甚至更远，直到 GMEM 或对等 GPU 内存）。当数据返回时，它会被缓存到 L1 中，替换（evicting）现有的一行，并并行地发送回发起请求的线程束。
   这就是我刚才描述的系统：
   
   图 13：L1 缓存模型
   让我们再深入一层，详细查看标签阶段（tag stage）和数据阶段（data stage）：

图 14：k 路组关联缓存组织的分解
当一个 GMEM（全局内存） 地址进入标签阶段时，命中/未命中（hit/miss）逻辑按如下方式展开：

1. 标签阶段接收 GMEM 地址。
2. 提取集合 ID 位（set id bits），并检查该集合中的所有缓存行（标签）。
3. 如果发现标签匹配（潜在的缓存命中）：
   ￮ 检查该行的有效性标志（validity flags）。
   ￮ 如果无效 → 视为缓存未命中（继续执行步骤 4）。
   ￮ 如果有效 → 从数据数组中提取所请求的分区（sectors），并交付至线程束的寄存器中。
4. 如果未发现匹配（缓存未命中），请求被路由至内存层级结构的其余部分（L2 及更高层级）。
   ￮ 当数据从 L2 返回时，它被存储在该集合中，并根据替换策略（例如伪 LRU 算法）驱逐（evicting）现有的某一行，同时并行地交付给发起请求的线程束。
   注意，L2 缓存与 L1 并没有太大区别，除了它是全局的（而非每个 SM 独立）、容量大得多（具有更高的关联度）、被划分为由横截（crossbar）连接的两个切片，并且支持更细致的持久化和缓存策略。
   至此，我们已经涵盖了理解后续章节所需的关键 GPU 硬件组件。
   📝 GPU 世代间的梯度：
   我之前提到过，理解 Hopper 是深入了解 NVIDIA GPU 未来和过去世代的绝佳基础。
   到目前为止，最大的世代跨越是从 Ampere → Hopper，引入了：
   • 分布式共享内存 (DSMEM)：在整个 GPC 的 SMEM 之间，实现加载、存储和原子操作的直接 SM 到 SM 通信。
   • TMA (张量内存加速器)：用于异步张量数据移动（GMEM ↔ SMEM, SMEM ↔ SMEM）的硬件单元。
   • 线程块集群 (Thread Block Clusters)：一种新的 CUDA 编程模型抽象，用于跨 SM 对块进行分组。
   • 异步事务屏障 (Asynchronous transaction barriers)：拆分式屏障，计数事务（字节数）而非仅仅是线程数。
   Ampere (例如 A100) 自身也引入了几个关键特性：
   • 张量核心 (Tensor Cores) 支持 tf32 和 bf16。
   • 异步拷贝 (GMEM → SMEM)，具有两种模式：绕过 L1 和访问 L1。
   • 异步屏障（在共享内存中由硬件加速）。
   • CUDA 任务图 (CUDA task graphs)：它是 PyTorch 中 CUDA graphs 的基础，减少了 CPU 启动和网格初始化开销。
   • 通过 CUDA 协作组 (CUDA Cooperative Groups) 暴露的线程束级规约指令（实现了单步、整数数据类型的线程束范围规约，无需 shuffle 模式）。