mcpField
W
warp-compile-time-optimizer

warp-compile-time-optimizer

github.com/nvidia开发3.8k99 次安装今天更新
代码仓库

关于 warp-compile-time-optimizer

用于测量和降低 NVIDIA Warp 工作负载冷启动编译耗时的 Skill。它通过隔离探针定位编译开销,并指导模块级调整,同时要求保留启动行为并验证结果。

主要功能

  1. 01测量冷启动与热启动编译
  2. 02识别重复的模块构建
  3. 03诊断 Warp 编译开销
  4. 04指导模块与选项调整
  5. 05对比测量并验证行为

适用场景

  1. 01排查 Warp 应用启动缓慢
  2. 02查找模块身份导致的重复编译
  3. 03将编译时间改动与基线比较

安装

复制下面的命令,在终端中运行。

$npx skills add https://github.com/nvidia/skills --skill warp-compile-time-optimizer

安装后在 Claude Code 中输入/warp-compile-time-optimizer确认已加载。

包含的文件

除 SKILL.md 外,安装时一起下载的参考资料和脚本。

BENCHMARK.md
config/
skillspector-baseline.yaml
evals/
config.yml
evals.json
evals/files/c01_particle_preprocessing/
pyproject.toml
evals/files/c01_particle_preprocessing/src/particle_prep/
__init__.py
__main__.py
pipeline.py
plugins.py
stages.py
evals/files/c04_batch_signals/
pyproject.toml
evals/files/c04_batch_signals/src/batch_signals/
__init__.py
__main__.py
conditioning.py
maintenance.py
runtime.py
evals/files/g01_pde_stepper/
pyproject.toml
evals/files/g01_pde_stepper/src/pde_stepper/
__init__.py
__main__.py
boundaries.py
operators.py
solver.py
evals/files/g02_stencil_bank/
pyproject.toml
evals/files/g02_stencil_bank/src/stencil_bank/
__init__.py
__main__.py
analysis.py
loss.py
stencils.py
evals/files/g04_image_pyramid/
pyproject.toml
evals/files/g04_image_pyramid/src/image_pyramid/
__init__.py

详细介绍

Warp 冷启动编译时间

从可运行的命令开始

探测程序会在子进程中运行目标程序,只写入临时目录,不需要网络、外部工具服务或 Warp 源码检出。

脚本用途参数
scripts/warp_compile_probe.py测量隔离的冷启动/热启动编译和启动。measure [OPTIONS] -- COMMAND...;使用 --help。

支持时,使用 run_script("scripts/warp_compile_probe.py", args=[...]); 否则使用下面的 Python 命令。目标程序必须运行至完成。

编译模型

Warp 编译的是模块,而不是单个内核。模块的身份由以下内容决定:

(live kernel & function set) x (module options) x (CUDA block_dim) x (generic instances)

每个身份都需要为整个模块生成代码并进行本机编译。

冷启动成本大致为:

number of distinct module identities you touch  x  size of each module

可通过两种方式降低成本:

  1. 停止身份频繁变更。模块加载后若发生变化,就会再次编译。
  2. 停止源码重复。模块在三个块维度下使用时,每个内核都会编译三次。

从仍需构建的模块中删除一个内核,只能节省一次构建的一部分。移除不必要的模块身份,则可节省完整的一次构建。

两者都不适用时,可并行执行互不依赖的 CUDA 构建(CS-13)。这只会改变工作的执行时机,不会减少编译量,因此应根据耗时来评估。

模块选项固定的时机

选项有两个设置时限:

  1. 模块创建时(Python 导入时)。 模块会从 warp.config 复制 enable_backward、 max_unroll、lineinfo、deterministic、deterministic_max_records 和 compile_time_trace。之后设置全局值会被该模块静默忽略。例外是 default_grid_stride。
  2. 首次加载时。 编译之前,使用 wp.set_module_options() 或 wp.get_module(name).options 修改现有模块。加载后修改选项会创建新的身份,并重新构建模块(CS-3)。
错过的时限现象成本
导入后设置 wp.config.*哈希未变,选项被静默忽略整个收益都被无声地错过
加载后设置模块选项哈希再次变化,模块构建两次额外构建一次,可在跟踪中看到

修改选项后,确认每个目标模块的哈希都发生了变化。哈希未变就表示选项并未生效。

保持行为不变

改变 Warp 编译代码的方式,而不是工作负载。

不要通过删除或合并内核来宣称性能提升。看似冗余的阶段可能用于保持所有权、别名行为、保留输出、数值边界或 API 行为。应在模块层面解决重复问题。

保留每次启动及其顺序、维度、数据类型、设备、块维度、梯度、数值模式、动态/插件行为和公共 API 签名。将定义移至模块作用域时,保留内核名称,因为日志、缓存产物和外部工具都会显示这些名称。

