diff options
| author | Julius Ikkala <julius.ikkala@gmail.com> | 2025-10-16 21:53:10 +0300 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2025-10-16 18:53:10 +0000 |
| commit | 537697d3aa21b418c348b1017005603668b6a4cb (patch) | |
| tree | 4ad71cfe4ab2c51f2071d8171ba7bd98929a776a /source/slang/slang-emit.cpp | |
| parent | 0257cb001b6f5c31e4ac0435d546fe638a17c48a (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.cpp | 8 |
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); } |
