GPU 架构

1 GPU 架构要建立什么图景

把服务器机箱打开，GPU 通常不是悬在系统之外的“算力数字”，而是插在主板上的加速卡，或安装在专用底板上的加速模块。卡上有 GPU 芯片、显存芯片、供电和散热部件；主板上还有 CPU、系统内存、存储和网卡。训练程序必须让这些部件接力：存储提供样本，CPU 准备数据，经互连送到 GPU，GPU 执行计算，再把必要的结果交回主机。理解 GPU 架构，首先要能沿着这条路径说清楚数据在哪里、由谁处理、经过什么接口。

设想 1 台双路服务器，装有 256 GB 系统内存、1 块 NVMe 固态盘、1 张 100 Gb/s 网卡和 1 张显存为 24 GB 的 GPU 卡。这些容量属于不同的资源池。模型在显存中需要 28 GB，并不会因为服务器还有 200 GB 空闲内存就自然装得下。程序可以采用分批处理或显式搬移，但随之增加的是传输时间和管理工作。显存容量决定能否容纳当前工作集，显存带宽决定工作集能以多快的速度被反复访问，计算资源决定数据到位后能以多快的速度运算。

CPU 擅长承担分支复杂、任务变化快的控制工作；GPU 通过许多并行执行资源处理结构相近的大批数据。例如，对 100 万个像素分别做相似运算，通常比让 GPU 处理少量、依赖前序结果的操作更容易发挥硬件能力。这不是说 GPU 只能做简单算术，而是说它需要足够多可以交错推进的工作。训练中的大矩阵乘法往往提供了这种机会，逐个解析文本文件则未必能提供。

连接方式也是架构图的一部分。PCIe 4.0 x16 每个方向的理论数据速率约为 31.5 GB/s，应用可用速率还会受到协议、内存和传输方式影响；100 Gb/s 网络的原始速率折合 12.5 GB/s，尚未扣除通信开销。卡内显存带宽可能达到数百 GB/s，甚至更高。把主机到设备的搬运误当成卡内访问，就会把不同数量级的通道混为一谈。

本课建立的是“整机供数—芯片执行—显存供数”的图景。遇到慢任务时，应先定位哪个阶段拖长了总时间，再讨论该阶段的资源边界。芯片上有多少计算单元只是其中的线索，服务器能否稳定供电、持续散热、及时送入样本，同样会决定最终吞吐。后面的示例采用公开常见的 SM、Warp 等名称；具体数量随芯片变化，不能把教学配置当成所有 GPU 的统一规格。

2 SM、Warp 与执行资源

GPU 芯片内部通常包含许多计算单元。这里使用 SM，即流式多处理器，表示能够驻留线程块、调度线程组并提供计算和片上存储资源的单元。程序把大量线程组织成线程块，运行时把线程块分配给具备空闲资源的 SM。线程块一旦驻留，就要占用寄存器、共享内存和线程名额，直到这个块完成。未被分配的块继续等待，因此“启动了很多线程”不等于“所有线程同时执行”。

在采用相关术语的常见架构中，1 个 Warp 包含 32 个线程。调度器以 Warp 为重要的推进单位，让线程执行相应指令。线程各自处理不同的数据，却常常走相同的指令路径。例如，编号为 i 的线程负责数组的第 i 个元素，连续编号的线程访问连续元素。这样的组织便于硬件把相邻地址的请求合并，也使同一条算术指令能够服务许多数据。

需要把“属于同一个 Warp”与“在同一时刻完成所有指令”区分开。一次乘加可能等待输入，一次显存读取可能迟迟未返回，调度器便可能选择另一个已经就绪的 Warp。GPU 正是利用大量待推进工作填补等待时间。若所有驻留 Warp 都在等显存，单纯增加算术单元数量不会让这些读取提前完成；若就绪 Warp 很少，执行资源也可能没有足够任务可发。

SM 里并非只有同一种“核心”。它通常具有常规整数或浮点运算资源、访存相关资源以及特定矩阵运算资源，还配有寄存器文件和共享内存。不同指令走不同资源，吞吐也不同。某个精度下的矩阵运算峰值，不能拿来直接预测普通标量加法的速度。入门阶段只需知道资源有分工，暂不展开指令发射端口、流水线和矩阵指令的内部细节。

