Skip to content

Commit 38a86e9

Browse files
Milky2018Milky2018
committed
Make Naga writer const reachability structural
1 parent 7c09f2d commit 38a86e9

2 files changed

Lines changed: 250 additions & 1 deletion

File tree

‎ir/wgsl_naga_compat_derived_view.mbt‎

Lines changed: 110 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -55,6 +55,7 @@ fn WgslIrNagaCompatPlanner::derived_module_view(
5555
builder.import_event(event)
5656
}
5757
builder.append_remaining()
58+
builder.import_entry_points()
5859
builder.view
5960
}
6061

@@ -79,14 +80,28 @@ fn WgslIrNagaDerivedModuleBuilder::import_event(
7980
}
8081
FullModule | AliasModule => {
8182
self.import_event_symbols_by_kind(symbols, None, Type)
82-
self.import_event_symbols_by_kind(symbols, None, Constant)
83+
self.import_full_module_constant_event_symbols(symbols)
8384
self.import_event_symbols_by_kind(symbols, None, Override)
8485
self.import_event_symbols_by_kind(symbols, None, GlobalVariable)
8586
self.import_event_symbols_by_kind(symbols, None, Function)
8687
}
8788
}
8889
}
8990

91+
///|
92+
fn WgslIrNagaDerivedModuleBuilder::import_full_module_constant_event_symbols(
93+
self : WgslIrNagaDerivedModuleBuilder,
94+
symbols : Array[WgslIrImportArenaSymbol],
95+
) -> Unit {
96+
for symbol in symbols {
97+
if symbol.kind() == Constant {
98+
self.import_named_non_generated_constant_if_present(
99+
symbol.generated_symbol_name(),
100+
)
101+
}
102+
}
103+
}
104+
90105
///|
91106
fn WgslIrNagaDerivedModuleBuilder::import_event_symbols_by_kind(
92107
self : WgslIrNagaDerivedModuleBuilder,
@@ -163,6 +178,9 @@ fn WgslIrNagaDerivedModuleBuilder::append_remaining(
163178
self.import_type(index)
164179
}
165180
for index in 0..<self.planner.shader_module.constants.items.length() {
181+
if self.planner.shader_module.constants.items[index].generated_import {
182+
continue
183+
}
166184
self.import_constant(index)
167185
}
168186
for index in 0..<self.planner.shader_module.overrides.items.length() {
@@ -214,6 +232,26 @@ fn WgslIrNagaDerivedModuleBuilder::import_named_constant_if_present(
214232
}
215233
}
216234

235+
///|
236+
fn WgslIrNagaDerivedModuleBuilder::import_named_non_generated_constant_if_present(
237+
self : WgslIrNagaDerivedModuleBuilder,
238+
name : String,
239+
) -> Unit {
240+
for index in 0..<self.planner.shader_module.constants.items.length() {
241+
if self.seen_constants.contains(index) {
242+
continue
243+
}
244+
match self.planner.shader_module.constants.items[index] {
245+
{ name: Some(constant_name), generated_import: false, .. } if constant_name ==
246+
name => {
247+
self.import_constant(index)
248+
return
249+
}
250+
_ => ()
251+
}
252+
}
253+
}
254+
217255
///|
218256
fn WgslIrNagaDerivedModuleBuilder::import_named_override_if_present(
219257
self : WgslIrNagaDerivedModuleBuilder,
@@ -351,6 +389,7 @@ fn WgslIrNagaDerivedModuleBuilder::import_constant(
351389
self.import_type_handle(constant.ty)
352390
self.seen_constants.add(index)
353391
self.view.constants.push(index)
392+
self.import_global_expression(constant.init)
354393
}
355394
None => ()
356395
}
@@ -490,6 +529,76 @@ fn WgslIrNagaDerivedModuleBuilder::import_function_inner(
490529
visiting.pop() |> ignore
491530
}
492531

532+
///|
533+
fn WgslIrNagaDerivedModuleBuilder::import_entry_points(
534+
self : WgslIrNagaDerivedModuleBuilder,
535+
) -> Unit {
536+
for index in 0..<self.planner.shader_module.entry_points.length() {
537+
match self.planner.filter {
538+
Some(filter) => if !filter.contains_entry_point(index) { continue }
539+
None => ()
540+
}
541+
self.import_entry_point(self.planner.shader_module.entry_points[index])
542+
}
543+
}
544+
545+
///|
546+
fn WgslIrNagaDerivedModuleBuilder::import_entry_point(
547+
self : WgslIrNagaDerivedModuleBuilder,
548+
entry_point : EntryPoint,
549+
) -> Unit {
550+
let visiting : Array[Int] = []
551+
for argument in entry_point.function.arguments {
552+
self.import_type_handle(argument.ty)
553+
}
554+
match entry_point.function.result {
555+
Some(result) => self.import_type_handle(result.ty)
556+
None => ()
557+
}
558+
for local_var in entry_point.function.local_variables.items {
559+
self.import_type_handle(local_var.ty)
560+
}
561+
match entry_point.workgroup_size_overrides {
562+
Some(overrides) =>
563+
for item in overrides {
564+
match item {
565+
Some(expr) => self.import_global_expression(expr)
566+
None => ()
567+
}
568+
}
569+
None => ()
570+
}
571+
match entry_point.mesh_info {
572+
Some(mesh_info) => {
573+
match mesh_info.max_vertices_override {
574+
Some(expr) => self.import_global_expression(expr)
575+
None => ()
576+
}
577+
match mesh_info.max_primitives_override {
578+
Some(expr) => self.import_global_expression(expr)
579+
None => ()
580+
}
581+
self.import_type_handle(mesh_info.vertex_output_type)
582+
self.import_type_handle(mesh_info.primitive_output_type)
583+
self.import_global(mesh_info.output_variable.index())
584+
}
585+
None => ()
586+
}
587+
match entry_point.task_payload {
588+
Some(handle) => self.import_global(handle.index())
589+
None => ()
590+
}
591+
match entry_point.incoming_ray_payload {
592+
Some(handle) => self.import_global(handle.index())
593+
None => ()
594+
}
595+
self.import_function_block(
596+
entry_point.function,
597+
entry_point.function.body,
598+
visiting,
599+
)
600+
}
601+
493602
///|
494603
fn WgslIrNagaDerivedModuleBuilder::import_global_expression(
495604
self : WgslIrNagaDerivedModuleBuilder,

‎ir/wgsl_naga_compat_derived_view_wbtest.mbt‎

Lines changed: 140 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -302,6 +302,146 @@ test "Naga writer arena uses derived constant import order" {
302302
debug_inspect(names.join(","), content="\"callee_dependency,source_first\"")
303303
}
304304

305+
///|
306+
test "Naga writer arena does not root unused full-module generated constants" {
307+
let source =
308+
#|const DEP: u32 = 1u;
309+
#|const USED: u32 = DEP;
310+
#|const UNUSED: u32 = 2u;
311+
#|
312+
#|fn selected() -> u32 {
313+
#| return USED;
314+
#|}
315+
#|
316+
let event = WgslIrImportArenaEvent(
317+
0,
318+
"root.wgsl",
319+
"dep.wgsl",
320+
"dep",
321+
"",
322+
[],
323+
FullModule,
324+
[
325+
WgslIrImportArenaSymbol(Constant, "DEP", "DEP"),
326+
WgslIrImportArenaSymbol(Constant, "USED", "USED"),
327+
WgslIrImportArenaSymbol(Constant, "UNUSED", "UNUSED"),
328+
WgslIrImportArenaSymbol(Function, "selected", "selected"),
329+
],
330+
)
331+
let module_ = parse_wgsl_module_to_ir_with_generated_imports_and_import_arena_events(
332+
source,
333+
[
334+
WgslIrGeneratedImportProvenance(
335+
"DEP",
336+
"dep.wgsl",
337+
"DEP",
338+
ImportedSourceSymbol,
339+
false,
340+
),
341+
WgslIrGeneratedImportProvenance(
342+
"USED",
343+
"dep.wgsl",
344+
"USED",
345+
ImportedSourceSymbol,
346+
false,
347+
),
348+
WgslIrGeneratedImportProvenance(
349+
"UNUSED",
350+
"dep.wgsl",
351+
"UNUSED",
352+
ImportedSourceSymbol,
353+
false,
354+
),
355+
WgslIrGeneratedImportProvenance(
356+
"selected",
357+
"dep.wgsl",
358+
"selected",
359+
ImportedSourceSymbol,
360+
false,
361+
),
362+
],
363+
[event],
364+
) catch {
365+
error => abort(error.message())
366+
}
367+
let planner = WgslIrNagaCompatPlanner::{
368+
shader_module: module_,
369+
filter: None,
370+
}
371+
let arena = planner.writer_arena()
372+
let names : Array[String] = []
373+
for slot in arena.constants {
374+
match slot.item.name {
375+
Some(name) => names.push(name)
376+
None => ()
377+
}
378+
}
379+
debug_inspect(names.join(","), content="\"USED,DEP\"")
380+
}
381+
382+
///|
383+
test "Naga writer arena imports entry point generated constant dependencies" {
384+
let source =
385+
#|const ENTRY_FLAG: u32 = 1u;
386+
#|const UNUSED_FLAG: u32 = 2u;
387+
#|
388+
#|@compute @workgroup_size(1)
389+
#|fn main() {
390+
#| if ENTRY_FLAG != 0u {
391+
#| return;
392+
#| }
393+
#|}
394+
#|
395+
let event = WgslIrImportArenaEvent(
396+
0,
397+
"root.wgsl",
398+
"dep.wgsl",
399+
"dep",
400+
"",
401+
[],
402+
FullModule,
403+
[
404+
WgslIrImportArenaSymbol(Constant, "ENTRY_FLAG", "ENTRY_FLAG"),
405+
WgslIrImportArenaSymbol(Constant, "UNUSED_FLAG", "UNUSED_FLAG"),
406+
],
407+
)
408+
let module_ = parse_wgsl_module_to_ir_with_generated_imports_and_import_arena_events(
409+
source,
410+
[
411+
WgslIrGeneratedImportProvenance(
412+
"ENTRY_FLAG",
413+
"dep.wgsl",
414+
"ENTRY_FLAG",
415+
ImportedSourceSymbol,
416+
false,
417+
),
418+
WgslIrGeneratedImportProvenance(
419+
"UNUSED_FLAG",
420+
"dep.wgsl",
421+
"UNUSED_FLAG",
422+
ImportedSourceSymbol,
423+
false,
424+
),
425+
],
426+
[event],
427+
) catch {
428+
error => abort(error.message())
429+
}
430+
let planner = WgslIrNagaCompatPlanner::{
431+
shader_module: module_,
432+
filter: None,
433+
}
434+
let arena = planner.writer_arena()
435+
let names : Array[String] = []
436+
for slot in arena.constants {
437+
match slot.item.name {
438+
Some(name) => names.push(name)
439+
None => ()
440+
}
441+
}
442+
debug_inspect(names.join(","), content="\"ENTRY_FLAG\"")
443+
}
444+
305445
///|
306446
test "Naga writer arena uses derived nested function-call order" {
307447
let source =

0 commit comments

Comments
 (0)