From 1fa4e486f598c5a7eed0db65f187ab95f890133c Mon Sep 17 00:00:00 2001 From: Darren Wihandi <65404740+fairywreath@users.noreply.github.com> Date: Mon, 31 Mar 2025 21:17:49 -0400 Subject: Fix compilation of global builtin variables inside generics (#6701) * Include generics' operands in call graph construction * add test --- source/slang/slang-ir-call-graph.cpp | 1 + tests/glsl/builtin-inside-generics.slang | 41 ++++++++++++++++++++++++++++++++ 2 files changed, 42 insertions(+) create mode 100644 tests/glsl/builtin-inside-generics.slang diff --git a/source/slang/slang-ir-call-graph.cpp b/source/slang/slang-ir-call-graph.cpp index 47b18be2e..a709bd418 100644 --- a/source/slang/slang-ir-call-graph.cpp +++ b/source/slang/slang-ir-call-graph.cpp @@ -88,6 +88,7 @@ void buildEntryPointReferenceGraph( case kIROp_GlobalParam: case kIROp_GlobalVar: case kIROp_SPIRVAsmOperandBuiltinVar: + case kIROp_Generic: addToWorkList({entryPoint, operand}); break; } diff --git a/tests/glsl/builtin-inside-generics.slang b/tests/glsl/builtin-inside-generics.slang new file mode 100644 index 000000000..fc5caeb6a --- /dev/null +++ b/tests/glsl/builtin-inside-generics.slang @@ -0,0 +1,41 @@ +//TEST:SIMPLE(filecheck=CHECK_SPIRV): -entry main -stage compute -target spirv -allow-glsl +//TEST:SIMPLE(filecheck=CHECK_GLSL): -entry main -stage compute -target glsl -allow-glsl +//TEST:SIMPLE(filecheck=CHECK_HLSL): -entry main -stage compute -target hlsl -allow-glsl +//TEST:SIMPLE(filecheck=CHECK_METAL): -entry main -stage compute -target metal -allow-glsl +//TEST:SIMPLE(filecheck=CHECK_WGSL): -entry main -stage compute -target wgsl -allow-glsl + +RWStructuredBuffer outputBuffer; + +T getGlobalInvocationID(T value) +{ + return T(gl_GlobalInvocationID.x) + value; +} + +T getWaveLaneIndex(T value) +{ + return T(WaveGetLaneIndex()) + value; +} + +// CHECK_SPIRV: GlobalInvocationId +// CHECK_SPIRV: SubgroupLocalInvocationId + +// CHECK_GLSL: gl_GlobalInvocationID +// CHECK_GLSL: gl_SubgroupInvocationID + +// CHECK_HLSL: WaveGetLaneIndex() +// CHECK_HLSL: SV_DispatchThreadID + +// CHECK_METAL: thread_position_in_grid +// CHECK_METAL: thread_index_in_simdgroup + +// CHECK_WGSL: global_invocation_id +// CHECK_WGSL: subgroup_invocation_id + +[shader("compute")] +void main() +{ + outputBuffer[0U] = getGlobalInvocationID(0U); + outputBuffer[1U] = getWaveLaneIndex(0U); +} + + -- cgit v1.2.3