From 83514bd25160a9af91abc1b9acd7e44657447526 Mon Sep 17 00:00:00 2001
From: Yong He <yonghe@outlook.com>
Date: Mon, 21 Sep 2020 08:27:10 -0700
Subject: Enable all dynamic dispatch tests on CUDA. (#1552)

* Enable all dynamic dispatch tests on CUDA.

* Fix expected cross-compile test results.
---
 source/slang/slang-emit-cuda.cpp | 11 ++++++++++-
 1 file changed, 10 insertions(+), 1 deletion(-)

(limited to 'source/slang/slang-emit-cuda.cpp')

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