说明

1. 找出真实耗时,并确认瓶颈是编译

询问用户实际在等待哪个命令,然后进行冷启动测量:

python scripts/warp_compile_probe.py measure --samples 3 \
    --json baseline.json -- <the user's command>

探测程序为每个样本分配独立的 WARP_CACHE_PATH、WARP_CACHE_ROOT 和 CUDA_CACHE_PATH 目录,启用模块计时器并记录启动情况。它会在系统临时目录下创建这些目录,并自行清理,因此隔离样本时无需将缓存写入项目,也无需删除任何内容来重新测量。用同样的方法隔离手写样本。切勿使用 wp.clear_kernel_cache() 或 wp.clear_lto_cache() 清除正在使用的缓存;清除操作并非隔离操作,还可能影响其他进程。

先阅读探测输出,再查看源码。如果编译只占总耗时的一小部分,就报告真正的瓶颈并停止。对于库和测试,使用能编译工作负载模块的最精简命令。

如果模块都只编译过一次,且没有重复哈希、块维度变体或 LTO,就不存在结构性反复构建。这只能排除冗余构建,并不能说明构建规模过大;仍需检查缓存复用(CS-2)、反向代码生成(CS-10)、展开(CS-11)、预编译头文件(CS-12),以及存在多个 CUDA 模块时的重叠(CS-13)。

每个样本还会在刚填充的缓存上重新运行该命令。如果暖缓存下的模块处理耗时并未接近零,应先诊断缓存复用(CS-2),再更改模块结构。

2. 必要时询问运行时取舍

直接进行顺序调整、生命周期分组和选项提取,无需询问。更改 fast_math、max_unroll 或 MathDx/平铺实现前,请先询问:

其中一些选项可以缩短编译时间,但可能会拖慢编译出的内核或改变数值结果。您是在优化快速编辑运行循环(通常可以接受内核变慢),还是优化生产环境启动(通常不能接受)?

如果用户无法联系:

  • 不要更改会影响数值结果的选项,也不要替换实现。
  • 在移除反向代码生成之类的功能前,先测试该功能是否在使用;报告移除了什么、测得的收益以及如何还原。
  • 在应用入口点设置全局 wp.config.* 选项,不要在库代码中设置。

在第 6 步的记录表中,记录用户拒绝的选项及其测得的收益。

让更改范围与证据范围相匹配

性能分析支持的是对受测应用的更改,而非对共享库所有使用方的更改。仓库内搜索也无法找到仓库外或未来的调用方。例如,仅使用前向计算的应用,并不能证明应在求解器库中禁用梯度,因为其他应用可能会对该库进行微分。

在导入库之前,将选项限定到受测进程:

import warp as wp

wp.config.enable_backward = False   # must precede the library import

import the_library

这样也会覆盖应用加载的每个模块。CS-10 介绍了容易被忽略的导入顺序问题。如果只有修改库才能解决,请将测量结果发给库的维护者,由他们决定应采用的约定。

3. 根据测量结果诊断,而不是阅读源代码

探针会打印每个已编译模块的身份,包括其名称、哈希值、设备和块维度,然后指出哪些模块构建了多次。将看到的情况与下表匹配:

探针显示的情况含义查看位置
一个模块名称,多个 哈希值身份不断变化:模块首次加载后,其内核集、选项或泛型实例发生了变化CS-1、CS-3、CS-6
一个模块名称,多个 block_dim 值整个模块会针对每个块维度重新编译(CUDA)CS-5
一个功能中有许多单内核模块每个模块的固定开销被重复产生CS-4
每个内核都有一个以哈希命名的模块在稳定的内核上使用了 module="unique"CS-9
模块耗时与原生编译耗时之间差距很大,并且存在 .lto 构件MathDx/LTO 设置CS-7
本应命中预热缓存的运行显示 (compiled)缓存未被复用CS-2
模块加载后出现 “Failed to find module”CPU JIT 首次使用时发生并发竞态CS-8
生成的源代码很大,但没有重复构建问题展开预算CS-11
模块中有未被任何内容微分的伴随代码反向代码生成CS-10
各处编译都很慢,或者在工具包 13 以下的 CUDA 上有几个较小模块预编译头已关闭,或其收益不足以抵消开销CS-12
多个独立模块各构建一次,且 overlap_factor 接近 1.0构建逐个运行;默认关闭并行加载CS-13
设置的选项没有产生任何变化,且该模块的哈希值未变该选项是在模块创建后赋值的,因此从未应用到该模块“模块选项何时固定”
以上情况均未出现没有重复构建;开销来自构建本身的规模步骤 6

references/mechanisms.md 中每个机制都有对应章节,介绍如何确认、解决方法、局限性和失效模式。只阅读测量结果对应的章节。

