Enable all dynamic dispatch tests on CUDA. (#1552)

* Enable all dynamic dispatch tests on CUDA.

* Fix expected cross-compile test results.

22 files changed, 40 insertions, 42 deletions

diff --git a/source/slang/slang-emit-c-like.cpp b/source/slang/slang-emit-c-like.cpp
index f234b0be6..1a424306a 100644
--- a/source/slang/slang-emit-c-like.cpp
+++ b/source/slang/slang-emit-c-like.cpp
@@ -999,7 +999,8 @@ bool CLikeSourceEmitter::shouldFoldInstIntoUseSites(IRInst* inst)
     // if target langauge doesn't support pointers.
     if(as<IRPtrTypeBase>(type))
     {
-        return !doesTargetSupportPtrTypes();
+        if (!doesTargetSupportPtrTypes())
+            return true;
     }
 
     // First we check for uniform parameter groups,
diff --git a/source/slang/slang-emit-c-like.h b/source/slang/slang-emit-c-like.h
index b89d5d1c4..05cffd053 100644
--- a/source/slang/slang-emit-c-like.h
+++ b/source/slang/slang-emit-c-like.h
@@ -254,7 +254,7 @@ public:
     UInt getCallablePayloadLocation(IRInst* inst);
 
         /// Emit modifiers that should apply even for a declaration of an SSA temporary.
-    void emitTempModifiers(IRInst* temp);
+    virtual void emitTempModifiers(IRInst* temp);
 
     void emitVarModifiers(IRVarLayout* layout, IRInst* varDecl, IRType* varType);
 
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index ecc3eb68b..b029a8aa6 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -79,6 +79,15 @@ static bool _isSingleNameBasicType(IROp op)
     }
 }
 
