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

* Enable all dynamic dispatch tests on CUDA.

* Fix expected cross-compile test results.

6 files changed, 16 insertions, 4 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: