Inline global constants for shader style CPU targets (#8686)

On the shader-host-callable target, test `gh-4874.slang` generates IR
that contains global constants referencing global params. These need to
get inlined into functions, as otherwise
`introduceExplicitGlobalContext()` will fail with "no outer func at use
site for global", making the test crash the compiler.

5 files changed, 19 insertions, 2 deletions

diff --git a/source/slang/slang-emit.cpp b/source/slang/slang-emit.cpp
index d5e89b1fe..bdd7dfe10 100644
--- a/source/slang/slang-emit.cpp
+++ b/source/slang/slang-emit.cpp
@@ -1270,7 +1270,13 @@ Result linkAndOptimizeIR(
     // For CUDA targets, always inline global constants to avoid dynamic initialization
     // of __device__ variables rejected by NVRTC. This runs independently of the broader
     // resource/existential type legalization, which remains disabled for CUDA.
-    if (target == CodeGenTarget::CUDASource || options.shouldLegalizeExistentialAndResourceTypes)
+    //
+    // We also need this pass on the CPU targets in shader mode, as global
+    // constants may reference global parameters, which can't be emitted as
+    // constants.
+    if (target == CodeGenTarget::CUDASource ||
+        (isCPUTarget(targetRequest) && isKernelTarget(target)) ||
+        options.shouldLegalizeExistentialAndResourceTypes)
     {
         inlineGlobalConstantsForLegalization(irModule);
     }
diff --git a/source/slang/slang-ir-legalize-global-values.cpp b/source/slang/slang-ir-legalize-global-values.cpp
index 55676f2d1..f7152d561 100644
--- a/source/slang/slang-ir-legalize-global-values.cpp
+++ b/source/slang/slang-ir-legalize-global-values.cpp
@@ -129,6 +129,7 @@ bool GlobalInstInliningContextGeneric::isInlinableGlobalInst(IRInst* inst)
     case kIROp_Neq:
     case kIROp_Eql:
     case kIROp_Call:
+    case kIROp_Load:
         return true;
     default:
         if (isInlinableGlobalInstForTarget(inst))
diff --git a/source/slang/slang-target.h b/source/slang/slang-target.h
index 8fa1f83f6..f9f653406 100644
--- a/source/slang/slang-target.h
+++ b/source/slang/slang-target.h
@@ -83,6 +83,9 @@ bool isCPUTarget(TargetRequest* targetReq);
 bool isWGPUTarget(TargetRequest* targetReq);
 bool isWGPUTarget(CodeGenTarget target);
 
+// Are we generating code for a Kernel-style target (as opposed to host-style target)
+bool isKernelTarget(CodeGenTarget codeGenTarget);
+
 /// A request to generate output in some target format.
 class TargetRequest : public RefObject
 {
diff --git a/source/slang/slang-type-layout.cpp b/source/slang/slang-type-layout.cpp
index 5bbcd2eb1..2daf33fdf 100644
--- a/source/slang/slang-type-layout.cpp
+++ b/source/slang/slang-type-layout.cpp
@@ -2712,6 +2712,12 @@ bool isWGPUTarget(TargetRequest* targetReq)
     return isWGPUTarget(targetReq->getTarget());
 }
 
+bool isKernelTarget(CodeGenTarget codeGenTarget)
+{
+    return ArtifactDescUtil::makeDescForCompileTarget(asExternal(codeGenTarget)).style ==
+           ArtifactStyle::Kernel;
+}
+
 SourceLanguage getIntermediateSourceLanguageForTarget(TargetProgram* targetProgram)
 {
     // If we are emitting directly, there is no intermediate source language
diff --git a/tests/bugs/gh-4874.slang b/tests/bugs/gh-4874.slang
index 403f6fc50..53a2aaf38 100644
--- a/tests/bugs/gh-4874.slang
+++ b/tests/bugs/gh-4874.slang
@@ -6,6 +6,7 @@
 //
 //TEST(compute):COMPARE_COMPUTE(filecheck-buffer=CHECK): -shaderobj
 //TEST(compute):COMPARE_COMPUTE(filecheck-buffer=CHECK): -shaderobj -vk
+//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=CHECK): -shaderobj -cpu
 //
 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
 RWStructuredBuffer<uint> outputBuffer;
@@ -47,4 +48,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
     // CHECK: 432
     // CHECK: 543
     // CHECK: 654
-}
\ No newline at end of file
+}