用 IR 传递测试包验证 Slang 编译器:`-dump-ir`、多目标 emit 与诊断测试的实战方法
2026/9/20 0:22:33 网站建设 项目流程
  • 编译器
  • 图形学
  • 编程语言

【免费下载链接】slang

Making it easier to work with shaders

项目地址:https://gitcode.com/GitHub_Trending/sl/slang
点击查看免费下载

导读

本文围绕 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 valuesslang-ir-link.cpp、slang-ir-check-recursion.cpp 等
SSA 构建与基础清理(SSA construction and basic cleanup)DCE、SCCP、Peephole、Simplify CFG、Single return、Init local varslang-ir-dce.cpp 等
特化与泛型(Specialization and generics)Specialize(specializeModule)、Bind existentials、AnyValue marshallingslang-ir-specialize.cpp 等
微分 / 自动微分(Differentiation)Autodiff 驱动、Forward mode、Reverse mode、Unzipslang-ir-autodiff-*.cpp
类型与值合法化(Type and value legalization)Lower optional type、Lower tuple types、Byte-address legalize、Lower enum typeslang-ir-lower-optional-type.cpp 等
内联与调用图(Inlining and call-graph)Inline、DLL export、DLL import、Defer buffer loadslang-ir-inline.cppslang-ir-dll-export.cpp
入口点与参数处理(Entry-point and parameter handling)Entry-point decorations、Collect entry-point uniforms、Transform params to constrefslang-ir-entry-point-*.cpp
布局与绑定(Layout and binding)Layout、Collect global uniforms、Translate global varying var、String hashslang-ir-collect-global-uniforms.cpp 等
循环变换(Loop transformations)Loop unroll、Loop inversion、Fuse satcoop、Synthesize active maskslang-ir-loop-unroll.cpp
目标特化降级(Target-specific lowering)GLSL/SPIR-V/Metal/WGSL legalize、CUDA immutable loadslang-ir-cuda-immutable-load.cpp 等
插桩(Instrumentation)Coverage instrument、Insert debug value store、Obfuscate locslang-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)与性能剖析对所有传递一视同仁地生效。其工作流程为:

  1. linkIR(slang-ir-link.cpp)把各翻译单元的 IR 模块链接到一起,得到后续所有传递共同修改的LinkedIR
  2. 依次执行一批验证(validation)、特化(specialization)与合法化(legalization);
  3. TargetRequest指明需要目标特有传递时插入它们(如 HLSL ray-payload 合法化、GLSL 入口点合法化、Metal 合法化、WGSL 合法化、CUDA immutable-load 处理等);
  4. 把链接并优化后的模块交还给调用方,由其转交给选定的发射后端。

一个关键例外:SPIR-V 的主合法化并不在linkAndOptimizeIR内部legalizeIRForSPIRV(定义于 slang-ir-spirv-legalize.cpp)由 SPIR-V 发射后端在linkAndOptimizeIR返回之后调用,以便能在地址空间特化之后运行。由于它是普通函数调用而非SLANG_PASS(...)-dump-ir输出中永远不会出现BEFORE legalizeSPIRV/AFTER legalizeSPIRV。这也是测试包中一个有趣的“负向”测试素材——-dump-ir只显示linkAndOptimizeIRSLANG_PASS调用点,直接函数调用运行的阶段在该流中不可见(尽管它确实执行了)。

另一个结构性事实:传递序列不是固定的清单。不同目标、不同优化级别、是否存在微分或覆盖率插桩,都会改变序列。多数可选传递由RequiredLoweringPassSet标志门控,该集合由calcRequiredLoweringPassSet计算,且其标志是累积的而非两次扫描间重置。

核心翻译规则:把“文档声明”变成“可观察测试”

本测试包最重要的方法论是声明到观察(claims to observations)的翻译规则。设计文档的每行分类表描述“传递本身做了什么”,但测试只能验证“读者能看到它做了什么”。二者并不等价——大量条目是分析类传递(构建旁路结构而不改动 IR),或者其产物仅供后续传递消费,没有任何独立的可观察足迹

规范(_prompt.md)给出了四种翻译模板:

  1. “传递 X 移除操作码 Y”→ 用-dump-ir-before X -dump-ir-after X -target <文本目标> -o -编译,用 FileCheck 断言 Y 出现在### BEFORE X:中、不出现在### AFTER X:中;
  2. “传递 X 是目标 T 特有的是”→ 对文档列出的目标编译到文本目标并 FileCheck 重写结果 Y;再对传递不运行的目标编译,FileCheck 断言原始操作码仍在;
  3. “传递 X 强制不变量 Z”→ 编写违反 Z 的源码,用DIAGNOSTIC_TESTFileCheck 诊断(如递归检查、缺 return);
  4. “传递 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++ 传递入口函数名(如specializeModuleeliminateDeadCodesimplifyIR)。头行格式为### 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.cppOptional<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 legalizeRWStructuredBuffer<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 模式中只能匹配裸子串kernelbuffer(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 returnslang-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 文本
特化与泛型specializeModulespecialize(%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 ShaderOpEntryPoint 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 传递工具模块CloneDominatorsUtilInsts infoInsts 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.slangwgsl-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 模式中禁止出现字面量[[...]](使用裸 tokenkernelbuffer(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” 一节为希望扩展传递管线的开发者提供了六步流程,与本测试包一脉相承:

  1. 在 source/slang 下创建slang-ir-<name>.{h,cpp}文件对,头文件声明公共入口函数(如void doMyPass(IRModule*, ...));
  2. 复用 slang-ir.h、slang-ir-clone.cpp、slang-ir-dominators.cpp 等工具,而不是编写一次性遍历;
  3. linkAndOptimizeIR中插入调用点,使用局部SLANG_PASS宏而非直接调用,这样传递才能参与 IR 转储与剖析;pre-link 传递放进generateIRForTranslationUnit,后端所属传递放进该后端的合法化编排器;
  4. 目标特有传递用TargetRequest门控,仅在特定构造出现时才需要的传递在RequiredLoweringPassSet中加标志;
  5. 在 tests 下添加测试,典型为COMPARE_COMPUTEINTERPRET,若传递会产生错误则加DIAGNOSTIC_TEST
  6. 提交前运行./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

项目地址:https://gitcode.com/GitHub_Trending/sl/slang
点击查看免费下载

相关推荐

创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考

需要专业的网站建设服务?

联系我们获取免费的网站建设咨询和方案报价,让我们帮助您实现业务目标

立即咨询