Feature/fix cuda function preamble (#1187)

* Fix tests/compute/global-init.slang by handling some other cases where functions are emitted.

* Fix comment.

6 files changed, 14 insertions, 12 deletions

diff --git a/source/slang/slang-emit-c-like.cpp b/source/slang/slang-emit-c-like.cpp
index 8e0c64d1a..d630d18c9 100644
--- a/source/slang/slang-emit-c-like.cpp
+++ b/source/slang/slang-emit-c-like.cpp
@@ -2579,6 +2579,8 @@ void CLikeSourceEmitter::emitSimpleFuncImpl(IRFunc* func)
         emitEntryPointAttributes(func, entryPointDecor);
     }
 
+    emitFunctionPreambleImpl(func);
+
     auto name = getName(func);
 
     emitType(resultType, name);
@@ -2963,6 +2965,8 @@ void CLikeSourceEmitter::emitGlobalVar(IRGlobalVar* varDecl)
     String initFuncName;
     if (varDecl->getFirstBlock())
     {
+        emitFunctionPreambleImpl(varDecl);
+
         // A global variable with code means it has an initializer
         // associated with it. Eventually we'd like to emit that
         // initializer directly as an expression here, but for
diff --git a/source/slang/slang-emit-c-like.h b/source/slang/slang-emit-c-like.h
index 611678865..8381f194f 100644
--- a/source/slang/slang-emit-c-like.h
+++ b/source/slang/slang-emit-c-like.h
@@ -328,6 +328,7 @@ public:
     virtual void emitOperandImpl(IRInst* inst, EmitOpInfo const& outerPrec);
     virtual void emitParamTypeImpl(IRType* type, String const& name);
     virtual void emitIntrinsicCallExprImpl(IRCall* inst, IRTargetIntrinsicDecoration* targetIntrinsic, EmitOpInfo const& inOuterPrec);
+    virtual void emitFunctionPreambleImpl(IRInst* inst) { SLANG_UNUSED(inst); }
 
         // Only needed for glsl output with $ prefix intrinsics - so perhaps removable in the future
     virtual void emitTextureOrTextureSamplerTypeImpl(IRTextureTypeBase*  type, char const* baseName) { SLANG_UNUSED(type); SLANG_UNUSED(baseName); }
diff --git a/source/slang/slang-emit-cpp.cpp b/source/slang/slang-emit-cpp.cpp
index 759f980d3..99cc2f61c 100644
--- a/source/slang/slang-emit-cpp.cpp
+++ b/source/slang/slang-emit-cpp.cpp
@@ -808,7 +808,7 @@ void CPPSourceEmitter::_emitSignature(const UnownedStringSlice& funcName, const
     const int paramsCount = int(funcType->getParamCount());
     IRType* retType = specOp->returnType;
 
-    emitSpecializedOperationDefinitionPreamble(specOp);
+    emitFunctionPreambleImpl(nullptr);
 
     SourceWriter* writer = getSourceWriter();
 
@@ -981,7 +981,7 @@ void CPPSourceEmitter::_emitGetAtDefinition(const UnownedStringSlice& funcName,
     {
         UnownedStringSlice typePrefix = (i == 0) ? UnownedStringSlice::fromLiteral("const ") : UnownedStringSlice();
 
-        emitSpecializedOperationDefinitionPreamble(specOp);
+        emitFunctionPreambleImpl(nullptr);
 
         writer->emit(typePrefix);
         emitType(specOp->returnType);
@@ -1075,7 +1075,7 @@ void CPPSourceEmitter::_emitConstructConvertDefinition(const UnownedStringSlice&
     IRType* srcType = funcType->getParamType(1);
     IRType* retType = specOp->returnType;
 
-    emitSpecializedOperationDefinitionPreamble(specOp);
+    emitFunctionPreambleImpl(nullptr);
 
     emitType(retType);
     writer->emit(" ");
@@ -1143,7 +1143,7 @@ void CPPSourceEmitter::_emitConstructFromScalarDefinition(const UnownedStringSli
     IRType* srcType = funcType->getParamType(1);
     IRType* retType = specOp->returnType;
 
-    emitSpecializedOperationDefinitionPreamble(specOp);
+    emitFunctionPreambleImpl(nullptr);
 
     emitType(retType);
     writer->emit(" ");
diff --git a/source/slang/slang-emit-cpp.h b/source/slang/slang-emit-cpp.h
index 9ea083199..082497510 100644
--- a/source/slang/slang-emit-cpp.h
+++ b/source/slang/slang-emit-cpp.h
@@ -84,9 +84,7 @@ protected:
     virtual SlangResult calcTypeName(IRType* type, CodeGenTarget target, StringBuilder& out);
     virtual SlangResult calcFuncName(const HLSLIntrinsic* specOp, StringBuilder& out);
     virtual SlangResult calcScalarFuncName(HLSLIntrinsic::Op op, IRBasicType* type, StringBuilder& outBuilder);
-    virtual void emitSpecializedOperationDefinitionPreamble(const HLSLIntrinsic* specOp) { SLANG_UNUSED(specOp); }
-
-
+    
     void _maybeEmitSpecializedOperationDefinition(const HLSLIntrinsic* specOp);
 
     void _emitForwardDeclarations(const List<EmitAction>& actions);
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index 83ad4a0f8..2d49e606c 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -501,9 +501,7 @@ void CUDASourceEmitter::emitSimpleFuncParamsImpl(IRFunc* func)
 
 void CUDASourceEmitter::emitSimpleFuncImpl(IRFunc* func)
 {
-    // Mark as run on device. Don't need to worry about entry point, as that is output separately to call the __device_ implementation
-    m_writer->emit("__device__ ");    
-
+    // Skip the CPP impl - as it does some processing we don't need here for entry points.
     CLikeSourceEmitter::emitSimpleFuncImpl(func);
 }
 
diff --git a/source/slang/slang-emit-cuda.h b/source/slang/slang-emit-cuda.h
index e75eb4e88..392e4f98d 100644
--- a/source/slang/slang-emit-cuda.h
+++ b/source/slang/slang-emit-cuda.h
@@ -47,6 +47,8 @@ protected:
     virtual void emitMatrixLayoutModifiersImpl(IRVarLayout* layout) SLANG_OVERRIDE;
     virtual void emitOperandImpl(IRInst* inst, EmitOpInfo const&  outerPrec) SLANG_OVERRIDE;
     virtual void emitCall(const HLSLIntrinsic* specOp, IRInst* inst, const IRUse* operands, int numOperands, const EmitOpInfo& inOuterPrec) SLANG_OVERRIDE;
+    virtual void emitFunctionPreambleImpl(IRInst* inst) SLANG_OVERRIDE { SLANG_UNUSED(inst); m_writer->emit("__device__ "); }
+
 
     //virtual bool tryEmitGlobalParamImpl(IRGlobalParam* varDecl, IRType* varType) SLANG_OVERRIDE;
     virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) SLANG_OVERRIDE;
@@ -58,8 +60,7 @@ protected:
     // CPPSourceEmitter overrides 
     virtual SlangResult calcTypeName(IRType* type, CodeGenTarget target, StringBuilder& out) SLANG_OVERRIDE;
     virtual SlangResult calcScalarFuncName(HLSLIntrinsic::Op op, IRBasicType* type, StringBuilder& outBuilder) SLANG_OVERRIDE;
-    virtual void emitSpecializedOperationDefinitionPreamble(const HLSLIntrinsic* specOp) SLANG_OVERRIDE { SLANG_UNUSED(specOp); m_writer->emit("__device__ "); }
-
+    
     SlangResult _calcCUDATextureTypeName(IRTextureTypeBase* texType, StringBuilder& outName);
 };