- 编译器
- 图形学
- 编程语言
【免费下载链接】slang
Making it easier to work with shaders
导读
本文围绕 Slang 编译器(GitHub_Trending/sl/slang)的 IR 传递(IR pass)目录展开,讲解如何把 IR 传递设计文档 中的每一条声明翻译成可自动执行的测试:-dump-ir-before/-dump-ir-after观察 IR 变换前后差异、多目标SIMPLEemit 观察目标代码形态、DIAGNOSTIC_TEST验证校验类诊断。读完本文,你将掌握这套测试体系的三种观察模式、九大传递分类的核心可观察足迹(footprint),以及哪些声明注定无法通过slangcCLI 验证。本文的主体依据是 05-ir-passes 测试包生成规范,辅以 IR 传递设计文档 与 测试生成通用规则 中的实现细节。
IR 传递在 Slang 编译管线中的位置
Slang 的编译流程中,AST 经过前端降级生成 IR(04-ast-to-ir 设计文档),随后在进入代码发射阶段(06-emit 设计文档)之前,会执行一系列 IR 变换(IR pass)。这些变换统称IR 传递目录(IR-pass catalog),它们正是本测试包要验证的对象。
按 IR 传递设计文档 的统计,source_commit时刻的源码树中约有 162 个slang-ir-*.cpp文件,且每个文件都出现在分类表中;单个文件常常提供多个传递入口点。设计文档将它们归为九大类别:
| 类别 | 代表性传递 | 源码文件 |
|---|---|---|
| 链接与验证(Linking and validation) | Link、Validate、Check recursion、Missing return、Use of uninitialized values | slang-ir-link.cpp、slang-ir-check-recursion.cpp 等 |
| SSA 构建与基础清理(SSA construction and basic cleanup) | DCE、SCCP、Peephole、Simplify CFG、Single return、Init local var | slang-ir-dce.cpp 等 |
| 特化与泛型(Specialization and generics) | Specialize(specializeModule)、Bind existentials、AnyValue marshalling | slang-ir-specialize.cpp 等 |
| 微分 / 自动微分(Differentiation) | Autodiff 驱动、Forward mode、Reverse mode、Unzip | slang-ir-autodiff-*.cpp |
| 类型与值合法化(Type and value legalization) | Lower optional type、Lower tuple types、Byte-address legalize、Lower enum type | slang-ir-lower-optional-type.cpp 等 |
| 内联与调用图(Inlining and call-graph) | Inline、DLL export、DLL import、Defer buffer load | slang-ir-inline.cpp、slang-ir-dll-export.cpp等 |
| 入口点与参数处理(Entry-point and parameter handling) | Entry-point decorations、Collect entry-point uniforms、Transform params to constref | slang-ir-entry-point-*.cpp等 |
| 布局与绑定(Layout and binding) | Layout、Collect global uniforms、Translate global varying var、String hash | slang-ir-collect-global-uniforms.cpp 等 |
| 循环变换(Loop transformations) | Loop unroll、Loop inversion、Fuse satcoop、Synthesize active mask | slang-ir-loop-unroll.cpp等 |
| 目标特化降级(Target-specific lowering) | GLSL/SPIR-V/Metal/WGSL legalize、CUDA immutable load | slang-ir-cuda-immutable-load.cpp 等 |
| 插桩(Instrumentation) | Coverage instrument、Insert debug value store、Obfuscate loc | slang-ir-coverage-instrument.cpp等 |
注:设计文档中“分类表”的类别标题(如
#linking-and-validation、#ssa-construction-and-basic-cleanup)同时是测试锚点(anchor),每个测试的//META: doc_ref必须解析到这些标题。
编排器linkAndOptimizeIR:一次编译到底跑多少传递
设计文档明确指出,所有后链接(post-link)传递的编排器是 slang-emit.cpp 中的linkAndOptimizeIR(其函数体包含约 180 处SLANG_PASS(...)调用点)。该宏在 slang-emit.cpp 中定义,将每个传递调用包装进wrapPass,从而让IR 转储(dump)与性能剖析对所有传递一视同仁地生效。其工作流程为:
- 用
linkIR(slang-ir-link.cpp)把各翻译单元的 IR 模块链接到一起,得到后续所有传递共同修改的LinkedIR; - 依次执行一批验证(validation)、特化(specialization)与合法化(legalization);
- 当
TargetRequest指明需要目标特有传递时插入它们(如 HLSL ray-payload 合法化、GLSL 入口点合法化、Metal 合法化、WGSL 合法化、CUDA immutable-load 处理等); - 把链接并优化后的模块交还给调用方,由其转交给选定的发射后端。
一个关键例外:SPIR-V 的主合法化并不在linkAndOptimizeIR内部。legalizeIRForSPIRV(定义于 slang-ir-spirv-legalize.cpp)由 SPIR-V 发射后端在linkAndOptimizeIR返回之后调用,以便能在地址空间特化之后运行。由于它是普通函数调用而非SLANG_PASS(...),在-dump-ir输出中永远不会出现BEFORE legalizeSPIRV/AFTER legalizeSPIRV头。这也是测试包中一个有趣的“负向”测试素材——-dump-ir只显示linkAndOptimizeIR的SLANG_PASS调用点,直接函数调用运行的阶段在该流中不可见(尽管它确实执行了)。
另一个结构性事实:传递序列不是固定的清单。不同目标、不同优化级别、是否存在微分或覆盖率插桩,都会改变序列。多数可选传递由RequiredLoweringPassSet标志门控,该集合由calcRequiredLoweringPassSet计算,且其标志是累积的而非两次扫描间重置。
核心翻译规则:把“文档声明”变成“可观察测试”
本测试包最重要的方法论是声明到观察(claims to observations)的翻译规则。设计文档的每行分类表描述“传递本身做了什么”,但测试只能验证“读者能看到它做了什么”。二者并不等价——大量条目是分析类传递(构建旁路结构而不改动 IR),或者其产物仅供后续传递消费,没有任何独立的可观察足迹。
规范(_prompt.md)给出了四种翻译模板:
- “传递 X 移除操作码 Y”→ 用
-dump-ir-before X -dump-ir-after X -target <文本目标> -o -编译,用 FileCheck 断言 Y 出现在### BEFORE X:中、不出现在### AFTER X:中; - “传递 X 是目标 T 特有的是”→ 对文档列出的目标编译到文本目标并 FileCheck 重写结果 Y;再对传递不运行的目标编译,FileCheck 断言原始操作码仍在;
- “传递 X 强制不变量 Z”→ 编写违反 Z 的源码,用
DIAGNOSTIC_TESTFileCheck 诊断(如递归检查、缺 return); - “传递 X 在发射中有用户可观察结果”→ 许多传递最干净的观察点在目标文本而非 IR 转储中(collect-global-uniforms 在 HLSL/CUDA 上创建
GlobalParams结构体、byte-address legalize 引入asuint/Load形态的访问、optional/result 降级在所有目标上产出扁平结构体),这类用多目标SIMPLE覆盖。
观察模式一:-dump-ir-before/-dump-ir-after捕捉传递前后差异
IR 转储对比是验证“某个命名传递确实改变了 IR”的首选手段。以下面这份真实测试为例(specialize-module-substitutes-generic-with-concrete-type.slang):
//TEST:SIMPLE(filecheck=CHECK):-target spirv-asm -dump-ir-before specializeModule -dump-ir-after specializeModule -o - -stage compute -entry main uniform int a; uniform RWStructuredBuffer<int> buf; T pickFirst<T>(T x, T y) { return x; } [numthreads(1, 1, 1)] void main(uint3 tid : SV_DispatchThreadID) { buf[0] = pickFirst(a, a + 1); } // Before the pass: `pickFirst` is generic (Func over type variables). // CHECK-LABEL: ### BEFORE specializeModule: // CHECK: func %pickFirst{{[^:]*}}: Func(% // After the pass: `pickFirst` carries a concrete `Int, Int, Int` signature. // CHECK-LABEL: ### AFTER specializeModule: // CHECK: func %pickFirst{{[^:]*}}: Func(Int, Int, Int)要点总结(结合 _prompt.md 的 “Lessons captured” 与 _common.md):
-dump-ir-before/-dump-ir-after接受 C++ 传递入口函数名(如specializeModule、eliminateDeadCode、simplifyIR)。头行格式为### BEFORE <name>:/### AFTER <name>:。名字不匹配时头不会出现(FileCheck 报 “expected string not found”),所以应以对代表性源码跑一次-dump-ir得到的真实名字为准。simplifyIR会运行多次,因此会产生多个### BEFORE simplifyIR:/### AFTER simplifyIR:块。锚定用CHECK-LABEL指向首次出现处,而不是断言“唯一一次”。- 必须配合文本目标与
-o -:-dump-ir转储写到 stderr,目标文本写到 stdout,二者不会混流;-o -保证测试可移植(绝对路径在跨平台套件中不可复现)。 - IR 转储有巨大的前导内容(core-module 导入、能力集、微分胶水、autodiff witness 表),应在
func %main或用户命名函数处锚定,而不是从头匹配。 - pre-link 的 DCE 会先移除局部未使用的辅助函数:文件作用域定义但从未调用的
unusedHelper在首个转储阶段就已消失。要观察BEFORE eliminateDeadCode/AFTER eliminateDeadCode的差异,IR 中必须存在 post-link DCE 能移除的东西——通常更简单的方式是观察 emit 文本中不再包含该辅助函数。 - 验证类负向测试:
spirv-legalization-has-no-dump-ir-header.slang(文档对应测试)正是利用“SPIR-V 合法化不走SLANG_PASS”这一事实,断言转储流中没有该头。
观察模式二:多目标SIMPLEemit 验证发射形态
大量合法化/降级传递的“足迹”体现在最终目标文本中。规范要求:一个传递如果影响多个目标的发射,就应在每个可行的文本目标上测试,而不是只测 HLSL + SPIR-V。每个目标使用独立的filecheck=<NAME>前缀。以下面这份测试为例(lower-optional-type-emits-struct-with-has-value.slang):
//TEST:SIMPLE(filecheck=HLSL):-target hlsl -stage compute -entry main //TEST:SIMPLE(filecheck=GLSL):-target glsl -stage compute -entry main //TEST:SIMPLE(filecheck=CUDA):-target cuda -stage compute -entry main //TEST:SIMPLE(filecheck=CPP):-target cpp -stage compute -entry main uniform int a; uniform RWStructuredBuffer<int> buf; Optional<int> makeOpt(int v) { if (v > 0) return v; return none; } [numthreads(1, 1, 1)] void main(uint3 tid : SV_DispatchThreadID) { Optional<int> r = makeOpt(a); if (r.hasValue) buf[0] = r.value; else buf[0] = -1; } // HLSL: hasValue // GLSL: hasValue // CUDA: hasValue // CPP: hasValue该测试验证slang-ir-lower-optional-type.cpp把Optional<T>降级为带value+hasValue字段的结构体(HLSL 上形如_slang_Optional_int_0),其形态标记 token 就是hasValue。类似的多目标“发射形状”断言遍布整个测试包,例如:
- DCE(dce-removes-unused-helper-from-emit.slang):文件作用域定义但从未调用的
unusedHelper不进入任何目标的 emit——CHECK-NOT: unusedHelper在所有文本目标上都成立,这是清理类最干净的观察模式; - Collect global uniforms:模块作用域的
uniform int a;在 HLSL 上成为GlobalParams_0的字段、GLSL 上成为block_GlobalParams_0uniform block 的字段、CUDA 上成为SLANG_globalParams的字段(collect-global-uniforms-builds-global-params-struct.slang); - GLSL legalize:
RWStructuredBuffer<T>重写为layout(std430, binding = N) buffer { T _data[]; }形态(glsl-legalize-emits-std430-ssbo-shape.slang); - CUDA immutable load:可变的全局内存读取(
StructuredBuffer<T>、ConstantBuffer<T>内容)被包进__ldg(&...);而顶层uniform位于__constant__的SLANG_globalParams,不应被包装(cuda-immutable-load-wraps-buffer-read-in-ldg.slang); - Metal legalize:入口点获得位置性
[[buffer(N)]]/kernel标记——注意按 _common.md 的要求,FileCheck 模式中只能匹配裸子串kernel、buffer(0),绝不能用字面量[[kernel]](FileCheck 会把[[...]]解析为正则变量); - WGSL legalize:插入
@binding(N) @group(N)与@compute @workgroup_size(...); - 循环展开(loop-unroll-removes-static-loop-on-glsl.slang):
[unroll]作用于编译期已知迭代次数的循环;HLSL 发射器会有意保留for(;;)字面形态并打上[unroll]标记(真正的展开交给下游 DXC),而在 GLSL/CUDA/SPIR-V 上由 Slang 自身完全展开——因此应挑一个观察无歧义的目标(如 GLSL 上断言不再有for关键字)进行验证。
观察模式三:DIAGNOSTIC_TEST验证校验类传递
对文档明确点名且以诊断为其表面后果的验证传递,使用诊断测试:
//DIAGNOSTIC_TEST:SIMPLE(diag=CHECK):-target hlsl -stage compute -entry main适用场景(均来自 _prompt.md):
- Check recursion(slang-ir-check-recursion.cpp):诊断不支持的递归,负向测试覆盖自递归(recursion-check-rejects-self-recursion.slang)、互递归(recursion-check-rejects-mutual-recursion.slang)与三函数环;
- Missing return(
slang-ir-missing-return.cpp):非 void 函数缺失返回路径时告警(missing-return-warns-on-non-void.slang); - Use of uninitialized values / Detect uninitialized resources:诊断未初始化使用;若在语义检查阶段而非 IR 阶段就能产生干净诊断,则该阶段应跳过;
- 其余校验类还包括 check-optional-none-usage、operator-shift-overflow、check-differentiability、mesh-output 读取、late 的
validateVectorsAndMatrices矩阵维度上限检查等。
规范特别强调:箭头(caret)位置必须以测试运行器报告的 “Suggested annotations” 为准,不要手数列号;诊断文本只注释编译器实际输出的内容,拿不准措辞时就断言稳定的错误码(如E41030)。
逐分类的可观察足迹详解
规范按设计文档的类别标题逐条给出了“可写测试的声明”(Observable claims)。下表汇总各分类的核心观察点:
| 分类 | 可观察声明 | 观察方式 |
|---|---|---|
| 链接与验证 | 递归检查、缺 return 告警、未初始化使用诊断 | DIAGNOSTIC_TEST |
| SSA 构建与基础清理 | DCE 移除未调用且未导出的顶层辅助函数(CHECK-NOT: helper);SPIR-V 文本中OpStore初始化局部变量(init-local-var) | 多目标SIMPLE/ emit 文本 |
| 特化与泛型 | specializeModule前specialize(%generic, T),后具体函数体(无generic残留);接口类型局部量从具体类型初始化时调用点被解析(无lookup_witness) | -dump-ir-before/-after+ emit |
| 类型与值合法化 | Optional<int>在 HLSL/GLSL/CUDA 上降级为含value/hasValue的结构体;Tuple<int,int>扁平化为_0/_1访问;RWStructuredBuffer<T>各目标发射形态 | 多目标SIMPLE(本类最丰富) |
| 内联与调用图 | [ForceInline]函数体出现在main的 emit 中而非helper(...)调用;DLL export 装饰影响 SPIR-VOpEntryPoint形态 | emit 文本 |
| 入口点与参数处理 | [numthreads(...)]在 HLSL 发射[numthreads]、GLSL 发射layout(local_size_x = ...)、SPIR-V 发射OpExecutionMode ... LocalSize;模块作用域uniform int a;被收集进GlobalParams_0(HLSL)/block_GlobalParams_0(GLSL)/SLANG_globalParams(CUDA) | 多目标SIMPLE |
| 布局与绑定 | SV_DispatchThreadID输入按目标消费方式不同(观察点在 emit 文本) | emit 文本(覆盖较弱,避免与 06-emit 重复) |
| 循环变换 | [unroll]常量界循环展开后无循环构造残留(GLSL/CUDA/SPIR-V);HLSL 保留[unroll]字面形态 | 单目标定点断言 |
| 目标特化降级 | GLSL SSBO 形态layout(std430, binding = N) buffer { T _data[]; };SPIR-VOpCapability Shader、OpEntryPoint GLCompute;Metalkernel/buffer(0);WGSL@binding/@group/@compute @workgroup_size;CUDA__ldg(&...)边界(buffer 正例、顶层 uniform 反例) | 单目标 / 多目标SIMPLE |
| 插桩与其他 | 除-trace-coverage-binding等标志门控或内部脚手架外,除非点名了表面可观察声明,否则跳过 | — |
值得注意的是入口点装饰的传播路径:设计文档明确说明,没有任何传递“降低”[numthreads(...)]——AST→IR 降级把它作为IRNumThreadsDecoration附加到入口点函数上,原样保留到发射,由各后端各自写出启动维度标记(HLSLnumthreads(8, 1, 1)、GLSLlocal_size_x = 8、SPIR-VOpExecutionMode %main LocalSize 8 1 1、WGSL@workgroup_size(8)。中间的传递只读取它,唯一例外是 SPIR-V 合法化会为缺失该装饰的 compute 入口点补默认值(1, 1, 1)。这正是本测试包要表达的“是入口点传递导致了这个结果”这一角度(entry-point-decorations-emit-numthreads-marker.slang)。
注定无法通过 slangc 验证的声明(Untested claims)
规范明确列出应记录在## Untested claims中的内容,测试生成时不得为它们编造测试:
- 传递运行的确切顺序——设计文档明确把顺序交给
linkAndOptimizeIR与各目标管线页面(target-pipelines 索引),文本层面的顺序断言既脆弱也非文档声明; - 各传递内部的 C++ 辅助函数与类结构(如每个传递中的
IRBuilder用法)——内部 API; - 可提升指令是否浮动到最外层支配者——IR 转储展示的是提升后的文本,而非提升决策本身;
- IR 传递工具模块(
Clone、Dominators、Util、Insts info、Insts stable names)——它们不是变换,没有自身的可观察效果; - pre-link 区域(见 04b-pre-link-passes.md)——属于另一个测试包的管辖范围;
- 覆盖率插桩标志(
-trace-coverage-binding、-trace-coverage-reserved-space)——命令行表面,本包不覆盖; - 微分传递族(
slang-ir-autodiff-*.cpp)——需要[Differentiable]函数加bwd_diff/fwd_diff调用才能产出可观察输出,而设计文档对该传递族只说明了家族本身、未给出本包可干净锚定的用户可观察后果,因此跳过并记录为设计文档缺口。
测试包结构、规模与质量要求
本测试包是整套体系中规模最大的一个。生成规范规划的规模是 20–40 个测试、覆盖大多数类别,文件上限 100;而按当前仓库的实际情况,05-ir-passes 测试包 已包含144 个.slang测试文件(其size_cap_files为 150),覆盖从any-value-marshalling-pack-unpack-names.slang到wgsl-legalize-multiple-binding-groups.slang的广泛场景,并配套 README.md 记录声明清单、功能覆盖表、未测试声明与文档缺口。
质量要求(结合 _prompt.md 与 _common.md):
- 每个测试的
doc_ref必须解析到 05-ir-passes.md(或规范列出的少量次级文档)中的某个锚点; -dump-ir类测试必须使用-target <文本目标> -o -;- 多目标
SIMPLE测试每个目标使用独立的filecheck=<NAME>前缀与对应 CHECK 前缀; - FileCheck 模式中禁止出现字面量
[[...]](使用裸 tokenkernel、buffer(0)、thread_position_in_grid); - CUDA 算术观察的操作数应来自
SV_DispatchThreadID等线程/调度 ID,而非uniform全局量; - Metal
[[buffer(N)]]下标是位置性的,不得断言由 HLSLregister(u3)或vk::binding(...)驱动的具体下标; - 校验类诊断测试使用运行器 “Suggested annotations” 提供的列位置;
- 任何测试都不依赖 GPU;
- 任何测试都不断言传递运行顺序;
- 不得通过查看未覆盖源码行来编写测试(行为声明必须来自文档);
## Doc gaps observed记录文档作出但本包无法测试的声明(尤其微分族与合法化长尾)。
给贡献者的补充:如何新增一个 IR 传递
设计文档的 “Adding a new pass” 一节为希望扩展传递管线的开发者提供了六步流程,与本测试包一脉相承:
- 在 source/slang 下创建
slang-ir-<name>.{h,cpp}文件对,头文件声明公共入口函数(如void doMyPass(IRModule*, ...)); - 复用 slang-ir.h、slang-ir-clone.cpp、slang-ir-dominators.cpp 等工具,而不是编写一次性遍历;
- 在
linkAndOptimizeIR中插入调用点,使用局部SLANG_PASS宏而非直接调用,这样传递才能参与 IR 转储与剖析;pre-link 传递放进generateIRForTranslationUnit,后端所属传递放进该后端的合法化编排器; - 目标特有传递用
TargetRequest门控,仅在特定构造出现时才需要的传递在RequiredLoweringPassSet中加标志; - 在 tests 下添加测试,典型为
COMPARE_COMPUTE或INTERPRET,若传递会产生错误则加DIAGNOSTIC_TEST; - 提交前运行
./extras/formatting.sh。
相关资源
- 测试包生成规范:05-ir-passes/_prompt.md
- IR 传递设计文档:05-ir-passes.md
- 测试生成通用规则:docs/generated/tests/_meta/prompts/_common.md
- 测试包 README(声明清单与覆盖表):05-ir-passes/README.md
- 相邻阶段:AST→IR 降级(04-ast-to-ir.md)、pre-link 传递(04b-pre-link-passes.md)、代码发射(06-emit.md)、各目标管线(target-pipelines/index.md)
- 核心实现:slang-emit.cpp(
linkAndOptimizeIR编排器)、slang-ir-link.cpp(链接)、slang-ir-specialize.cpp(特化驱动)、slang-ir-dce.cpp(死代码消除)、slang-ir-lower-optional-type.cpp(Optional 降级)、slang-ir-cuda-immutable-load.cpp(CUDA__ldg包装)
- 编译器
- 图形学
- 编程语言
【免费下载链接】slang
Making it easier to work with shaders
相关推荐
Slang 编译器 AST 到 IR 的 Lowering 阶段:映射规则与基于 `-dump-ir` 的测试验证
Slang 编译器 AST 到 IR 的 Lowering 阶段:映射规则与基于 dump ir 的测试验证 导读 本文围绕 Slang(Shader Lang
编译器图形学编程语言Slang 编译器 IR 杂项操作码测试集:以文档断言为锚的 IR 回归测试方法论
Slang 编译器 IR 杂项操作码测试集:以文档断言为锚的 IR 回归测试方法论 本文围绕 Slang 仓库中的测试生成提示词 docs/generated/
编译器图形学编程语言Slang 编译器 AST→IR Lowering 阶段白盒表征测试全解析:coverage/lower-to-ir 测试束实战指南
Slang 编译器 AST→IR Lowering 阶段白盒表征测试全解析:coverage/lower to ir 测试束实战指南 导读:本文深入解析 Sla
编译器图形学编程语言
创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考