4. 有意选择模块边界

仅在以下条件都相同时,才将内核分组到同一模块中:

  • 生命周期:它们同时定义、加载和失效;
  • 选项集:它们需要相同的 fast_math、enable_backward、max_unroll 和 MathDx 设置;
  • 在 CUDA 上,块维度稳定。

生命周期相同但稳定块维度不同的内核不应共用一个模块,因为每个块维度都会导致一次编译。生命周期独立的内核也应分开。

块维度在运行时变化的内核(例如根据输入大小选择)没有稳定的映射,因此应将它们保留在独立模块中,而不是让整个共享模块额外生成一个变体。

优先采用能够消除一次构建的最小改动。调整顺序和提前设置选项,比重新规划模块布局成本更低、风险更小;只有当每个模块的固定开销或块维度重复构建成为主要问题时,才重新分组。

wp.set_module_options() 的目标是调用它的 Python 模块,而不是显式指定了 module="pkg.name" 的内核。应使用真正的 Python 模块,或在命名模块加载前更新它的选项:

wp.set_module_options({"enable_backward": False})  # at module scope

wp.get_module("pkg.name").options.update({"enable_backward": False})

不要将 wp.get_module() 传给 wp.set_module_options(module=...),也不要在未指定 module="unique" 的情况下使用 @wp.kernel(module_options={...})。每内核的 enable_backward=False 有一个适用于 tile 模块的例外,详见 CS-10。每次更改选项后,都要确认模块哈希值。

5. 验证
python scripts/warp_compile_probe.py measure --samples 3 \
    --json candidate.json -- <the same command>
python scripts/warp_compile_probe.py compare baseline.json candidate.json

compare 会拒绝启动拓扑发生变化的情况,并将结果落在 max(1% of baseline, 2 x baseline MAD) 范围内的情况视为无法得出结论。

对于 BUILDS OVERLAPPED,应根据编译耗时而非模块计时器总和来判断调度变化。必需的预热轮次会提供该计时数据。请参阅 references/measurement.md。

然后检查探测器无法看到的内容:

  • 对比数值输出,并运行项目的测试或入口点。
  • 更改 enable_backward 或边界后,验证梯度路径。
  • 更改 fast_math、max_unroll、MathDx 或实现后,测量稳态运行时。
  • 测试未覆盖的动态内核、数据类型、配置和模式。
6. 报告结果

阅读 references/measurement.md 中的“报告结果”。报告:

用通俗语言描述每项优化:说明行为、证据和效果。例如,写“在首次加载前移动模块选项,以避免不必要的重建”,而不要写“应用了 CS-3”。将 CS-* 标签视为内部导航辅助信息,而非面向用户的说明。

选项测量结果未采用的原因采用方法
pkg.solver 上的 enable_backward=False冷启动 −38%活跃的 tape 会遍历这些内核在入口点设置,然后重新检查伴随量
max_unroll=4−2%,处于噪声范围内更改了生成的代码,却没有带来可测量的收益—
  • 优化前后的中位数和样本数量;
  • 减少幅度和剩余耗时,并结合最初的问题规模进行评估;
  • 已修复、排除、测量后放弃或尚未涉及的机制;
  • 未验证的权衡和行为;
  • 上述记录表中已放弃或未完全验证的选项;
  • 放宽任何阻碍修复的约束后测得的数值。

只陈述证据所支持的内容。“没有结构性开销”并不意味着“最优”或“不可再降低”。实际可行时,应测量已放弃的调整项;任何估算都应标明尚未测试。在报告不存在可行修复方案之前,请检查 CS-13。若考虑拆分模块,应先测量一个使用相同选项的单内核模块,以确定重复的固定开销。

故障排查

调试探测器之前,请直接运行目标命令。有关缓存和噪声问题,请参阅 references/measurement.md;有关特定机制的故障,请参阅 references/mechanisms.md。

限制

以下两条规则优先于任何性能提升:

  • 每个冷样本都必须隔离 Warp 和 CUDA 缓存。
  • 当加载可能针对 CPU 时,令 max_workers <= 1,包括 device=None 和混合设备列表。并发的 CPU 首次加载可能导致内核丢失;重试也无法保证安全。仅加载 CUDA 不受影响。

测量结果因环境而异:冷启动耗时会随 CPU、GPU、驱动程序、工具链和 Warp 版本变化。机制可以迁移;数值则不能。作出广泛结论之前,请阅读 references/mechanisms.md 中的已知未知项。

参考文件

  • references/mechanisms.md:十三种编译时机制,每种都包含确认信号、修复方法、适用范围限制和故障模式。请阅读测量结果指向的章节。
  • references/measurement.md:测量规程、各项指标的含义及其局限、报告指南、日志示例和手动测量方法。

评论

支持 Markdown