+void CUDASourceEmitter::emitTempModifiers(IRInst* temp)
+{
+    CPPSourceEmitter::emitTempModifiers(temp);
+    if (as<IRModuleInst>(temp->getParent()))
+    {
+        m_writer->emit("__device__ ");
+    }
+}
+
 SlangResult CUDASourceEmitter::_calcCUDATextureTypeName(IRTextureTypeBase* texType, StringBuilder& outName)
 {
     // Not clear how to do this yet
@@ -324,7 +333,7 @@ String CUDASourceEmitter::generateEntryPointNameImpl(IREntryPointDecoration* ent
 
 void CUDASourceEmitter::emitGlobalRTTISymbolPrefix()
 {
-    m_writer->emit("__device__");
+    m_writer->emit("__device__ ");
 }
 
 void CUDASourceEmitter::emitCall(const HLSLIntrinsic* specOp, IRInst* inst, const IRUse* operands, int numOperands, const EmitOpInfo& inOuterPrec)
diff --git a/source/slang/slang-emit-cuda.h b/source/slang/slang-emit-cuda.h
index 18c6e86a2..9378453ba 100644
--- a/source/slang/slang-emit-cuda.h
+++ b/source/slang/slang-emit-cuda.h
@@ -34,6 +34,7 @@ public:
     static UnownedStringSlice getVectorPrefix(IROp op);
 
     virtual RefObject* getExtensionTracker() SLANG_OVERRIDE { return m_extensionTracker; }
+    virtual void emitTempModifiers(IRInst* temp) SLANG_OVERRIDE;
 
     CUDASourceEmitter(const Desc& desc) :
         Super(desc),
diff --git a/source/slang/slang-ir-link.cpp b/source/slang/slang-ir-link.cpp
index cc4c7e68d..f40f83fb8 100644
--- a/source/slang/slang-ir-link.cpp
+++ b/source/slang/slang-ir-link.cpp
@@ -1480,7 +1480,7 @@ LinkedIR linkIR(
             cloneValue(context, bindInst);
         }
     }
-    if (target == CodeGenTarget::CPPSource)
+    if (target == CodeGenTarget::CPPSource || target == CodeGenTarget::CUDASource)
     {
         for (IRModule* irModule : irModules)
         {
diff --git a/source/slang/slang-ir.cpp b/source/slang/slang-ir.cpp
index 28fa70edf..e9d90adfe 100644
--- a/source/slang/slang-ir.cpp
+++ b/source/slang/slang-ir.cpp
@@ -5274,6 +5274,7 @@ namespace Slang
         case kIROp_RTTIType:
         case kIROp_Func:
         case kIROp_Generic:
+        case kIROp_Var:
         case kIROp_GlobalVar: // Note: the IRGlobalVar represents the *address*, so only a load/store would have side effects
         case kIROp_GlobalConstant:
         case kIROp_GlobalParam:
diff --git a/tests/bugs/vk-image-atomics.slang.glsl b/tests/bugs/vk-image-atomics.slang.glsl
index 383b396ed..16dffd8dc 100644
--- a/tests/bugs/vk-image-atomics.slang.glsl
+++ b/tests/bugs/vk-image-atomics.slang.glsl
@@ -9,10 +9,8 @@ out vec4 _S1;
 
 void main()
 {
-    const uint _S2 = uint(1);
-
-    uint _S3;
-    _S3 = imageAtomicAdd(t_0, ivec2(uvec2(0)), _S2);
-    _S1 = vec4(_S3);
+    uint _S2;
+    _S2 = imageAtomicAdd(t_0, ivec2(uvec2(0)), 1);
+    _S1 = vec4(_S2);
     return;
 }
diff --git a/tests/compute/dynamic-dispatch-10.slang b/tests/compute/dynamic-dispatch-10.slang
index 19f734c47..db0ad8fc1 100644
--- a/tests/compute/dynamic-dispatch-10.slang
+++ b/tests/compute/dynamic-dispatch-10.slang
@@ -1,5 +1,5 @@
 //TEST(compute):COMPARE_COMPUTE:-cpu -xslang -disable-specialization
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
 
 // Test dynamic dispatch code gen for specializing a generic with
 // an existential value.
diff --git a/tests/compute/dynamic-dispatch-11.slang b/tests/compute/dynamic-dispatch-11.slang
index 18b416697..c21f88b7f 100644
--- a/tests/compute/dynamic-dispatch-11.slang
+++ b/tests/compute/dynamic-dispatch-11.slang
@@ -1,7 +1,7 @@
 // Test using interface typed shader parameters with dynamic dispatch.
 
 //TEST(compute):COMPARE_COMPUTE:-cpu -xslang -disable-specialization
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
 
 [anyValueSize(8)]
 interface IInterface
diff --git a/tests/compute/dynamic-dispatch-12.slang b/tests/compute/dynamic-dispatch-12.slang
index 67598033d..cd122ec56 100644
--- a/tests/compute/dynamic-dispatch-12.slang
+++ b/tests/compute/dynamic-dispatch-12.slang
@@ -1,7 +1,7 @@
 // Test using interface typed shader parameters with dynamic dispatch.
 
 //TEST(compute):COMPARE_COMPUTE:-cpu
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda
 
 [anyValueSize(8)]
 interface IInterface
diff --git a/tests/compute/dynamic-dispatch-13.slang b/tests/compute/dynamic-dispatch-13.slang
index 723f42f52..4e523e483 100644
--- a/tests/compute/dynamic-dispatch-13.slang
+++ b/tests/compute/dynamic-dispatch-13.slang
@@ -1,7 +1,7 @@
 // Test using interface typed shader parameters wrapped inside a `StructuredBuffer`.
 
 //TEST(compute):COMPARE_COMPUTE:-cpu
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda
 
 [anyValueSize(8)]
 interface IInterface
diff --git a/tests/compute/dynamic-dispatch-2.slang b/tests/compute/dynamic-dispatch-2.slang
index 1bbc8edf4..8fa96db23 100644
--- a/tests/compute/dynamic-dispatch-2.slang
+++ b/tests/compute/dynamic-dispatch-2.slang
@@ -1,5 +1,5 @@
 //TEST(compute):COMPARE_COMPUTE:-cpu -xslang -disable-specialization
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
 
 // Test dynamic dispatch code gen for static member functions
 // of associated type.
diff --git a/tests/compute/dynamic-dispatch-3.slang b/tests/compute/dynamic-dispatch-3.slang
index 2e1e2076a..851251ffd 100644
--- a/tests/compute/dynamic-dispatch-3.slang
+++ b/tests/compute/dynamic-dispatch-3.slang
@@ -1,5 +1,5 @@
 //TEST(compute):COMPARE_COMPUTE:-cpu -xslang -disable-specialization
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
 
 // Test dynamic dispatch code gen for static member functions
 // of associated type.
diff --git a/tests/compute/dynamic-dispatch-4.slang b/tests/compute/dynamic-dispatch-4.slang
index bc60df1b0..b2d20020e 100644
--- a/tests/compute/dynamic-dispatch-4.slang
+++ b/tests/compute/dynamic-dispatch-4.slang
@@ -1,5 +1,5 @@
 //TEST(compute):COMPARE_COMPUTE:-cpu -xslang -disable-specialization
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
 
 // Test dynamic dispatch code gen for generic-typed local variables.
 
diff --git a/tests/compute/dynamic-dispatch-5.slang b/tests/compute/dynamic-dispatch-5.slang
index c0a022325..88c26eee7 100644
--- a/tests/compute/dynamic-dispatch-5.slang
+++ b/tests/compute/dynamic-dispatch-5.slang
@@ -1,5 +1,5 @@
 //TEST(compute):COMPARE_COMPUTE:-cpu -xslang -disable-specialization
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
 
 // Test dynamic dispatch code gen for general `This` type.
 [anyValueSize(8)]
diff --git a/tests/compute/dynamic-dispatch-6.slang b/tests/compute/dynamic-dispatch-6.slang
index 99f6c9e82..8becc9381 100644
--- a/tests/compute/dynamic-dispatch-6.slang
+++ b/tests/compute/dynamic-dispatch-6.slang
@@ -1,5 +1,5 @@
 //TEST(compute):COMPARE_COMPUTE:-cpu -xslang -disable-specialization
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
 
 // Test dynamic dispatch code gen for generic-typed return values.
 [anyValueSize(8)]
diff --git a/tests/compute/dynamic-dispatch-7.slang b/tests/compute/dynamic-dispatch-7.slang
index a13062aa0..eb79f03ef 100644
--- a/tests/compute/dynamic-dispatch-7.slang
+++ b/tests/compute/dynamic-dispatch-7.slang
@@ -1,5 +1,5 @@
 //TEST(compute):COMPARE_COMPUTE:-cpu -xslang -disable-specialization
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
 
 // Test dynamic dispatch code gen for associated-typed return values
 // and local variables.
diff --git a/tests/compute/dynamic-dispatch-8.slang b/tests/compute/dynamic-dispatch-8.slang
index a2a93525e..440be385e 100644
--- a/tests/compute/dynamic-dispatch-8.slang
+++ b/tests/compute/dynamic-dispatch-8.slang
@@ -1,5 +1,5 @@
 //TEST(compute):COMPARE_COMPUTE:-cpu -xslang -disable-specialization
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
 
 // Test dynamic dispatch code gen for extential type parameters.
 
diff --git a/tests/compute/dynamic-dispatch-9.slang b/tests/compute/dynamic-dispatch-9.slang
index 786bb217b..d69707d5d 100644
--- a/tests/compute/dynamic-dispatch-9.slang
+++ b/tests/compute/dynamic-dispatch-9.slang
@@ -1,5 +1,5 @@
 //TEST(compute):COMPARE_COMPUTE:-cpu -xslang -disable-specialization
-//DISABLE_TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
+//TEST(compute):COMPARE_COMPUTE:-cuda -xslang -disable-specialization
 
 // Test dynamic dispatch code gen for initializing an extential value
 // from a generic value.
diff --git a/tests/cross-compile/sv-coverage.slang.glsl b/tests/cross-compile/sv-coverage.slang.glsl
index 04728e81b..2a8ff0734 100644
--- a/tests/cross-compile/sv-coverage.slang.glsl
+++ b/tests/cross-compile/sv-coverage.slang.glsl
@@ -9,11 +9,8 @@ in vec4 _S2;
 
 void main()
 {
-    uint _S3 = uint(gl_SampleMaskIn[0]);
-    uint _S4;
-
-    _S4 = _S3 ^ uint(1);
+    uint _S3 = uint(gl_SampleMaskIn[0]) ^ uint(1);
     _S1 = _S2;
-    gl_SampleMask[0] = int(_S4);
+    gl_SampleMask[0] = int(_S3);
     return;
 }
diff --git a/tests/slang-extension/atomic-float-byte-address-buffer-cross.slang.glsl b/tests/slang-extension/atomic-float-byte-address-buffer-cross.slang.glsl
index 6e21156aa..f5f575287 100644
--- a/tests/slang-extension/atomic-float-byte-address-buffer-cross.slang.glsl
+++ b/tests/slang-extension/atomic-float-byte-address-buffer-cross.slang.glsl
@@ -38,12 +38,9 @@ layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;void main()
 
     float delta_0 = ((anotherBuffer_0)._data[(uint(idx_0 & 3))]);
 
-
-    uint _S13 = uint(idx_0 << 2);
-
 #line 21
-    float _S14;
-    RWByteAddressBuffer_InterlockedAddF32_0(_S13, 1.00000000000000000000, _S14);
+    float _S13;
+    RWByteAddressBuffer_InterlockedAddF32_0(uint(idx_0 << 2), 1.00000000000000000000, _S13);
     RWByteAddressBuffer_InterlockedAddF32_1(uint(int(tid_0 >> 2) << 2), delta_0);
 
 #line 13
diff --git a/tests/vkray/intersection.slang.glsl b/tests/vkray/intersection.slang.glsl
index 09d7e63a5..66846d993 100644
--- a/tests/vkray/intersection.slang.glsl
+++ b/tests/vkray/intersection.slang.glsl
@@ -9,12 +9,10 @@
 #define tmp_direction _S4
 #define tmp_tmin _S5
 #define tmp_tmax _S6
-#define tmp_ray _S7
-#define tmp_sphere _S8
-#define tmp_thit _S9
-#define tmp_hitattrs _S10
-#define tmp_dithit _S11
-#define tmp_reportresult _S12
+#define tmp_thit _S7
+#define tmp_hitattrs _S8
+#define tmp_dithit _S9
+#define tmp_reportresult _S10
 
 struct Sphere_0
 {
@@ -83,13 +81,9 @@ void main()
     float tmp_tmax = gl_RayTmaxNV;
     ray_1.TMax_0 = tmp_tmax;
 
-    RayDesc_0 tmp_ray = ray_1;
-
-    Sphere_0 tmp_sphere = U_0._data.gSphere_0;
-
     float tmp_thit;
     SphereHitAttributes_0 tmp_hitattrs;
-    bool tmp_dithit = rayIntersectsSphere_0(tmp_ray, tmp_sphere, tmp_thit, tmp_hitattrs);
+    bool tmp_dithit = rayIntersectsSphere_0(ray_1, U_0._data.gSphere_0, tmp_thit, tmp_hitattrs);
 
     float tHit_2 = tmp_thit;
     SphereHitAttributes_0 attrs_1 = tmp_hitattrs;