我的课程笔记，欢迎关注：https://github.com/BBuf/how-to-optim-algorithm-in-cuda/tree/master/cuda-mode 。
这节课介绍了两个LinkedIn开发的 Liger Kernel的核心优化：RMSNorm和Fused Linear Cross Entropy。本笔记详细记录了课程中的优化的数学原理、实现方法和测试验证过程，包括一些脚本的解读。RMSNorm部分展示了反向传播的推导过程和内存优化技巧；Fused Linear Cross Entropy部分展示了如何通过检查点(checkpointing)、分块(chunking)和前向梯度(gradient-in-forward)等技术来减少内存使用。此外，课程还分享了在Triton框架开发上面的优化kernel中的一些实践经验，比如连续性(Contiguity)问题和索引越界问题的处理。通过这些优化，Liger Kernel能够将多GPU训练吞吐量提升20%并减少60%的内存使用，是Triton在工业界很好的一次实践。
第28课，Liger Kernel
Liger Kernel(https://github.com/linkedin/Liger-Kernel) 是一个专门为 LLM 训练设计的 Triton kernels 集合，由LinkedIn的工程师开发和维护。它能有效地将多 GPU 训练吞吐量提高 20%，并将内存使用量减少 60%。目前已经实现了与 HuggingFace 兼容的 RMSNorm、RoPE、SwiGLU、CrossEntropy、FusedLinearCrossEntropy 等功能，未来还会有更多。Liger Kernel可以直接与 Flash Attention、PyTorch FSDP 和 Microsoft DeepSpeed 配合使用。我们欢迎社区贡献，共同收集最佳的 LLM 训练kernel。
课程笔记之 RMSNorm
这张Slides介绍了本节课的大纲。具体包括了LLM（大语言模型）训练中的性能瓶颈问题，为什么选择使用Triton框架，并介绍如何实现RMS Norm和Fused Linear Cross Entropy，这可以减少很多内存，也会提到如何测试减少的内存。同时提供了三个重要的优化技巧：收敛性测试、连续性优化和地址范围处理。最后给Liger kernel打个广告。
这里讲到大模型的瓶颈不仅包含显存导致的OOM，还包含效率问题，有一个误区就是GPU利用率越高速度越快，实际上这是错误的理解，GPU利用率高只是说明GPU很忙，推荐阅读 https://arthurchiao.art/blog/understanding-gpu-performance/ 这篇文章来理解这个问题。此外，Profiler是理解一切性能问题的基础，推荐cuda-mode的lecture1和lecture16来学习如何使用Profiler。
然后作者对一个LLama模型做了一个在线的Profile，我们可以看到在内存变化阶段Cross entropy有一个峰值突变，它消耗了很多内存。如下图所示：
由于使用了Checkpointing技术，在前向和反向阶段的每个Transformer Block上也能观察到内存的升降，因为计算下一个Transformer Block的时候会释放当前Transformer Block占用的内存。这里的重点是Cross Entropy的内存消耗，来源是具体化logits的过程中产生的峰值内存，因为vocab size很大。
接下来作者介绍了一下kernel trace部分，从这部分我们可以看到LLama模型有很多elmentwise ops和很多cuda kernel launch的overhead。作者也提到这个kernel trace由于是用FSDP训练LLamaT，所以我们可以在每个Transformer Block看到2次all gather和一次reduce scatter。FSDP具体原理图可以参考：https://zhuanlan.zhihu.com/p/485208899 ，然后可以参考【翻译】使用PyTorch FSDP最大化训练吞吐量 。
简单总结一下，FSDP训练LLamaT可以观察到两个明显的瓶颈，第一个是Cross Entropy的大量内存消耗，第二个是elmentwise ops和很多cuda kernel launch的overhead。
这张Slides介绍了为什么要选择使用Triton（一个GPU编程框架）的几个主要原因：它比CUDA更容易编程，开发内核速度更快；它采用类似Numpy的向量化思维方式而不是传统的线程思维；对AI研究人员更友好，便于他们理解和扩展；作为Python原生框架不需要处理多个文件类型；并且依赖关系简单，在大多数情况下都能正常工作。总的来说，这些优势使Triton成为一个更现代、更易用的GPU编程解决方案。
这张Slides说明使用Triton写RMS Forward 会很简单，但是写Backward会比较难，下面会展示一些作者总结的技巧。
这张slides介绍了Backward Pass(Backprop)的基础知识，以"Backprop 101"为标题，主要强调了在学习Backward Pass时应该按元素思考，因为标量微积分比向量微积分更容易推导，同时建议复习微积分基础知识和矩阵-矩阵乘法公式。slides最后给出了矩阵乘法 Y = XW 的Backward Pass推导结果，包括对X的梯度 ∂L/∂X = (∂L/∂Y)W^T 和对W的梯度 ∂L/∂W = X^T(∂L/∂Y)，这些基础知识对于理解和实现更复杂操作（如RMS Norm）的Backward Pass非常重要。
这张Slides展示了RMSNorm(Root Mean Square Normalization)操作的Backward Pass(backprop)推导过程。主要包含两个关键公式：
Forward Pass公式：yi = (xi * wi) / sqrt((1/n) * ∑xk²)，表示对输入xi进行归一化 Backward Pass公式：dxi = ∂o/∂xi = ∑k (∂o/∂yk * ∂yk/∂xi)，使用链式法则计算梯度
这里特别强调了链式法则的应用原因：因为输入xi会影响所有的输出yi，所以在计算梯度时需要考虑xi对所有yi的影响并求和。这是RMSNormBackward Pass计算中的核心思想。
这张Slides展示了RMSNormBackward Pass的详细数学推导，特别强调了需要分开处理k=i和k≠i两种情况。其中引入了RMS（Root Mean Square）变量来简化表达式，最终得到了当k=i时的偏导数公式。通过数学变换，将一个复杂的表达式简化为更简洁的形式：(wi - 1/(RMS^2) * 1/n * xi^2 * wi)/RMS。这个推导过程对于实现RMSNorm的Backward Pass计算非常重要，它为后续的代码实现提供了理论基础。
这张Slides进一步展示了k≠i时的RMSBackward Pass的数学推导。
这张Slides把k=i和k≠i的两种情况合并起来获得了RMSBackward Pass的完整数学推导。并且我们可以从单个元素的推导推广到向量，这样就可以在Triton中方便的实现了。
这张Slides展示了Liger-Kernel实现RMSNorm的时候使用的2个技巧，Inplac Tensor reuse和Cache rms，可以从 https://github.com/linkedin/Liger-Kernel/blob/main/src/liger_kernel/ops/rms_norm.py 源码中看到。
作者提供了一个jupyter的demo，用来展示对Liger-Kernel的测试过程，包含正确性，性能，内存等方面的测试。 https://colab.research.google.com/drive/1CQYhul7MVG5F0gmqTBbx1O1HgolPgF0M?usp=sharing ，我们接下来解读一下这个测试。
Live Demo: RMSNorm: 确认正确性和性能的测试解读
现在你已经学会了如何推导RMSNorm的Backward Pass以及内存节省技巧，实现本身相对来说比较直接。因此，我们将重点关注测试，并使用来自Liger Kernel的现有实现。
为什么需要测试？
假设我们已经有了一个可运行的RMSNorm版本，在将其部署到生产环境之前，我们需要验证以下几点：
正确性：确保kernel的精度与原始实现一致。任何偏差都可能影响模型收敛或导致严重错误。 性能：确认kernel在时间和内存使用上都比原始版本更高效。如果没有这些改进，用Triton重新实现就失去了意义。
正确性测试
准备一个纯PyTorch实现，比如使用HuggingFace提供的版本。
我们需要用不同的输入形状和数据类型来测试实现。除了像2的幂这样的规则形状外，测试不规则形状也很重要，以确保能正确处理边界情况。
设置容差可能比较棘手。通常对于fp32，使用atol = 1e-7和rtol = 1e-5。对于bf16，使用atol = 1e-3和rtol = 1e-2。但在实践中，即使kernel是精确的，有时也可能需要进一步放宽容差。
稍后，我们将讨论另一种测试方法来验证kernel不会对端到端的收敛产生负面影响。
import torch
import torch.nn as nn
# Copy from HuggingFace
class LlamaRMSNorm(nn.Module):
def __init__(self, hidden_size, eps=1e-6):
"""
LlamaRMSNorm是等价于T5LayerNorm的实现
参数:
hidden_size: 隐藏层维度大小
eps: 用于数值稳定性的小常数
"""
super().__init__()
# 初始化可学习的缩放参数
self.weight = nn.Parameter(torch.ones(hidden_size))
# 保存epsilon值用于避免除零
self.variance_epsilon = eps
def forward(self, hidden_states):
# 保存输入数据类型
input_dtype = hidden_states.dtype
# 转换为float32以提高精度
hidden_states = hidden_states.to(torch.float32)
# 计算方差
variance = hidden_states.pow(2).mean(-1, keepdim=True)
# 标准化操作
hidden_states = hidden_states * torch.rsqrt(variance + self.variance_epsilon)
# 应用可学习参数并恢复原始数据类型
return self.weight * hidden_states.to(input_dtype)
import torch
from liger_kernel.transformers.rms_norm import LigerRMSNorm
input_data = [
(4, 16, 32, torch.float32, 1e-6, 1e-4),
(8, 32, 64, torch.float32, 1e-6, 1e-4),
(16, 64, 128, torch.float32, 1e-6, 1e-4),
(3, 9, 13, torch.float32, 1e-6, 1e-4),
# T4 GPU doesn't support bf16 :(
# (16, 64, 128, torch.bfloat32, 1e-3, 1e-2),
]
for bs, sl, hd, dtype, atol, rtol in input_data:
# h
_tensor = torch.randn(bs, sl, hd, device="cuda", dtype=dtype)
h1 = _tensor.clone().requires_grad_(True)
h2 = _tensor.clone().requires_grad_(True)
# do
do = torch.randn(bs, sl, hd, device="cuda", dtype=dtype)
# llama
llama_rms = LlamaRMSNorm(hidden_size=hd).to("cuda").to(dtype)
llama_o = llama_rms(h1)
llama_o.backward(do.clone(), retain_graph=True)
# triton
triton_rms = LigerRMSNorm(hidden_size=hd).to("cuda").to(dtype)
triton_o = triton_rms(h2)
triton_o.backward(do.clone(), retain_graph=True)
assert torch.allclose(llama_o, triton_o, atol=atol, rtol=rtol) is True
# print(h1.grad, h2.grad)
assert torch.allclose(h1.grad, h2.grad, atol=atol, rtol=rtol) is True
性能测试
我们需要测试两个维度:速度和内存。但是应该使用什么输入形状来测试呢?你可以使用训练时的实际输入形状来测试。例如,在微调LLaMA 3-8B模型时,我们通常使用batch size为4,hidden size为2048。我们将序列长度作为变量。
这样,测试结果就能反映出我们在生产环境训练中可以预期的实际收益。这里使用了Triton提供的自动测试工具。
import os
import torch
import torch.nn as nn
import triton
@triton.testing.perf_report(
[
triton.testing.Benchmark(
x_names=["seq_len"],
x_vals=[2**i for i in range(8, 11)], # 256, 512, 1024
xlabel="seq len",
line_arg="provider",
line_vals=["liger", "huggingface"],
line_names=["Liger", "Hugging Face"],
styles=[("blue", "solid"), ("orange", "solid")],
ylabel="time (ms)",
plot_name="rmsnorm-full-speed-benchmark",
args={"batch_size": 4, "hidden_size": 2048, "dtype": torch.float32, "mode": "full"},
),
]
)
def bench_speed_rms_norm(batch_size, seq_len, hidden_size, dtype, provider, mode, eps=1e-5, device="cuda"):
x_shape = (batch_size, seq_len, hidden_size)
triton_rms = LigerRMSNorm(hidden_size=hidden_size).to("cuda")
llama_rms = LlamaRMSNorm(hidden_size=hidden_size).to("cuda")
x = torch.randn(x_shape, dtype=dtype, device="cuda")
dy = torch.randn_like(x)
x.requires_grad_(True)
x = x.view(batch_size * seq_len, hidden_size)
dy = dy.view(batch_size * seq_len, hidden_size)
quantiles = [0.5, 0.2, 0.8]
def full():
if provider == "liger":
y = triton_rms(x)
elif provider == "huggingface":
y = llama_rms(x)
y.backward(dy, retain_graph=True)
ms, min_ms, max_ms = triton.testing.do_bench(
full, quantiles=quantiles, grad_to_none=[x], rep=500
)
return ms, max_ms, min_ms
bench_speed_rms_norm.run(show_plots=True, print_data=True)
def test_memory(func, _iter):
total_mem = []
for _ in range(_iter):
torch.cuda.memory.reset_peak_memory_stats()
func()
mem = torch.cuda.max_memory_allocated() / (2**20)
total_mem.append(mem)
return sum(total_mem) / len(total_mem)
@triton.testing.perf_report(
[
triton.testing.Benchmark(
x_names=["seq_len"],
x_vals=[2**i for i in range(8, 11)], # 256, 512, 1024
xlabel="seq len",
line_arg="provider",
line_vals=["liger", "huggingface"],
line_names=["Liger", "Hugging Face"],
styles=[("blue", "solid"), ("orange", "solid")],
ylabel="Memory (MB)",
plot_name="rmsnorm-full-memory-benchmark",
args={"batch_size": 4, "hidden_size": 2048, "dtype": torch.float32, "mode": "full"},
),
]
)
def bench_memory_rms_norm(batch_size, seq_len, hidden_size, dtype, provider, mode, eps=1e-5, device="cuda"):
x_shape = (batch_size, seq_len, hidden_size)
triton_rms = LigerRMSNorm(hidden_size=hidden_size).to("cuda")
llama_rms = LlamaRMSNorm(hidden_size=hidden_size).to("cuda")
x = torch.randn(x_shape, dtype=dtype, device="cuda")
dy = torch.randn_like(x)
x.requires_grad_(True)
x = x.view(batch_size * seq_len, hidden_size)
dy = dy.view(batch_size * seq_len, hidden_size)
quantiles = [0.5, 0.2, 0.8]
def full():
if provider == "liger":
y = triton_rms(x)
elif provider == "huggingface":
y = llama_rms(x)
y.backward(dy, retain_graph=True)
mem = test_memory(full, 10)
return mem
bench_memory_rms_norm.run(show_plots=True, print_data=True)
RMSNorm测试总结
我们可以清楚地看到,Triton实现在速度和内存使用上都优于原始实现,并且我们也验证了其正确性!由于Google Colab的GPU限制,我们只进行了部分测试。在Liger-Kernel的实际测试中,我们还使用更大的输入尺寸验证了bf16的性能。完整版本请参考 https://github.com/linkedin/Liger-Kernel 。
课程笔记之Fused Linear Cross Entropy
这张Slides展示了Transformer模型里面 Linear Cross Entropy 的Forward Pass和Backward Pass过程。图中左侧显示了Forward Pass(Forward)流程，从输入(input)通过 lm_head 层产生激活值(Activations)，然后输出(output)与目标值(target)计算交叉熵(Cross Entropy)。右侧显示了Backward Pass(Backward)流程，展示了梯度(Gradients)的传递方向。图片底部指出了一个问题：大词汇表尺寸(Large Vocab Size)是这个模型面临的主要挑战。
这5张slides展示了完整的线性层和交叉熵梯度计算的推导过程。首先介绍了线性层的Forward Pass(y = Wx)和Backward Pass(∂o/∂x = W^T∂y)，接着对交叉熵损失函数l = -∑yⱼlog(exp(xⱼ)/∑exp(xᵢ))求偏导，将其分解为两项分别推导：一项是针对包含xₖ的项，另一项是针对其他所有项。经过复杂的代数运算和化简，最终得到了简洁的梯度表达式∂l/∂xₖ = -yₖ + softmax(xₖ)，并讨论了yₖ=1和yₖ=0两种特殊情况下的结果，将这个复杂的梯度计算问题优化为目标值与softmax的差值形式。有了最后的这个等式，我们就可以比较方便的在Triton中计算向量的交叉熵梯度了。
这里需要注意下第二点，由于Cross Entropy是最后一层，它的输出一定是一个标量，所以我们可以在forward的时候就计算梯度。
这张Slides展示了Fused Linear Cross Entropy中的梯度检查点(Gradient Checkpointing)技术。在模型训练过程中，左侧展示了正向传播路径：从input经过lm_head层得到output，然后与target计算交叉熵；右侧展示了Backward Pass路径。关键点在于，在Backward Pass时会重新计算前向过程(Forward Recomputation)，而不是保存激活值，这样可以节省存储空间。图中用"×"表示激活值(Activations)，用"△"表示梯度(Gradients)。底部说明文字强调了这个策略的核心：在Backward Pass时重新计算前向过程，避免了需要持久化存储激活值。
这张Slides展示了Fused Linear Cross Entropy中的gradient-in-forward优化技术。与之前的梯度检查点方法不同，这里在Forward Pass过程中就同时计算梯度，避免了需要重新计算前向过程。图中显示了从input通过lm_head层到output，再与target计算交叉熵的流程，其中lm_head层同时包含了激活值(用"×"表示)和梯度(用"△"表示)。底部说明文字强调了这种方法的优势：通过在Forward Pass时就计算梯度，可以消除重新计算前向过程的需求，从而提高计算效率。（可以在前向过程中计算梯度的原因是因为Cross Entropy的输出是一个Scalar，所以上游的梯度稳定为1。这样就避免了在Backward Pass计算梯度和重计算了）
这张Slides展示了Fused Linear Cross Entropy中的Chunking（分块）技术。图中显示input被虚线分成多个块（chunks），表明输入数据被分成若干小块进行处理。这种方法的核心思想是：每次只处理输入数据的一个块，因此在任意时刻只需要在内存中保存当前块的激活值（用"×"表示）和梯度（用"△"表示）。底部说明文字解释了这种策略的优势：通过逐块处理输入数据，可以显著减少内存使用，因为同一时刻只需要保存一小块数据的激活值和梯度信息，而不是全部数据。需要注意的是，对输入进行分chunk处理由于Cross Entropy的梯度计算存在Softmax操作，具体见上面几张Slides，当我们对输入进行分块之后我们需要像Online Softmax算法那样逐chunk更新缩放系数，最后才能对 hidden_states 得到正确的梯度。
Live Demo: FusedLinearCrossEntropy: 确认内存减少
我们已经讨论了检查点(checkpointing)、分块(chunking)和前向梯度(gradient-in-forward)这些概念,它们可以避免具体化logits以及重新计算。由于实现比较复杂,所以这部分就留给大家自学。在这个notebook中,我们将验证FusedLinearCrossEntropy是否真的可以在保持相当速度的同时减少峰值内存使用。
FusedLinearCrossEntropy基准测试
比较FusedLinearCrossEntropy和Hugging Face实现在速度和内存方面的表现
import os
import torch
import triton
from liger_kernel.transformers.fused_linear_cross_entropy import (
LigerFusedLinearCrossEntropyLoss,
)
class TorchLMHeadCE(torch.nn.Module):
"""Ground truth implementation of the linear fused with torch based cross entropy loss.
:param H: hidden size
:param V: vocab size
:param ignore_index: index to ignore
:param reduction: reduction method
"""
def __init__(self, H: int, V: int, dtype: torch.dtype, ignore_index: int = -100):
super().__init__()
# 定义线性层,不带偏置
self.lin = torch.nn.Linear(
in_features=H, out_features=V, bias=False, dtype=dtype
)
# 定义交叉熵损失函数
self.ce_loss = torch.nn.CrossEntropyLoss(
ignore_index=ignore_index, reduction="mean"
)
def forward(self, x, y):
# 前向传播:先通过线性层,再计算交叉熵损失
logits = self.lin(x)
return self.ce_loss(logits, y)
class LigerLMHeadCE(torch.nn.Module):
def __init__(self, H: int, V: int, dtype: torch.dtype, ignore_index: int = -100):
super().__init__()
# 定义线性层,不带偏置
self.lin = torch.nn.Linear(
in_features=H, out_features=V, bias=False, dtype=dtype
)
# 定义Liger的融合交叉熵损失函数
self.ce_loss = LigerFusedLinearCrossEntropyLoss(
ignore_index=ignore_index, reduction="mean"
)
def forward(self, x, y):
# 前向传播:直接使用权重矩阵、输入和目标计算融合的交叉熵损失
return self.ce_loss(self.lin.weight, x, y)
def test_memory(func, _iter):
# 用于测试内存使用的辅助函数
total_mem = []
for _ in range(_iter):
# 重置CUDA峰值内存统计
torch.cuda.memory.reset_peak_memory_stats()
func()
# 获取最大分配内存(MB)
mem = torch.cuda.max_memory_allocated() / (2**20)
total_mem.append(mem)
# 返回平均内存使用量
return sum(total_mem) / len(total_mem)
@triton.testing.perf_report(
[
triton.testing.Benchmark(
x_names=["BT"],
x_vals=[2**i for i in range(10, 13)], # 1024, 2048, 4096
xlabel="B x T",
line_arg="provider",
line_vals=["liger", "huggingface"],
line_names=["Liger", "Hugging Face"],
styles=[
("blue", "solid"),
("orange", "solid"),
],
ylabel="GPU memory usage (MB)",
plot_name="fused-linear-cross-entropy-memory-benchmark",
args={"H": 4096, "V": 128256, "dtype": torch.float32},
)
]
)
def bench_memory_cross_entropy(BT, H, V, provider, dtype, device="cuda"):
# 打印基准测试参数
print(
f"Running benchmark with BT={BT}, H={H}, V={V}, dtype={dtype} provider={provider}"
)
# 初始化PyTorch和Liger模型
torch_lm_head_ce = TorchLMHeadCE(H=H, V=V, dtype=dtype).to(device)
liger_lm_head_ce = LigerLMHeadCE(H=H, V=V, dtype=dtype).to(device)
# 生成随机输入和目标
_input = torch.randn(BT, H, requires_grad=True, dtype=dtype, device=device)
target = torch.randint(V, (BT, 1), dtype=torch.long, device=device).squeeze(1)
def fwd():
# 根据provider选择不同的前向实现
if provider == "liger":
return liger_lm_head_ce(_input, target)
elif provider == "huggingface":
return torch_lm_head_ce(_input, target)
def full():
# 完整的前向+反向传播
y = fwd()
y.backward()
# 测试内存使用并返回结果
mem = test_memory(full, _iter=10)
return mem
bench_memory_cross_entropy.run(show_plots=True, print_data=True)
@triton.testing.perf_report(
[
triton.testing.Benchmark(
x_names=["BT"],
x_vals=[2**i for i in range(10, 13)], # 1024, 2048, 4096
xlabel="B x T",
line_arg="provider",
line_vals=["liger", "huggingface"],
line_names=["Liger", "Hugging Face"],
styles=[
("blue", "solid"),
("orange", "solid"),
],
ylabel="Time (ms)",
plot_name="fused-linear-cross-entropy-speed-benchmark",
args={"H": 4096, "V": 128256, "dtype": torch.float32},
)
]
)
def bench_speed_cross_entropy(BT, H, V, provider, dtype, device="cuda"):
print(
f"Running benchmark with BT={BT}, H={H}, V={V}, dtype={dtype} provider={provider}"
)
torch_lm_head_ce = TorchLMHeadCE(H=H, V=V, dtype=dtype).to(device)
liger_lm_head_ce = LigerLMHeadCE(H=H, V=V, dtype=dtype).to(device)
_input = torch.randn(BT, H, requires_grad=True, dtype=dtype, device=device)
target = torch.randint(V, (BT, 1), dtype=torch.long, device=device).squeeze(1)
def fwd():
if provider == "liger":
return liger_lm_head_ce(_input, target)
elif provider == "huggingface":
return torch_lm_head_ce(_input, target)
def full():
y = fwd()
y.backward()
quantiles = [0.5, 0.2, 0.8]
ms, min_ms, max_ms = triton.testing.do_bench(full, quantiles=quantiles, rep=100)
return ms, min_ms, max_ms
bench_speed_cross_entropy.run(show_plots=True, print_data=True)
FusedLinearCrossEntropy测试总结
我们可以观察到我们的实现在内存使用上有显著优势，因为我们在任何时候都不会具体化完整的logits。 速度略微变慢，但由于lm_head + cross_entropy只执行一次，而transformer block要执行N次，这个开销是可以接受的。这使我们能够增加batch size、序列长度，或者关闭梯度检查点。
Triton 操作的是physical view Contiguous()非常重要
这张Slides主要介绍了收敛性测试中逐层比较的重要性，指出仅进行单元正确性和性能测试对于生产环境是不够的，因为实际生产中可能会遇到连续性、张量形状和数据类型的差异。因此建议通过模拟真实的生产训练环境来验证模型输出(logits)、权重(weights)和损失值(loss)，Slides底部还提供了一个Google Colab链接用于比较Triton内核补丁版本与原始模型的逐层对比，这些内容强调了在模型部署到生产环境之前进行全面测试的重要性(https://colab.research.google.com/drive/1e52FH0BcE739GZaVp-3_Dv7mc4jF1aif?usp=sharing)。这个脚本比较简单，我们也可以让ChatGPT写，这里就不单独看了。
这张Slides主要讲解了在CUDA编程中"连续性(Contiguity)"这个容易被忽视但非常重要的问题。它指出连续性问题可能会导致难以调试的静默bug，需要花费大量时间来解决。幻灯片通过张量(Tensor)的例子，展示了逻辑视图(logical view)和物理视图(physical view)之间的区别，并用一个具体的带有stride的张量表示方式来说明这个概念。图中展示了一个2x2的张量在逻辑上和物理存储上的不同表现形式，以及相应的大小(sizes)和步长(strides)参数。对于Triton来说，它是在物理视图(physical view)上进行计算的，所以需要特别注意contiguous()的问题。https://colab.research.google.com/drive/1llnAdo0hc9FpxYRRnjih0l066NCp7Ylu?usp=sharing#scrollTo=1jTVlU1NC-TN 这个juypter展示了Liger Kernel的RoPE实现，因为忽略了对输入进行contiguous()操作导致单独的单元测试总是可以通过的，但是进行模型训练时始终出现loss发散的问题，这个问题也是通过上面的收敛性测试中逐层比较发现的。
Triton 中的index越界bug
这张Slides讲的是Triton的program_id是int32来表示的，然后在开发Cross Entropy时没有考虑到这一点，导致在较大的Vocab Size时index会越界。https://colab.research.google.com/drive/1WgaU_cmaxVzx8PcdKB5P9yHB6_WyGd4T?usp=sharing#scrollTo=X_Dn9wzVNpMC 这个juypter展示了这个问题，修复的方案是把program_id转换为int64。不过，因为32位寻址可能会导致性能很慢，所以需要非常谨慎的处理这个问题。例如在PyTorch中，针对这两种不同的数据类型会通过C++模板来处理，它们的实现会共享一个kernel，但是可以避免这个index溢出的问题。
Liger Kernel 相关开源信息和反响
这几张Slides总结一下，LinkedIn开发的Liger Kernel是一个专为LLM训练优化的GPU高效运行时kernel，它能够将多GPU训练吞吐量提升20%并减少60%的内存使用，支持多种Hugging Face兼容功能，可与Flash Attention、PyTorch等主流框架协同工作。该项目在开源社区获得了积极反响，开发者反馈显示其性能表现优异。项目的成功离不开众多贡献者的支持，包括LOGO设计、训练灵感来源、测试数据集提供等方面的帮助，以及来自CUDA/Triton社区的大力支持，充分体现了开源协作的力量。
课程笔记总结
这节课介绍了两个核心优化：RMSNorm和Fused Linear Cross Entropy。本笔记详细记录了课程中的优化的数学原理、实现方法和测试验证过程，包括一些脚本的解读。RMSNorm部分展示了反向传播的推导过程和内存优化技巧；Fused Linear Cross Entropy部分展示了如何通过检查点(checkpointing)、分块(chunking)和前向梯度(gradient-in-forward)等技术来减少内存使用。此外，课程还分享了在Triton框架开发上面的优化kernel中的一些实践经验，比如连续性(Contiguity)问题和索引越界问题的处理。通过这些优化，Liger Kernel能够将多GPU训练吞吐量提升20%并减少60%的内存使用，是Triton在工业界很好的一次实践。
