大模型模型优化深度学习【免费下载链接】Liger-KernelEfficient Triton Kernels for LLM Training项目地址https://gitcode.com/gh_mirrors/li/Liger-Kernel点击查看免费下载本指南围绕 Liger-Kernel 性能优化技能liger-kernel-perf中的 Optimization Profile优化画像文档模板展开逐字段讲解如何为一个现存 Triton 内核生成结构化的性能画像从内核身份登记、GPU 硬件探测、基线基准数据到 NCU 硬件级剖析、瓶颈分类、当前配置审计与优化机会清单。你将掌握该模板的每个字段含义、数据来源、填写方法与在 3 阶段优化流水线中的交接作用可直接照模板产出属于自己的内核优化报告。一、模板定位Profile 是 3 阶段优化流水线的“跨阶段契约”optimization-profile.md是.agents/skills/liger-kernel-perf/技能包中的一份输出模板由该技能包的 Profiler剖析阶段产出作为 Optimizer优化阶段的输入依据。整个技能包定义了 Liger-Kernel 内核性能优化的完整闭环Stage 1 Profile剖析创建optimization/{kernel}/工作区、备份原始内核、运行基线基准、探测 GPU 架构、可选 NCU 剖析、分析内核代码、分类瓶颈最终生成 profile 并保存到optimization/{kernel}/profile.md即本模板的填充结果Stage 2 Optimize优化读取 profile 与原始内核先做参数调优BLOCK_SIZE / num_warps / num_stages 手动扫描禁用triton.autotune再依据瓶颈分类套用优化策略目录 optimization-strategies.md 中的技术逐个生成变体optimization/{kernel}/{kernel}_vN.py与实验笔记Stage 3 Finalize收尾将获胜变体应用回生产内核运行完整测试硬性门槛、checkstyle、生成三方对比图与最终报告。在交互模式下每个阶段之间设有人工检查点用户说出“just optimize it”“optimize autonomously”等指令时进入自主模式全程端到端自动运行。Profile 正是 Stage 1 的交付物也是 Stage 2 决策的唯一数据入口——因此本模板被称为“跨阶段契约”cross-stage contract。注本文后续所有命令与路径均以当前仓库Liger-Kernel为基准。该技能仅支持 NVIDIA GPU。二、字段速览Profile 的七大区块模板由 7 个区块构成下文逐一展开Kernel Identity内核身份GPU InformationGPU 信息Baseline Benchmarks基线基准NCU Profiling SummaryNCU 剖析摘要Bottleneck Classification瓶颈分类Current Configuration Analysis当前配置分析Identified Optimization Opportunities 与 Recommended Strategy Order优化机会与推荐策略顺序模板头部约定“Generated by Profiler Agent (Stage 1)”与日期表明这是一份阶段化产出物便于后续阶段快速定位来源与时效性。三、Kernel Identity内核身份登记本区块用一张表格固定记录被优化内核的基本事实其中每个字段都在 Profiler 的前置校验Pre-Flight Validation中被验证过字段含义校验规则Kernel Name内核名如rms_norm、cross_entropy用户请求中的target_kernelSource File生产内核路径必须存在src/liger_kernel/ops/{kernel}.pyBenchmark Script基准脚本路径必须存在benchmark/scripts/benchmark_{kernel}.pyTest File测试文件路径必须存在test/transformers/test_{kernel}.pyTier Classification分层 描述见下方 Tier 分类Tier 分类是后续瓶颈判断的重要先验来自内核代码结构分析对应 Profiler 的 Step 5Tier 1element-wise 逐元素无跨行归约每个 program 处理一行。典型如swiglu、geglu、relu_squaredTier 2reduction 归约存在跨列归约可能需要基于 SM 的并行。典型如rms_norm、layer_normTier 3fused/complex 融合/复杂多趟算法、梯度前置gradient-in-forward等技巧。典型如fused_linear_cross_entropy、cross_entropy。从源码结构看src/liger_kernel/ops/rms_norm.py 的前向用每行一个 program 的(n_rows,)网格反向则用(sm_count,)网格配合rows_per_program做 dW 归约正是 Tier 2 “需要 SM 级并行”的典型结构而cross_entropy在 src/liger_kernel/ops/cross_entropy.py 中把 softmax、损失与梯度计算融进单内核并复用输入张量存储梯度属于 Tier 3 的“梯度前置”手法。四、GPU Information硬件探测与目标 GPU 核对4.1 探测命令Profiler 通过以下命令读取 GPU 属性对应其 Step 2python -c import torch; props torch.cuda.get_device_properties(0); print(fGPU: {props.name}, SM: {props.major}.{props.minor}, SMs: {props.multi_processor_count}, Mem: {props.total_mem // (1024**3)}GB)4.2 架构分类标准SM 版本架构代号代表 GPUSM 8.xAmpereA100、A10 等SM 9.0HopperH100、H200 等SM 10.0BlackwellB200 等模板中的Hardware Match字段用于核对用户指定目标 GPUAmpere / Hopper / Blackwell / auto-detect与实际探测硬件是否一致若不一致记录 mismatch但继续以实际硬件为准开展剖析。这个“目标与实机核对”的设计是因为后续参数调优如 num_stages的推荐值高度依赖具体架构。五、Baseline Benchmarks基线基准数据5.1 数据来源与采集基线数据来自仓库自带的基准脚本运行方式Profiler Step 3cd benchmark/scripts python benchmark_{kernel}.py --overwrite脚本将结果写入benchmark/data/all_benchmark_data.csv随后把该内核相关行复制为optimization/{kernel}/benchmarks/v0_baseline.csv作为 v0 基线快照。5.2 速度表median ms模板的速度表按x_value基准变量如 hidden_size、n_cols分组每种规模记录三种模式的耗时中位数x_valueMode{kernel} (Liger){baseline_provider}Speedup{x_val_1}forward………x{x_val_1}backward………x{x_val_1}full………x{x_val_2}forward………x……………其中baseline_provider通常为huggingface或torch。以 rms_norm 为例基准变量为hidden_size推荐测试规模[1024, 2048, 4096, 8192, 16384]固定配置{M: 4096, eps: 1e-6, dtype: torch.float32}见 rms-norm-profile.md 的 Benchmarking 节。5.3 内存表median MBx_value{kernel} (Liger){baseline_provider}Savings{x_val_1}………%{x_val_2}………%内存统计的是 full 模式前向反向的峰值占用Savings 反映相对基线的省内存比例。Liger 内核的省内存主要来自两大手法融合如 cross_entropy 融合 softmaxloss梯度避免中间张量落 HBM与以算换存如 rms_norm 只缓存每行 1 个 rstd 标量而非整行中间量。5.4 Observations基线观察摘要模板要求提炼三条基线观察Slowest mode耗时最长的模式forward / backward / full及其原因——例如归约内核的反向需要额外一次 dW 归约与 HBM 回写通常最慢Largest x_value tested最大测试规模Performance scaling trend性能随规模的变化趋势线性、亚线性、是否在某规模出现拐点这是判断内核适用输入域的依据。六、NCU Profiling Summary硬件级剖析摘要6.1 可用性探测which ncu 2/dev/null echo NCU_AVAILABLE || echo NCU_NOT_AVAILABLE6.2 可用的典型剖析命令读取目标内核的torch.autograd.Function类构造最小复现脚本Profiler Step 4例如 rms_normncu --set full --target-processes all -o optimization/{kernel}/ncu_profile \ python -c import torch from liger_kernel.ops.rms_norm import LigerRMSNormFunction x torch.randn(4096, 4096, devicecuda, dtypetorch.bfloat16, requires_gradTrue) w torch.ones(4096, devicecuda, dtypetorch.bfloat16) y LigerRMSNormFunction.apply(x, w, 1e-6, 0.0, llama, False) y.sum().backward() 其他内核需相应调整导入路径、函数参数与张量形状。6.3 提取指标Metric含义SM OccupancySM 占用率低值常由寄存器压力或网格过小导致Compute Throughput计算吞吐占峰值百分比Memory Throughput内存吞吐占峰值百分比L1 / L2 Cache Hit Rate缓存命中率低 L2 命中提示缓存策略可优化Warp Stall Reasonstop 3最突出的线程束停顿原因直接暴露等待点模板用两个条件区块available / unavailable处理两种情况NCU 不可用时明确声明“Bottleneck classification is based on code heuristics only”仅基于代码启发式避免把推断当实测。七、Bottleneck Classification瓶颈分类与证据链7.1 分类判据分类NCU 判据代码启发式判据Memory-bound内存吞吐 60% 峰值且计算吞吐 40%简单逐元素运算、大张量、每次 load 算术少Tier 1 几乎总是Compute-bound计算吞吐 60% 峰值且内存吞吐 40%每元素复杂数学exp/log/sigmoid 链、归约、每 load 多运算Tier 3 重计算Balanced两者均在 30-60%Tier 2 归约内核常落此区间7.2 Evidence 表可追溯的决策依据模板用“信号-观察-支持结论”三列表格把分类落到实处SignalObservationSupports{evidence_signal_1}{evidence_obs_1}{evidence_supports_1}………每条证据都应可追溯到 NCU 指标、基准趋势或代码特征。例如内存吞吐 85% vs 计算吞吐 22%NCU 信号→ 支持 memory-bound或“每次 load 仅一次乘法”的代码结构 → 支持 memory-bound。八、Current Configuration Analysis当前配置审计8.1 参数表ParameterCurrent ValueHow SetNotesBLOCK_SIZE…calculate_settings() / 硬编码 / autotune…num_warps………num_stages………Grid Dimensions1D / 2D / block-row……当前仓库的默认启发式实现在 src/liger_kernel/ops/utils.py 的calculate_settings(n)def calculate_settings(n): # reference: unsloth kernels utils MAX_FUSED_SIZE 65536 BLOCK_SIZE triton.next_power_of_2(n) if BLOCK_SIZE MAX_FUSED_SIZE: raise RuntimeError( fCannot launch Triton kernel since n {n} exceeds the recommended Triton blocksize {MAX_FUSED_SIZE}. ) num_warps 4 if BLOCK_SIZE 32768: num_warps 32 if not is_hip() else 16 elif BLOCK_SIZE 8192: num_warps 16 elif BLOCK_SIZE 2048: num_warps 8 return BLOCK_SIZE, num_warps三个审计要点BLOCK_SIZE 取 n 的 2 次幂受MAX_FUSED_SIZE 65536上限保护防止寄存器溢出若 BLOCK_SIZE 远超实际 n_cols 会造成掩码加载与浪费num_warps 按规模阶梯取值2048→8、8192→16、32768→32HIP 下 16num_stages 从未设置——Triton 默认 1。这是模板 Notes 列最重要的缺口也是优化策略目录中“最常被忽视的参数”。8.2 Memory Access Patterns访存模式Stride pattern连续或跨步访问。Liger 多数内核通过col_offsets tl.arange(0, BLOCK_SIZE)实现 warp 内连续访问需重点检查反向内核或多维网格中是否出现破坏合并的跨步指针算术Cache modifiers是否使用eviction_policyevict_first流式、evict_last复用In-place operations就地操作。cross_entropy 将梯度写入输入张量即为此类Saved for backward vs recomputed反向是保存还是重算中间量。rms_norm 保存每行 1 个rstd标量避免重算rsqrt属于“小量保存、避免重算”的典型取舍见 rms-norm-profile.md 的 Key Patterns 节。8.3 Compute Patterns计算模式Precision-sensitive opssigmoid、rsqrt、exp、log等精度敏感运算通常需 FP32 参与Casting modeLiger 通过casting_mode参数llama/gemma/none区分不同模型的精度策略Algorithm type是否使用在线online/chunked算法——cross_entropy已用在线 softmax 单趟完成 max/sum 统计constexpr parameterstl.constexpr参数可让编译器裁掉死代码分支、静态展开循环是降低寄存器压力的关键工具。九、优化机会、策略顺序与迭代预估9.1 Identified Optimization Opportunities基于内核代码审计列出具体机会Profiler Step 5.5常见类别代码中遗留的 TODO/FIXME 性能注释缺失 block-row 模式变体多行/每 program适配小 n_cols 大 n_rows次优的 BLOCK_SIZE 启发式不必要的内存往返可重算、可融合、可就地。9.2 Recommended Strategy Order依据瓶颈分类给出策略顺序。策略全目录见 optimization-strategies.md决策框架优先级总是先做参数调优Section 1BLOCK_SIZE / num_warps / num_stages 手动扫描Memory-bound → 访存合并、缓存修饰符、降内存流量、数据复用、block-row 模式Compute-bound → 算法改进、去冗余计算、张量核利用、降寄存器压力、2D 网格平铺Balanced → 交织应用上述两组架构特化Ampere / Hopper / Blackwell叠加在最优变体之上。9.3 为什么不能使用 triton.autotune这是参数调优一节的关键约束优化策略目录明确给出三条架构性理由前后向上下文耦合Liger 内核用torch.autograd.Functionforward()经calculate_settings()算出 BLOCK_SIZE/num_warps 存入ctxbackward()取出后以相同配置启动内核autotune 会把配置变成编译期常量反向内核无法得知前向为当前输入自动选择了哪套配置多后端不兼容NPU/Ascend 完全不支持num_warps、num_stages而 autotune 依赖这两个参数推理预热延迟每个新形状触发一次全配置基准单形状首调用惩罚 100ms对变长序列的 LLM 推理不可接受。dyt.py、grpo_loss.py中被注释掉的 autotune 装饰器即此前探索被放弃的证据。因此正确做法是替换calculate_settings为经基准标定的更优启发式并让 num_stages 也经ctx传到反向内核。9.4 架构相关参数建议与迭代预估架构num_stagesnum_warpsBLOCK_SIZEAmpereSM 8.x2 通常最优4-1632 增调度开销1024-4096HopperSM 9.03-5TMA/异步管线受益可容忍 322048-8192BlackwellSM 10.04-5 起步同 Hopper可更大256KB SMEM模板末尾的Estimated Iterations用于预估需要的变体数量默认max_variants 8及理由帮助用户对优化成本形成预期。整个优化循环的停止条件为预算耗尽、连续 2 个变体提升 1%收益递减、或达成用户设定的target_metric。十、Profile 的产出物与后续交接模板末尾注明“Profile saved tooptimization/{kernel}/profile.md”。Stage 1 完成后交互模式向用户展示 GPU 信息、基线摘要表、瓶颈分类及证据、策略顺序与迭代预估等待确认后进入 Stage 2自主模式直接进入 Stage 2。Stage 2 的 Optimizer 在生成每个变体前都必须完整读取该 profile并在实验笔记variant-notes.md中引用瓶颈分类作为假设依据Stage 3 的 Finalizer 则在报告中回引 profile 的基线数据。因此一份字段完整、证据可追溯的 profile直接决定了后续优化迭代的质量上限。十一、快速上手为你的内核生成第一份 Profile选定目标内核如rms_norm、cross_entropy、swiglu确认 src/liger_kernel/ops 下内核文件、benchmark/scripts/benchmark_{kernel}.py、test/transformers/test_{kernel}.py三者齐备建立工作区创建optimization/{kernel}/将内核快照为original_{kernel}.py探测 GPU 与跑基线执行上文 GPU 探测命令与python benchmark_{kernel}.py --overwriteNCU 可用则剖析按 6.2 的最小复现脚本执行并提取 5 项指标审计内核代码确定 Tier 分类、calculate_settings来源、访存/计算模式与具体优化机会分类瓶颈按第七节判据得出 memory-bound / compute-bound / balanced并填写证据表按模板产出optimization/{kernel}/profile.md附推荐策略顺序与迭代预估。按此流程产出的 profile可直接作为后续参数调优与诊断驱动优化的起点对rms_norm而言其反向 SM 网格归约、rstd 缓存与calculate_settings未设 num_stages 这三个特征通常就是第一批优化变体的入手点。赞分享大模型模型优化深度学习【免费下载链接】Liger-KernelEfficient Triton Kernels for LLM Training项目地址https://gitcode.com/gh_mirrors/li/Liger-Kernel点击查看免费下载相关推荐Liger-Kernel 性能剖析指南用 Profiler Agent 诊断 Triton Kernel 瓶颈并生成优化策略Liger Kernel 性能剖析指南用 Profiler Agent 诊断 Triton Kernel 瓶颈并生成优化策略 本指南围绕 Liger Kern大模型模型优化深度学习Liger Kernel Perf面向 Liger-Kernel 的 Profile–Optimize–Finalize 三阶段性能优化技能实战指南Liger Kernel Perf面向 Liger Kernel 的 Profile–Optimize–Finalize 三阶段性能优化技能实战指南 本文基于大模型模型优化深度学习突破训练瓶颈OpenRLHF Liger-Kernel GPU内核优化终极指南突破训练瓶颈OpenRLHF Liger Kernel GPU内核优化终极指南 OpenRLHF是首个结合Ray vLLM分布式架构与统一Agent设计范人工智能大模型强化学习RLHF分布式训练微调上一篇5 步解锁 Emby 全部 Premiere 功能硬件转码 主题定制完整指南下一篇.ncm、.qmcflac 出了 App 就放不了用 Unlock Music Electron 批量转成 MP3创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考