控制流也会改变执行效率。如果同一个 Warp 中有些线程满足条件，另一些不满足，硬件可能需要分阶段处理不同路径，只让相应线程参与。这里最容易形成的误解是“写了 if，速度必然减半”。实际影响取决于条件是否一致、每条路径有多长以及编译结果。对于边界判断，往往只有末尾少数线程受影响；对于每个元素都随机选择长分支的程序，代价可能明显得多。

观察架构图时，可以连续问 3 个问题：任务能分成多少彼此可推进的块，每个块占用多少驻留资源，已经驻留的 Warp 有多少能够马上工作。第 1 个问题关系到并行工作总量，第 2 个关系到装得下多少工作，第 3 个关系到执行资源是否真的忙起来。后面讨论占用率时，必须保留这 3 层含义，不能用某个百分比替代全部判断。

3 显存层次与带宽墙

从线程的视角看，数据并非只有“在 GPU 上”和“不在 GPU 上”两种状态。线程反复使用的局部值可能放在寄存器里；同一线程块协作复用的数据可以放在共享内存中；缓存帮助减少部分重复访问；容量更大的显存保存数组、模型参数和中间结果。它们在容量、访问代价和共享范围上不同。把热数据尽量留在合适的片上位置，才可能避免反复经过较慢、较窄的远端路径。

寄存器很快，却不是无限的。线程需要保留过多局部值时，编译器可能把部分值溢出到设备内存支持的局部内存中。这里的“局部”表示每个线程私有，不表示它一定就在芯片旁边。溢出可能增加加载和存储指令，也可能经过缓存。共享内存则由线程块显式组织使用，使用过多会减少能够同时驻留的块数；不正确的同步还会使程序给出错误结果或无法完成。

显存读取的效率还取决于线程访问的地址排列。若 32 个线程各读连续的 4 字节元素，硬件通常能用较少的内存事务覆盖请求；若它们分别跳到相隔很远的地址，就可能产生更多事务。程序看起来都读取了 128 字节有效数据，实际搬运量却可能不同。缓存命中、地址对齐和具体事务粒度也会影响结果，因此有效数据量和硬件传输字节数应分别理解。

用算术强度把计算与搬运放到同一个关系里。算术强度表示每搬运 1 字节数据完成多少次浮点运算，单位可以写成 FLOP/B。若某核函数从显存读入 2 个单精度数，再写出 1 个单精度数，共涉及约 12 字节，却只完成 2 次浮点运算，那么它的算术强度约为 0.167 FLOP/B。这里按显存访问的简化模型计数，缓存复用和额外读写会让真实情况发生变化。

设可持续显存带宽为 600 GB/s，这个核函数受带宽约束时，运算吞吐上界约为 600 × 0.167，即 100 GFLOP/s。即便相应运算类型的计算峰值达到 20 TFLOPS，也不能据此认为还有接近 200 倍的可用加速空间。数据只来得及支持约 0.1 TFLOPS，计算资源必须等待。这就是带宽墙：限制来自每次运算需要搬多少数据，以及通道每秒能送多少数据。

这个估算使用十进制 GB/s。容量里常见的 GiB、MiB 则分别基于 2 的幂，计算时应明确单位，避免把单位差异当成性能损失。理论带宽也不等于可持续带宽，时钟、访存模式和系统状态都会改变实际表现。入门时可以先测连续访问的基线，再拿同一设备、相同条件下的业务核函数作比较。

提高算术强度常常比增加线程更有价值。例如，让读入的矩阵小块参与多次计算，或把相邻操作合并以减少中间数组写回，都可能降低显存流量。但这些调整会增加寄存器或共享内存需求，还可能改变数值计算顺序。优化并不是把所有数据塞进最快的层次，而是让有限的片上资源换回足够多的数据复用，并确认结果仍然正确。

4 动手：对照一次核函数占用与吞吐

这次动手使用已有 GPU 开发环境，对照同一个逐元素核函数在不同线程块配置下的表现。示例用 CUDA 风格表达线程编号和启动方式，目的是把线程组织与硬件资源连接起来。没有对应设备时，也可以先按给定实验记录完成分析；不要把下面的示例数据写成自己机器的实测值。实际运行需要可用驱动、匹配的编译工具和足够的显存。

使用单精度数组 x、y，计算 y[i] = a × x[i] + y[i]。每个元素含 2 次浮点运算；融合乘加指令通常仍按 2 次计数。设置 N = 2^26，每个数组为 256 MiB，合计 512 MiB。读 x、读 y、写 y 的逻辑数据量约为 805 MB。数组应明显大于设备的末级缓存，否则测得的可能主要是缓存表现，而不是显存的持续供数能力。

