2 files changed, 42 insertions, 0 deletions
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<uint> outputBuffer;
+
+T getGlobalInvocationID<T: IInteger>(T value)
+{
+    return T(gl_GlobalInvocationID.x) + value;
+}
+
+T getWaveLaneIndex<T: IInteger>(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);
+}
+
+