diff options
| -rw-r--r-- | source/slang/slang-emit.cpp | 9 | ||||
| -rw-r--r-- | tests/compute/type-legalize-global-with-init.slang | 1 |
2 files changed, 9 insertions, 1 deletions
diff --git a/source/slang/slang-emit.cpp b/source/slang/slang-emit.cpp index 6725ac3de..e1689ccfc 100644 --- a/source/slang/slang-emit.cpp +++ b/source/slang/slang-emit.cpp @@ -1264,10 +1264,17 @@ Result linkAndOptimizeIR( legalizeEmptyArray(irModule, sink); + // 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) + { + inlineGlobalConstantsForLegalization(irModule); + } + // We don't need the legalize pass for C/C++ based types if (options.shouldLegalizeExistentialAndResourceTypes) { - inlineGlobalConstantsForLegalization(irModule); // The Slang language allows interfaces to be used like // ordinary types (including placing them in constant diff --git a/tests/compute/type-legalize-global-with-init.slang b/tests/compute/type-legalize-global-with-init.slang index 7316cad1d..e0de2ac03 100644 --- a/tests/compute/type-legalize-global-with-init.slang +++ b/tests/compute/type-legalize-global-with-init.slang @@ -6,6 +6,7 @@ // //TEST(compute):COMPARE_COMPUTE: -shaderobj //TEST(compute):COMPARE_COMPUTE: -vk -shaderobj +//TEST(compute):COMPARE_COMPUTE: -cuda // //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer RWStructuredBuffer<uint> outputBuffer; |