From 54faa55c0bd4c4beede7337a76ed3a56d1eb4f15 Mon Sep 17 00:00:00 2001
From: "Harsh Aggarwal (NVIDIA)" <haaggarwal@nvidia.com>
Date: Thu, 18 Sep 2025 21:16:44 +0530
Subject: Fix CUDA global variable initialization with constructor calls
 (#8340)

Fix CUDA global variable initialization with constructor calls

Resolves CUDA compilation failure where global variables with struct
constructor
initialization generated illegal `__device__` variable runtime
initialization.

  **Problem:**
  ```cuda
  // Generated invalid CUDA code:
  __device__ static const Stuff_0 gStuff_0 = Stuff_x24init_0(args...);
// Error: "dynamic initialization is not supported for a __device__
variable"

  Root Cause Discovered:
Through extensive debugging, found that
moveGlobalVarInitializationToEntryPoints
pass only handled kIROp_GlobalVar instructions, but global constants
with
  constructor calls appeared as kIROp_Call instructions at module scope.

  Solution:
1. IR Pipeline Fix: Extended moveGlobalVarInitializationToEntryPoints to
detect
and transform module-level constructor calls into proper global
variables with
  entry-point initialization
2. Field Access Fix: Enhanced kIROp_FieldExtract logic to emit correct
->
  syntax for pointer types and address-of operations
3. Constructor Emission: Added CUDA-specific handling for constructor
calls

  Architecture:
- Transforms let %gStuff = call %Constructor(...) into kernel context
initialization
- Moves runtime initialization from global scope to entry-point
execution
  - Follows CUDA best practices for global state management

  Files:
- source/slang/slang-ir-explicit-global-init.cpp: Extended IR
transformation pass
- source/slang/slang-emit-c-like.cpp: Enhanced field access and foldable
value logic
- source/slang/slang-emit-cuda.cpp: Added CUDA-specific field extraction
handling

  Result:
  // Now generates proper CUDA code:
  struct KernelContext_0 { Stuff_0 gStuff_1; };
  // Runtime initialization in entry point:
  kernelContext_1.gStuff_1 = constructor_call();

  Fixes: tests/compute/type-legalize-global-with-init.slang
---
 tests/compute/type-legalize-global-with-init.slang | 1 +
 1 file changed, 1 insertion(+)

(limited to 'tests/compute')

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;
-- 
cgit v1.2.3