__global__ void saxpy(const float* x, float* y, float a, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

// x、y 已在设备端分配并初始化；在同一流上记录计时事件。
int threads = 256;  // 分别改为 128、256、512
int blocks = (n + threads - 1) / threads;
saxpy<<<blocks, threads>>>(x, y, 2.0f, n);

先记录 GPU 型号、显存容量、驱动和编译版本，确认没有其他作业共享设备。保持数组大小、数据类型和编译选项不变，只修改每块线程数。完成预热后，使用设备事件包住多次核函数执行，再等待结束，按执行次数求平均值。每种配置做多轮，报告中位数和波动范围。主机直接测异步启动调用，往往只量到提交工作的时间，不能代表 GPU 完成计算的时间。

计时区间内不放主机到设备复制，也不放数组分配。重复更新 y 时应限制重复次数并重新初始化，避免数值不断放大，尤其不要用已经溢出的结果证明算法正确。正确性检查另做：可以把 x 初始化为 1、y 初始化为 0，单次执行后抽查或完整比对期望值 2；还要检查启动错误和同步时报告的执行错误。性能记录只有在结果正确时才有意义。

下面是教学记录，数字用于练习解释，不是特定产品承诺。占用率来自分析工具在相同测量范围内给出的实际活跃 Warp 统计，带宽按逻辑数据量除以核函数时间计算。

256 线程配置略快，但差异是否可靠，还要看波动范围。若多轮结果的变化已经达到 0.05 ms，就不能把 0.02 ms 的差距写成确定收益。进一步收集寄存器使用量、共享内存用量、显存读写吞吐和缓存信息，才有依据解释为什么占用率上升却只带来有限变化。工具分析可能扰动时间，因此正式计时和详细指标采集可以分开进行，保持输入和配置一致。

5 占用率与延迟隐藏取舍

占用率通常描述某个 SM 上活跃 Warp 数与该架构允许的最大活跃 Warp 数之比。它衡量的是驻留程度，并不直接表示浮点单元忙碌程度，也不表示任务完成百分比。理论占用率依据资源限制计算，实际占用率还受到工作量、执行阶段和块分布影响。分析工具中的具体口径应与报告中的名称对应，不能把理论上限和运行平均值放在同一列比较。

用 1 个简化 SM 建立直觉：最多支持 2048 个线程、64 个 Warp，具有 65536 个 32 位寄存器和 64 KiB 共享内存。假设每块 256 个线程、使用 8 KiB 共享内存，块数上限至少为 8。若每线程使用 32 个寄存器，寄存器允许驻留 8 个块，线程名额和共享内存也都允许 8 个块，因此最多活跃 64 个 Warp，理论占用率为 100%。

若每线程改为 64 个寄存器，1 个块便消耗 16384 个寄存器，寄存器总量只够 4 个块，对应 32 个 Warp，占用率降为 50%。若每线程需要 96 个寄存器，则只能容纳 2 个完整块，对应 16 个 Warp，占用率降为 25%。资源不能按小数个线程块使用，所以占用率往往呈台阶变化。真实芯片还存在分配粒度和其他限制，简化算式不能替代工具的架构专用计算。

更多 Warp 为什么有时会更快？当某些 Warp 等待读取结果时，其他 Warp 可能已有操作数，可以继续发出指令。这叫延迟隐藏。但它需要“其他工作已经就绪”，不是只要驻留就有效。所有 Warp 同时等待同一条拥塞通道时，再增加 Warp 可能只会增加队列长度。对于已经接近显存带宽上限的逐元素运算，进一步提高占用率通常不能按比例提高带宽。

反过来，为了提高占用率而强行减少寄存器，也可能更慢。原本留在寄存器中的值若被反复写回、读入，会增加内存流量；重新计算中间值也会增加指令。某个矩阵小块需要较多寄存器，却能复用数据很多次，降低占用率可能换来更高吞吐。判断时要比较完整的资源交换：少了多少驻留 Warp，增加或减少多少流量，最终完成相同任务用了多长时间。

可以按“足够并行—检查等待—比较时间”的顺序行动。先确认网格里有足够多的块覆盖整张卡，再看等待主要来自访存、依赖还是其他资源，然后对比配置变化后的任务时间。若小任务只有 8 个块，却运行在有许多 SM 的 GPU 上，即使单个繁忙 SM 的局部占用率很高，整张卡仍可能大部分闲置。整卡工作分布和单个 SM 驻留程度必须同时观察。

因此，占用率适合充当诊断线索和可解释的资源约束，不适合设为脱离业务的最高分目标。方案里写“占用率达到 100%”还不够，应该写清楚输入规模、核函数时间、有效带宽与正确性结果。只有这些证据相互一致，才能判断资源配置的改变到底是在隐藏等待，还是把代价挪到了别处。

6 常见坑：把峰值 TFLOPS 当业务吞吐

TFLOPS 描述每秒万亿次浮点运算，但必须同时说明数据精度、运算类型和计数条件。有些数字来自专用矩阵路径，有些对应常规浮点运算，还有些依赖特定稀疏结构。业务如果没有使用相应的数据格式、指令和数据组织，就无法直接兑现那个峰值。比较 2 张卡时，应先确保比较的是同一口径，再考虑软件是否能稳定调用对应资源。

训练吞吐则常用每秒样本数或每秒 token 数表达，还受到序列长度、批大小、模型结构和精度影响。2 个任务都叫“训练”，每个样本需要的运算量和搬运量也可能差得很远。把某张卡的峰值除以另一个任务的样本耗时，不会得到可靠的采购结论。即便已经估算每个样本的浮点运算量，也要考虑实际采用的算法和算子实现是否一致。

看 1 个端到端例子：每个训练步串行花 8 ms 读取与预处理、4 ms 传输、20 ms 做 GPU 计算、8 ms 做同步，总计 40 ms。若只让 GPU 计算阶段提速到 10 ms，总时间变为 30 ms，吞吐提高约 33%，而不是提高 100%。这是没有阶段重叠的教学模型；真实系统若存在流水重叠，需要根据时间线分析关键路径，不能把所有阶段耗时直接相加。

另一个坑是把监控里的“GPU 利用率 100%”解释为“全部计算资源达到峰值”。有些利用率指标主要表达采样窗口里是否有核函数在运行。长时间等待显存的核函数也可能让这个值很高。反过来，间歇启动许多短核函数时，采样窗口可能把启动间隙放大。指标名称相似，统计对象却不同，应先看工具定义，再判断它能回答什么问题。

机器形态还会改变持续性能。同一类芯片装在不同功耗预算、风道和插槽条件下，可能维持不同的工作频率。若机箱进风受阻，温度升高后频率下降，业务会表现为开始较快、持续运行后变慢。若 PCIe 插槽实际只协商到 x8，输入传输也可能落后于预期。记录卡的标称能力时，应同时记录服务器中的实际连接、功耗和温度状态。

容量错误同样会被误判成算力问题。批大小加大后触发显存不足，或者框架开始把部分数据放到主机内存，吞吐可能骤降，而不是平滑变化。分配失败、主机设备复制增多和核函数变慢是不同证据，处理方向也不同。先找时间线发生突变的位置，比直接更换更高峰值的 GPU 更能解释问题。

评价硬件时，至少保留固定输入条件下的端到端吞吐、关键阶段时间、显存峰值和运行稳定性。峰值计算能力可以帮助解释上界，却不能取代任务测量。采购或扩容结论应落在“这套配置完成这类任务需要多久”，而不是落在没有精度和负载条件的单个 TFLOPS 数字上。

7 案例：显存带宽打满导致算力空转

某台服务器使用 1 张 24 GB GPU 卡，模型参数能放入显存。任务每步包含多段逐元素操作：读入数组、缩放、加偏置，再应用简单变换，每段都写出中间数组。监控显示 GPU 利用率长期很高，但算术吞吐远低于该卡相应精度的峰值。工程师最初怀疑线程数不够，把每块线程数从 128 增至 512，端到端速度却几乎没有变化。

先确认输入数据已经在设备端，计时区间没有混入磁盘读取和主机复制；再观察 GPU 时间线，发现主要时间落在连续的逐元素核函数中，而不是 CPU 提交间隙。核函数规模足以覆盖整卡，实际占用率也不低。这些证据不能证明程序已经最优，但可以把“没有足够工作”和“主机供数太慢”暂时排到后面，集中检查卡内数据流量。

以其中的缩放相加核函数为例，N = 2^26，逻辑流量约 805 MB，实测耗时约 1.30 ms，有效带宽约 619 GB/s。同一设备上，较大数组的连续读写基线约为 650 GB/s。两者口径可比时，这说明该核函数已经接近这台机器在相似访问方式下的供数能力。这里的“打满”表示接近可持续基线，不要求达到说明书理论数字。

每次调用完成约 1.34 亿次浮点运算，除以 1.30 ms，约为 103 GFLOP/s。这个数与 0.167 FLOP/B 乘以 619 GB/s 的估算相符。较低的浮点吞吐并不矛盾：显存每秒送来的数据，只够支持这么多运算。此时计算资源空转更准确地说是部分执行机会被等待占据，不能据此推导每个计算单元的精确闲置比例。

接下来检查多段操作是否必须把中间值写回显存。若缩放、加偏置和简单变换之间没有需要跨线程协作的依赖，可以尝试合并核函数，让线程把中间值暂留在寄存器中。假设原先 3 段操作每段都读取并写回同一个数组，仅这条数组链就涉及 6N 次元素传输；理想合并后降为 2N 次。其他输入数组的流量仍需另算，不能把整个程序都按 3 倍收益承诺。

合并之后，重新核对输出误差、寄存器使用和执行时间。寄存器需求可能上升，占用率可能下降；但只要减少的显存流量更有价值，任务仍会变快。若数值计算顺序改变，还应按照业务允许的误差范围检查结果。优化成功的证据是相同输入得到合格输出，并且总时间减少，不是占用率曲线变得更好看。

如果合并并不适用，还可以检查数据布局、无效重复读取和不必要的中间精度转换。若真实算法已经没有可省流量，再考虑更高显存带宽的硬件，并用小规模实测验证收益。更多显存容量不会自动增加带宽，更多计算峰值也不会自动解决供数限制。最终记录应明确指出当前瓶颈是卡内显存路径，并说明证据成立的输入规模和访问条件。

8 把 GPU 资源边界写进训练方案

训练方案需要把“选哪张卡”展开成容量、吞吐和整机约束。先估计模型参数、梯度、优化器状态、激活值和临时工作区的显存需求。它们不能只按模型文件大小相加：训练保留的状态往往更多，激活值还随批大小与序列长度变化。以目标软件版本实际测量显存峰值，并给输入波动、工作区选择和后续调整保留明确余量，比只写“24 GB 应该够”更可执行。

容量预算还应覆盖主机侧。数据预取、解码缓存、多个加载进程都消耗系统内存；内存不足触发交换后，GPU 可能周期性等待。NVMe 的顺序读取标称值也不代表大量小文件的读取速度。若训练样本来自网络存储，100 Gb/s 网卡只是链路边界，存储服务的并发能力和数据组织仍会影响实际供数。方案应把样本从哪里来、如何准备和多久能送达写清楚。

吞吐预算从测量出的训练步时间出发。固定模型、精度、批大小与序列长度，给出预热后的样本吞吐和波动，再按时间线解释 CPU 准备、传输、GPU 计算和通信的贡献。对于主要受显存带宽限制的算子，记录有效带宽和复用机会；对于计算密集算子，记录实际使用的运算类型。这样才能说明增加计算资源或增加供数能力为何可能有效。

服务器形态要落到实际槽位和连接。确认卡的物理长度、厚度、供电接口、功率预算和风道是否匹配，也确认 PCIe 代际、通道宽度以及卡挂在哪个 CPU 下。双路服务器中，准备数据的 CPU 和 GPU 若隔着额外的处理器间路径，传输表现可能变化。入门阶段不必展开拓扑调优算法，但必须知道插在不同槽位可能不是同一种数据路径。

多卡训练还要单列卡间通信。不能因为每张卡都有较大显存带宽，就假设跨卡同步也同样快。通信可能经过专用互连、PCIe 或网络，容量也不会自动变成对单个程序透明的统一大显存。方案至少应说明卡数、连接方式、每步通信量和测得的同步时间。具体并行策略与通信优化留到后续课程展开。

验收条件应同时包括正确性、资源和时间。例如，固定输入范围下没有显存分配失败，运行指定时长后吞吐仍达到目标，温度和频率没有持续异常，输出满足误差要求。监控要能区分显存容量接近上限、显存带宽接近基线、主机内存不足和数据加载停顿。否则值班人员看到同样的“训练变慢”，无法选择对应动作。

最后写清退路：输入超出预算时减小微批大小，数据供给跟不上时检查加载与存储，持续降频时检查供电和散热，带宽受限时先评估流量与复用。每个动作都应有触发证据和复测条件。资源边界的作用，是让方案在真实机器上能够解释、实施和调整，而不是把某个标称硬件数字变成没有前提的性能保证。

9 从GPU 架构走向GPU 编程模型（H33）

现在可以把 GPU 看成受多类资源共同约束的执行系统。主机准备工作，互连传送数据，线程块分配到 SM，Warp 在等待与就绪之间切换，寄存器和共享内存决定部分驻留边界，显存持续为计算供数。任何阶段都可能限制总吞吐，但并非所有限制都能靠增加线程或更换更高计算峰值的芯片解决。

面对新核函数，先画数据路径，再列资源账本。输入从哪里读，结果写到哪里，每个输出大致需要多少字节、多少运算；任务能拆成多少块，每块需要多少线程、寄存器和共享内存。即使暂时不能精确预测运行时间，这些问题也能帮助排除不合理假设。例如，工作集放不下是容量问题，持续读取很慢可能是访问方式或带宽问题，块数太少则可能是并行工作不足。

还应分清程序员指定的组织与硬件实际的安排。线程块大小是启动配置，Warp 是硬件调度相关的组织，实际驻留块数则由资源与架构共同决定。程序写下的块编号通常不能当成 SM 编号使用，也不能依赖不同线程块按照编号先后执行。让块之间形成隐含的等待关系，既可能不正确，也可能导致某些任务永远无法推进。

H33 将把这些硬件直觉落到编程模型上：如何定义核函数，如何计算线程索引，如何组织线程块与网格，何时需要同步，如何安排设备内存的分配与复制。本课没有展开 Warp 内部的精细调度、指令流水细节或专用矩阵指令编排，因为先掌握工作组织、数据位置和测量范围，才能正确理解那些更细的机制。

进入下一课前，可以用自己的话解释 3 个判断：占用率提高为什么未必加速，显存剩余很多为什么仍可能受带宽限制，GPU 利用率很高为什么端到端吞吐仍可能不理想。回答时要引用明确的资源或时间范围，而不是只说“并行度不够”或“硬件不行”。能够提出可测量的解释，就已经具备从架构走向程序的基础。

放进方案里

本训练方案使用 1 台双路服务器，配置 256 GB 系统内存、NVMe 本地数据盘、100 Gb/s 网卡和 1 张 24 GB 显存的 GPU 卡。上线前确认实际 PCIe 连接达到预期代际与通道宽度，核对供电和风道。模型、梯度、优化器状态、激活值及临时工作区的显存峰值，以目标软件版本和最大输入配置实测，要求峰值不超过 20 GB，余量用于输入波动与工作区变化；超限时先减小微批大小并重新验收。

性能验收固定模型版本、精度、批大小、序列长度及数据来源，记录预热后的每秒样本数、训练步中位耗时和波动范围。时间线分别覆盖数据准备、主机到设备传输、GPU 计算与同步，注明是否存在重叠。逐元素热点同时记录有效显存带宽、寄存器用量和占用率；占用率仅用于解释驻留限制，吞吐目标以业务实测为准，不以峰值 TFLOPS 推算交付能力。

若显存带宽接近同机连续访问基线而计算吞吐偏低，优先检查中间数组写回、重复读取和可行的核函数合并，并在修改后复核输出误差与总时间。若 GPU 出现等待间隙，则检查数据加载、系统内存压力、传输与同步。连续运行 2 小时，确认吞吐达标、无显存分配失败，且没有持续异常降频；保存温度、功耗、频率及关键阶段时间。输入规模或软件版本变化后重跑基线，避免沿用已经失效的容量与速度结论。

验收附件保留线程块配置、编译器报告的每线程寄存器数及溢出字节数，防止升级后新增局部内存访问却仍沿用旧结论。主机到设备传输另用固定大小缓冲区测量，注明是否使用页锁定内存及其所在的 CPU 节点。若协商链路或传输带宽低于基线，先核查槽位与跨处理器路径；恢复后重新测训练步，确认改善落到端到端吞吐。

10 小结：下一步GPU 编程模型（H33）

GPU 的能力来自大量可交错推进的工作，也受显存、片上资源与整机供数约束。SM、Warp 和占用率解释工作如何驻留与推进，算术强度和带宽解释数据是否来得及供应，端到端时间线解释这些能力最终如何变成训练吞吐。

下一课 H33 将编写和组织 GPU 程序，把线程索引、线程块、同步和内存搬运与本课的硬件图景对应起来。带着“数据在哪里、谁在等待、测量覆盖哪里”这 3 个问题继续学习。