大小
右图:使用什么并行策略的组合需要根据批次大小决定 横轴:单张显卡分到的批次大小 纵轴:计算时间与通信时间的比值
蓝线:仅FSDP,当批次大小>850时,才会计算>通信限制 绿线:FSDP+MP,实际大模型中训练最优解 橙线:MP only(张量并行TP+流水线并行PP),因为张量并行完每一层算完都要通信,batch size变大,计算量变大,通信量也等比例变大。两者绑定
大小
右图:使用什么并行策略的组合需要根据批次大小决定 横轴:单张显卡分到的批次大小 纵轴:计算时间与通信时间的比值
蓝线:仅FSDP,当批次大小>850时,才会计算>通信限制 绿线:FSDP+MP,实际大模型中训练最优解 橙线:MP only(张量并行TP+流水线并行PP),因为张量并行完每一层算完都要通信,batch size变大,计算量变大,通信量也等比例变大。两者绑定
以分布在64个加速器上的75亿参数模型为例,其内存占用极其庞大,且总内存随GPU数量线性增长,这显然不可接受。
ZERO-1: 只切分优化器状态 不同GPU上对应的数据用A,B,C,D区分 A0-A3 :GPU0上数据的完整梯度张量被切分成4段(同一层,张量的前后)
重要捋清——ZERO 1工作原理 GPU0用自己那部分数据(样本A)进行前向和反向计算,算出完整的梯度张量A。然后将这个完整梯度切分成4份(A0-A3)。因为GPU0只负责维护1/4的优化器状态(对应第0块参数),所以它只留下A0,把A1-A3通过 Reduce-Scatter 发送给其他GPU,同时接收其他GPU发来的B0、C0、D0。GPU0将A0+B0+C0+D0相加,得到自己负责那1/4参数的全局梯度碎片。接着,GPU0用这个梯度碎片加上自己本地的1/4优化器状态,只更新自己负责的那1/4参数。最后,所有GPU通过 All-Gather 互相交换自己更新好的参数碎片,拼凑出完整的最新参数,用于下一轮前向计算。gather是把每部分优化器计算出的全局最新参数gather到每个GPU上。
所有组件
优化器状态、梯度、模型参数 ZeRO-2:每个人手里拿着完整的图纸(参数),负责拼凑属于自己的那一小块拼图(梯度)。拼好自己这块后,把不需要的边角料扔掉。
ZERO-3 计算前借全部,算完立刻扔,下次用再借
模型参数
线性层的W,b,训练完成后就固定。而激活值是前向传播的输出结果,反向传播计算完后就释放。
手动8.1毫秒,PyTorch 1.1毫秒,CUDA 1.8毫秒,而torch.compile仅需 1.47毫秒。
不同类别的使用场景: (1)manual: 几乎不使用 (2)pytorch: 日常开发、模型训练和推理的默认首选。常用的有优化就直接写 (3) CUDA: 需要极致性能,且 PyTorch 没有对应优化算子,或者需要极特殊的硬件控制(如 FlashAttention v1 的早期实现) (4)triton time: 需要自定义融合算子(Custom Fused Kernel),但不想写复杂的 CUDA C++ 时 (5) torch.compile: 自动捕获计算图进行算子融合。如果你不想手写 Triton,直接在你的 PyTorch 模型外面套一层 torch.compile(model),往往就能获得接近 Triton 手写的性能。
tl.store
x= t1.load加载数据-》store, 一次加载一次写回。
数据写回HBM “内存融合体现在 Triton 内核中 tl.load 到 tl.store 之间的所有计算代码(手动档)。在这段区间内,数据被加载到寄存器后,连续完成 x^3、加法、乘法、exp、除法、tanh 等所有数学运算,中间没有发生任何 HBM 读写,最后只通过一次 tl.store 写回最终结果。
如果对比纯 PyTorch 的实现,差异极其明显。PyTorch 的逐元素操作会为每一个数学运算符启动一个独立的 CUDA 内核。例如 x**3 是一个内核,x + term 是另一个内核。这导致数据在 HBM 和 SM 之间被反复搬运了十几次,受限于内存墙,性能极差。
而 Triton 的这段代码,通过内核融合,把十几次 HBM 读写压缩成了 1次读和1次写。数据一次性加载到 SM,在寄存器中算完再写回。这彻底消除了中间结果的显存开销和内核启动开销,这就是它比手写 PyTorch 快几倍甚至十倍的根本原因。”
实现gelu
Python 封装函数(triton_gelu)负责检查 is_cuda 和 is_contiguous、分配输出张量、计算网格大小并启动内核;Triton 内核(triton_gelu_kernel)则负责实际的 GPU 计算。
pid
每个pid分配一段block,只计算block内的数据就好
最后跨线程归约。
同一个块内的线程间通信(开销很小),而手写CUDA可能需要全局同步。必须手写 __shfl_xor_sync、__syncthreads(),极易出错,需要跨线程块通信。
为什么我们的手动实现会比编写cuda版本慢呢,cuda版本会将数据从GPU发送回CPU然后x将驻留在GPU中。而手动版本中我们在GPU中分配它并向CUDA实现一样处理,但它上不会一直待在流多处理器中。所以一旦我们执行x平方,那是一个CUDA内核。这个乘法操作将从全局内存读取向量到流多处理器,执行计算然后写回。所以这全是DRAM到流多处理器的通信成本,而不是CPU到GPU的通信成本。当然,如果是CPU设备,那么除了DRAM传输成本还有CPU传输成本
手动实现:一次实现就一次启动内核+DRAM读写,有几个基础操作就几次。而CUDA实现将基础计算融合,总共只需实现一次DRAM。
(description: str, run: Callable, num_warmups: int = 1, num_trials: int = 3):
description: 定义任务的描述, sleep run: 把要执行的动作包装成一个带固定参数的函数代入 预热次数 num_trails: 正式测量的重复次数
if torch.cuda.is_available(): torch.cuda.synchronize() # 等待 CUDA 线程 完成同步
如果不加这两行,可能CPU开始执行run()后,异步往后执行计时就结束,而实际上GPU还在算。
GPU寄存器由于其高速访问特性,可以在较低操作强度下达到高吞吐量,但容量有限
寄存器文件离GPU内核最近,访问速度最快,但是容量最小。让分块数据尽可能复用。GPU带宽比CPU大,但是单词访问延迟更高
共享内存(
SM内部,和L1共用一块SRAM。block共享,速度和L1一样快,但是block用完就要销毁。
3)MoE相较于Dense模型的区别是什么,两种架构运用场景是什么?
核心区别在于计算方式和参数使用效率。
(1)Dense计算方式为全参数激活,MoE是稀疏化计算。 (2)Dense 总参数量 ≈ 激活参数量,计算量(FLOPs)与参数量成正比;而MoE参数量与计算量解耦,可以使用参数量更大的模型。推理时每token的计算量更少,速度更快。
(3)适用场景:Dense模型:适用于资源受限或对部署简单性、性能可预测性要求高的场景,如端侧设备、小型应用,以及对模型结构简单性和可解释性有要求的场景;
MoE与Dense的核心区别在于稀疏激活。Dense模型全参数参与计算,计算成本与参数规模强相关;MoE通过路由只激活部分专家,实现了总参数量与计算量的解耦,能以更低的计算成本获得更大的模型容量。Dense模型适合资源受限、追求部署简单的场景;MoE则适合大规模、高吞吐、追求极致性能的前沿应用。
)目前 标准Attention 的改进以及其对应的原理是什么?
主线一:优化KV Cache,降低推理内存占用 这条路线通过减少Key和Value矩阵的存储量来提升推理效率。
Multi-Query Attention (MQA):让所有查询头(Query Head)共享同一组Key和Value头。这极大地压缩了KV Cache,但可能造成模型质量下降。
Grouped-Query Attention (GQA):这是MQA和标准多头注意力(MHA)的折中。将查询头分成G组,每组共享一组Key和Value头。GQA在减少KV Cache和保持模型质量之间取得了很好的平衡,是当前许多主流大模型(如Llama 3)的默认选择。
主线二:优化计算过程,提升硬件利用率。 lashAttention:其核心洞察是标准Attention的瓶颈并非计算量(FLOPs),而是显存访问(Memory Access)。FlashAttention通过分块(Tiling) 和重计算(Recomputation) 技术,将Attention计算过程拆分成小块,使其能放入GPU的快速SRAM中,从而大幅减少对慢速HBM的读写次数
为解决这些问题,Kimi团队提出了几种工程优化方法: 混合并行设计:合理安排不同并行方式,减少通信开销; 截断重要性采样修正:防止少数专家过载; 自适应并行调度器:根据实时指标GPU利用率、内存、步长时间自动调整tensor、pipeline、专家、sequence并行策略。
通信开销:GPU要把每个token给每个专家计算,得出打分最高的top-k后才把token真正交过去计算。all-to-all通信。直接切断并行会让专家缺少重要信息。
混合并行策略:含义:不只用一种并行策略,而是根据硬件拓扑(比如NVLink、InfiniBand带宽)和模型结构,把张量并行(TP)、流水线并行(PP)、数据并行(DP)、专家并行(EP)和序列并行(SP)组合起来。
自适应并行:打破以往静态的并行策略,根据动态指标实时调整并行策略。
Z = ∑ i = 1 V exp ( z i )
Z是softmax的分母项,L_cross是体现不同token的区分程度,但是softmax越大越不稳定,因此Lz_loss加了一个约束项、下面的z(x)是笔误,应该改成Z接近1。 在训练过程中softmax的分母 Z会一起被加到损失中,训练平稳softmax的分母,使整个训练更平稳。
softmax分分母Z接近1,softmax值就接近e^zi原值了,这样训练会更平稳是吗。如果小于1,会放大分子,大于1又会让softmax值很小
深度还是宽度
深:网络层数 宽:隐藏层维度
提问:残差连接的作用
没有残差时:第一:每层都必须直接学完整映射y=H(x),H(x)为理想输出有残差时y=H(x)-x,只需要学和理想值之间的差值。第二:当某一层学习不好时,深层网络里的信息更容易丢失,有残差信息流更稳定。第三:没有残差:多层导数不断相乘,容易出现梯度消失或梯度爆炸
PyTorch
pytorch随机种子:张量、神经网络、pytorch numpy: 数据集的随机划分 random: pytorch的random模块,Python 原生的随机操作,如 random.shuffle()(打乱列表)、random.choice() 一般情况下三个都要设置
进行标准高斯分布初始化参数
对W设定一个初始值,服从标准高斯分布初始化权重0-1之间。神经网络每层都由随机变量加权求和,这些和由方差决定。 y=sum D(xi*wi),var(y)=Dvar(x)var(w),若x的方差都相等, 后两项为方差1,则var(y)=D,层数多了会产生数值爆炸。若权重初始化值很小,又会梯度消失。
2 × B × D × D。
L是标量,但对向量求导后,导数会变成和张量一样的维度。(BK)*(KD)=(BD),总计算次数为2KBD
∂ L ∂ W 2 = h 1 T ∂ L
当W2发生微小变化时,Loss怎么改变,即求alpha L/alpha W2. 又因为W2和L有关,所以使用链式法则。
回顾一下前向传播 FLOPs 的计算,第一层 x @ w1 需要 2 * B * D * D FLOPs;第二层 h1 @ w2 需要 2 * B * D * K FLOPs,总计为 num_forward_flops = (2 * B * D * D) + (2 * B * D * K)
正常的矩阵运算:C(MP)=A(MN)B(NP),对于C中的每个元素,每个元素计算次数为N次乘法,N-1次加法,共MP个元素,所以总计算量2NMP。
对于两层非线性:x=(BD),w1(DD). h1(BD), w2 =(DK). 第一次计算量2BDD,第二层2BDK,总计算量为两层相加
线性模型,并将其扩展为一个简单的两层的线性模型:
线性模型y = xw, 为了增加模型的表达能力,增加线性层层数。h1 = xw1, h2 = h1*w2, 拿输出和w2比得到loss
nn.Parameter
torch.randn是实现一个维度为(input,output)维度的张量,nn.Parameter把该张量参数注册到参数列表,在梯度变化时自动更新
y.transpose(-2, -1)
把y的倒数第2个维度和倒数第1个维度转置
ZeRO 显存优化
让每张卡只存一部分优化器状态、梯度和参数。优化器状态:每个优化器额外存储:一阶矩估计、二阶矩估计
流水线并行(PP
把不同层给不同设备训练。加入transformer有八层,串行时每个算4层。算前4层时,后4层在等待。改成并行计算吞吐能大幅度提升。
张量并行(TP)
把单个矩阵切碎给多卡用
以及 Transformer 如何用 Self-Attention 解决它们
(1) attention一次性把所有向量矩阵(token)输入模型,实现并行计算。self-attention的核心是矩阵运算,把所有词存入矩阵一次性计算。
(2)每个token之间是直接连接的,可以直接(一步)计算每两个token的注意力分数,而RNN不同距离的token存在远近距离,因此会遗忘。
(头数)x 128(head_dim
头*head_dim= 隐藏层维度 有多个大小为隐藏层维度的隐藏层,这里隐藏层数是32
LayerNorm
两个Layernorm的作用: (1) 稳定Q,K,V的分布 (2) 稳定FFN的输入
干净
这里公式只给出了前向传播,
post-Norm 前向传播时梯度要先被LN求导。
反向传播时\alpha L/y*(1+\alpha LN\x), L对y的求导路径是干净的。
结果只依赖于两个 token 的相对位置差
sin: x+PE, 会改变原始的数值表达,PE是远近的绝对位置 RoPE: R*X, 选择 相对位置n-m,同样会有很多维度表示变化频率不同的远近。因为有多个维度,所以不用担心360度重复的问题。
iii 是维度索引
假如有维度0,维度1
维度0:变化快,区分相邻单词 维度1:变化慢,区分远距离单词
AllReduce
W1 按列切分:每个 GPU 拿到不同的中间特征(升维后的结果)。
W2 按行切分:每个 GPU 用自己那部分继续算最终输出。
这样中间完全不需要通信,只在最后把各 GPU 算出的最终结果加起来(做一次 AllReduce)就行。
M
M=1000,000
d×d
Q,K,V权重矩阵维度:d_model * d_model Q,K,V维度: (N,d) 序列长度*隐藏层维度
Online Softmax 算法将其合并为一遍扫描,FlashAttention 正是基于此实现了 Attention 的高效融合
online softmax: 两遍扫描变成一遍 flashattention: 在GPU上的高效硬件实现
?
Tokenizer 会直接决定序列长度,因此影响 attention 计算量(序列长度为L,则计算量为L^2)、KV cache、显存和有效上下文长度(切分的太碎计算成本高);同时它也决定 embedding 的参数共享方式(词拆分成token的形式影响学习能力),从而影响稀有词、新词、多语言、数字和代码的学习效率。切得太碎会增加计算成本,切得太粗会产生大量低频 token。因此 tokenizer 虽然不直接决定模型上限,但会显著影响 LLM 的训练效率、推理成本和最终任务表现。
?
一个好的 tokenizer 并不是单纯追求最大压缩率,而是要在序列长度、词表大小、覆盖能力和语义稳定性之间做 trade-off。粒度太细会导致 token 数过多,增加 Transformer 的计算和 KV cache 开销;粒度太粗则会导致词表爆炸、稀有 token 数据不足以及新词泛化能力下降。因此主流 tokenizer 通常采用 subword 方法,让高频词或高频片段成为长 token,而低频词继续拆分,从而兼顾压缩效率和泛化能力。
并行归约
并行归约是指:在并行计算中,将分布在 N 个不同处理单元(线程/核心) 上的大量数据,通过某个满足结合律(Associative)和交换律(Commutative)的二元运算符(如 +、*、max、min),“多对一”地合并成最终的一个(或一小批)结果的过程。
串行思维:for i in range(N): sum += arr[i](需要 N 步,依次等待)。
并行思维:N 个人各拿一个数,两两配对相加,不断合并,直到剩下最后一个人拿着总和。 2^k=N,经过k轮运算,每轮处理一半的数据,处理完N个样本。复杂度为O(log(N))
并行归约的本质是利用树形拓扑将对一维数组的累加复杂度从 O(N) 降到 O(log N)
Softmax + 交叉熵的简洁梯度
要求掌握这里结论的推导过程
LogSumExp
推导logsumupexp和logsoftmax的区别 logsum就是原softmax的分母部分
LSE(z)=m+logj∑ezj−m
zj=zj-m+m, LSE=logsum(e^(zj-m+m)) LSE=loge^msum(e^(zj-m)) = m +logsum(e^(zj-m))
dtype 容差
dtyper容差:不同精度允许参考误差不同 极端验证:测大值、小数、极端情况是否正确
exp 溢出、除零、消减或长归约误差
溢出:exp指数超出float 除零错误 消减误差:大数相减,小精度丧失 长规约误差:多个小数相加,后面的小数省略产生误差。
广播、转置或非连续 stride
广播:多个变量访问一块内存 转置:矩阵转置 stride是矩阵的访问步长,矩阵元素在gpu中连续存储,用stride控制如何访问(如访问第几行第n个元素),当stride和尺寸不一致,访问不连续
理论最小访存量
这个问题分析:一个 kernel 到底是算力瓶颈还是内存瓶颈。 假设矩阵A(MK),B(KN),C(MN),则共需要读写MN+MK+KN次。 好的kernel希望一次加载,多次复用,计算强度AI=FLOPS/memory. 优化后一个tile,可以供多个计算步骤使用。
能否分块?
分块 tileling 不能分块的情况: (1). 前面一个对后面一个存在强依赖关系,前面结果是后面的输入 (2) 数据量太小,分块后通信开销比不分块大 (3)计算需要全局信息。分块后前面信息丢失,无法计算完整结果
归约
矩阵乘法:A=MK,B=KN, C=AB=M*N,则维度K被规约
分布式均值
每个GPU计算的样本数量不一定相同,每个GPU有一个均值计算所有GPU就要分布式均值。将一个batch切分,放到多个GPU上并行运行。accuracy,loss等都需要所有batch计算后的整体均值
Online Softmax 的分块合并
把大块attention切分成小块attention减少显存开销。 原来大块attention需要保存S^2矩阵在显存中,计算小块attention,不用保存完整attention.用完就丢,节省显存
Dh\sqrt{D_h}Dh
每个head的Q,K维度为S*D, attention score =QK^T,每个token和哪个token有关。
QK^T得到矩阵维度为S*S,但每个元素的长度都是Dh,由Dh次乘法加法得到,用于归一化
M
M是因果掩码/上三角掩码,不允许模型看到之后的值,把未来的值都用-无穷掩盖。
X
X乘权重W,维度是W(H,H),乘完后Q,K,V维度为BSH。 1. 因为Transformer使用多头注意力机制,有N个头,每个都维度为D。所以BSND 2. 因为attention需要计算QK^T, 希望每个head单独计算,因此希望维度为BNSD。(每个head单独计算的Q,K维度为S*D).因为一次batch和head.都希望作为并行维度,所以前面要再乘B,N.
使用transpose可以交换指定轴的维度transpose(1,2),交换1,2维维度
反向传播中的 FP16 小梯度可能下溢为 0
混合精度训练(FP16/BF16)中的 Loss Scaling(损失缩放)技术,反向传播时FP16能表示的最小为6*10^-8,若最小梯度为10^-8,则下溢为0。
把梯度先放大,反向传播后再除S恢复正确数值。
不能被二进制浮点精确表
无限循环小数只能被二进制有限表示。当绝对数值越大,精度越差。 如1.m2^e, m是尾数,e是指数,如1.2310^4.e越大,数值范围越大,尾数固定,由于指数导致的范围很大,可表示的精度就越小
7 LayerNorm
归一化层,LayerNorm是归一化公式
7.8 RMSNorm
Layernorm的简化版本,RMSEnorm计算更快 LayerNorm:需要: 求mean 求variance 标准化
RMSNorm:只需要: 求平方和 开根号
减少: reduction次数 memory访问 kernel计算
LayerNorm(xi)=γiσ2+ϵxi−μ+β
归一化公式
残差
x+loss,防止梯度爆炸或消失问题,使整体变化平稳。
Adam
反向传播时需要参数更新,SGD是传统参数更新方法。然而只考虑当前梯度,当batchsize太小时,会出现噪声而导致更新不准确。(batch中一次全是猫,一次全是狗)。
动量把历史值通过权重考虑进去,Adam在动量的基础上又考虑了多个参数更新时量级不一致的问题。如w1=10000eta,w2=0.00001eta, 可能一次更新某个参数基本无变化。Adam通过一阶m(梯度方向),二阶v(步长大小),解决不同参数的更新问题
world_size: int
共有多少个进程
shard
切分数据分配给进程