summaryrefslogtreecommitdiff
path: root/source/slang/slang-emit.cpp
diff options
context:
space:
mode:
authorJulius Ikkala <julius.ikkala@gmail.com>2025-10-16 21:53:10 +0300
committerGitHub <noreply@github.com>2025-10-16 18:53:10 +0000
commit537697d3aa21b418c348b1017005603668b6a4cb (patch)
tree4ad71cfe4ab2c51f2071d8171ba7bd98929a776a /source/slang/slang-emit.cpp
parent0257cb001b6f5c31e4ac0435d546fe638a17c48a (diff)
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.
Diffstat (limited to 'source/slang/slang-emit.cpp')
-rw-r--r--source/slang/slang-emit.cpp8
1 files changed, 7 insertions, 1 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);
}