07 GPU、资源估算与 FlashAttention
课程锚点：GPUs Slides5–68；FlashAttention PDF Pages10–41；范围板书的显存/运算数/能耗。课件在某些页把FLOPs/FLOPS写法混用，答题先写单位。
7.1 GPU 为什么适合这些任务
CPU倾向复杂控制与低延迟任务，GPU用大量并行算术单元追求吞吐。GPU包含多个SM；thread组成block，block调度到SM；NVIDIA常见warp是连续编号32个thread。warp内分支分歧会让部分lane在执行某分支时被屏蔽，降低有效利用。
CUDA是host/device分工与编程生态：CPU准备数据、launch kernel，GPU执行。grid→block→thread，常见一维索引i=blockIdx.x*blockDim.x+threadIdx.x，末尾需i<n边界判断。
PyTorch用库kernel表达算子；CUDA更直接控制thread、内存和同步；Triton让你主要按tile编程，但仍受硬件带宽、片上容量、调度与布局约束。Tensor core擅长受支持dtype和形状的矩阵乘加，不意味着所有非矩阵算子也同速。
7.2 内存层次与三个不同瓶颈
寄存器最贴近thread；shared memory是片上block可协作使用的空间；L1/L2为缓存层级；HBM/显存容量大但搬运代价更高。寄存器溢出会spill，tile太大可能降低occupancy。
- compute-bound：运算吞吐先饱和；减少FLOPs或更好用tensor core有价值。
- memory-bandwidth-bound：搬运数据速度先饱和；降低bytes、提高复用有价值。
- launch/调度/通信/CPU瓶颈：小kernel、网络或数据流水线限制；不能全塞进前两类。
7.3 FLOPs 从矩阵尺寸算
[m,k]@[k,n]→[m,n]，每个输出k次乘、k-1次加，约2mkn FLOPs；若题目按一次FMA计1，需说明口径。
每token经过一个无biasD→F→D FFN约4DF FLOPs；SwiGLU三次线性约6DF，暂忽略激活。MHA单层序列投影8BTD²，attention两个matmul约4BT²D，因果三角如有效省算可约减半attention项；实际kernel是否跳过上三角另看实现。
粗略dense训练公式C≈6 P D_tok，P为主要非embedding dense参数数、D_tok为训练token数。动机：前向约2P/token，反向对输入和参数两类梯度再约4P/token。长序列attention、embedding/head、重计算、稀疏路由和通信会使此估计偏离；不要把6PD当所有模型精确恒等式。MoE要用实际活跃路径估compute，全部参数估主要权重存储。
7.4 显存计算要分账
推理：权重 + KV + 临时激活/工作区 + 通信/图等运行时预留 + 其他开销。
权重本体P×bits/8。8B BF16=16×10^9 bytes≈14.90 GiB；27B BF16=54×10^9 bytes≈50.29 GiB。8B 4-bit理想打包=4×10^9 bytes，但量化scale/zero point、分组、未量化层、临时反量化及KV另计。
训练还需梯度、优化器状态和backward保存的激活。典型混合精度Adam记账：2bytes模型权重+2bytes梯度+4bytes FP32主权重+8bytes一二阶矩=16bytes/参数。这是一种常用假设；梯度也可能FP32，某些实现无独立master weight，低精度优化器、分片和offload也会改变。题目先列假设再算。
有足够显存装下权重不表示可训练，也不表示能服务目标长度/并发；多卡的显存求和不是自动形成一整块连续显存，仍需并行切分策略。
7.5 Roofline 模型与课件勘误
Arithmetic intensity I=C/Q，单位FLOP/byte，Q为指定内存层级的流量。性能上界：
达到compute roof还需I≥R_peak/BW。例如100TFLOP/s、1TB/s，拐点100FLOP/byte。I=10时带宽屋顶约10TFLOP/s。这里使用十进制单位且是示例硬件，不是某型号实测承诺。
ReLU向量每元素1次读取+1次写回。若把比较按1 operation计：FP32每元素8bytes，I=1/8=.125 op/byte；FP16每元素4bytes，I=.25。课件正文“8/4 FLOP/byte”把倒数写反，旁边小表为正确值。严格说比较不一定计标准浮点加乘FLOP，所以最好说明operation口径。
低精度在相同操作下减少bytes，有可能提高I；硬件peak也可能变化，不能只改I却默认所有性能参数不变。量化做得好不必然“SM装满”，矩阵形状、batch、反量化和KV带宽都可能限制。
7.6 Tiling、coalescing、重计算
矩阵乘法朴素实现反复从global读同一输入；分块后把输入tile搬到shared/register，供多个乘加复用，增加每byte带来的运算。T×T简单tile复用分析常给HBM读取约T倍改善，但依赖布局和调度。
coalescing让相邻lane访问相邻/适当对齐地址，减少无效事务；与共享缓存复用是不同概念。算子fusion把多个操作合进kernel，避免中间张量反复写回HBM。重计算用一些额外FLOPs换掉昂贵的保存/读回；activation checkpointing在backward重跑一部分forward，主要省激活显存。
tile太大可能超shared/register限制，太小又复用不足。边界tile有masked-off元素；尺寸稍微增大会多出一整行/列tile，性能不一定随FLOPs平滑增长。
课件wave quantization例：tile256×128，1792²输出需要7×14=98 tiles；1793²需要8×15=120 tiles。简化每SM同驻1tile、108SM时，从1波变2波，最后一波仅12个tile。真实硬件可有每SM多block，具体拐点依kernel资源而定；wave quantization不是权重量化。
7.7 FlashAttention 的问题起点
朴素attention：把S=QKᵀ/√d写HBM→读S算softmax→把A写HBM→再读A乘V。T²中间矩阵带来大量HBM流量。FlashAttention按块在SRAM计算scores，在线维护归一化和加权输出，避免完整S/A落HBM。FlashAttention原论文。
标准精确FlashAttention不是“少看一部分token”的近似方法；浮点归约顺序改变可能有小数值差别。它不把标准dense attention的二次算术复杂度变线性。训练反向可借保存的统计重新计算块内概率，省内存而增加部分算术。
7.8 Online softmax 要能从头推
对某个query，一部分scores已经处理，维护最大值m、归一化和ℓ与未除ℓ的加权分子u：
ℓ=Σ_old exp(s-m)，u=Σ_old exp(s-m)v。
新块scores s_b到达，先取m'=max(m,max(s_b))。旧项需要统一换基到新最大值：
如果m改变却不重缩放旧ℓ/u，结果就错了。初值可取m=-∞、ℓ=0、u=0，但第一个有效块和全mask块要避免-∞减-∞。
例scores[0,ln2,ln3]，V=[1,2,4]，分成前2项/后1项。整行softmax为[1/6,2/6,3/6]，输出17/6≈2.833333。先处理前2项：m=ln2、ℓ=1.5、u=2.5；加入最后一项后m'=ln3，旧乘2/3，得到ℓ'=2、u'=17/3，输出仍17/6。
这段推导说明“softmax分母要看全行”不阻止分块：所需的全局信息可以用少量可合并统计表示。
7.9 FLOPs 到 kWh
若C=10^18 FLOPs，有效吞吐R=100TFLOP/s，则t=10,000s≈2.7778h。若整机平均功率W=1kW，电量约2.7778kWh。若数据中心PUE=1.2，估算总电量再乘1.2。
R要用有效吞吐，不能直接用宣传peak；W要明确是GPU还是整机、单卡还是集群。也可用E=C×每FLOP能耗估算，但每FLOP能耗依硬件、精度、利用率和访存而变。FLOPs与kWh之间没有无需条件的固定换算。
