CUDA support improvements (#1168)

* Add test result for compile-to-cuda

* Add RAII for some CUDA types to simplify usage.

* First pass handling of some instrinsics on CUDA (for example transcendentals)

* CUDA working with built in intrinsics.

* Add missing CUDA prelude intrinsics.

* CUDA matches CPU output on simple-cross-compile.slang

* First pass at hlsl-scalar-float-intrinsic.slang test.

* Fix smoothstep impl on CUDA and CPU.

* Fixed step intrinsic on CUDA/CPU.

* Added operator[] to Matrix for C++, to allow row access.
Needs a fix for CUDA.

* Fixed warning on clang build.

8 files changed, 256 insertions, 105 deletions

diff --git a/source/core/slang-nvrtc-compiler.cpp b/source/core/slang-nvrtc-compiler.cpp
index 1bb2669b8..bc7d1f4f6 100644
--- a/source/core/slang-nvrtc-compiler.cpp
+++ b/source/core/slang-nvrtc-compiler.cpp
@@ -276,7 +276,10 @@ SlangResult NVRTCDownstreamCompiler::compile(const CompileOptions& options, RefP
         cmdLine.addArg("-I");
         cmdLine.addArg(include);
     }
-    
+
+    {
+        cmdLine.addArg("-std=c++14");
+    }
 
     nvrtcProgram program = nullptr;
     nvrtcResult res = m_nvrtcCreateProgram(&program, options.sourceContents.getBuffer(), options.sourceContentsPath.getBuffer(), 0, nullptr, nullptr);
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index 7c88e530f..22a846eb7 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -1240,7 +1240,9 @@ __generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> sin(vector<T,
 __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> sin(matrix<T,N,M> x);
 
 // Sine and cosine
-__generic<T : __BuiltinFloatingPointType> void sincos(T x, out T s, out T c);
+__generic<T : __BuiltinFloatingPointType>
+__target_intrinsic(glsl, "$1 = sin($0); $2 = cos($0);")
+void sincos(T x, out T s, out T c);
 __generic<T : __BuiltinFloatingPointType, let N : int> void sincos(vector<T,N> x, out vector<T,N> s, out vector<T,N> c);
 __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> void sincos(matrix<T,N,M> x, out matrix<T,N,M> s, out matrix<T,N,M> c);
 
diff --git a/source/slang/hlsl.meta.slang.h b/source/slang/hlsl.meta.slang.h
index db0fc2285..0abae51b0 100644
--- a/source/slang/hlsl.meta.slang.h
+++ b/source/slang/hlsl.meta.slang.h
@@ -1316,7 +1316,9 @@ SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> si
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> sin(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Sine and cosine\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType> void sincos(T x, out T s, out T c);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
+SLANG_RAW("__target_intrinsic(glsl, \"$1 = sin($0); $2 = cos($0);\")\n")
+SLANG_RAW("void sincos(T x, out T s, out T c);\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> void sincos(vector<T,N> x, out vector<T,N> s, out vector<T,N> c);\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> void sincos(matrix<T,N,M> x, out matrix<T,N,M> s, out matrix<T,N,M> c);\n")
 SLANG_RAW("\n")
@@ -1577,7 +1579,7 @@ for (int aa = 0; aa < kBaseBufferAccessLevelCount; ++aa)
 
     sb << "};\n";
 }
-SLANG_RAW("#line 1504 \"hlsl.meta.slang\"")
+SLANG_RAW("#line 1506 \"hlsl.meta.slang\"")
 SLANG_RAW("\n")
 SLANG_RAW("\n")
 SLANG_RAW("\n")
diff --git a/source/slang/slang-emit-cpp.cpp b/source/slang/slang-emit-cpp.cpp
index 848ebd6e0..db442d131 100644
--- a/source/slang/slang-emit-cpp.cpp
+++ b/source/slang/slang-emit-cpp.cpp
@@ -403,16 +403,16 @@ SlangResult CPPSourceEmitter::calcTypeName(IRType* type, CodeGenTarget target, S
         {
             auto vecType = static_cast<IRVectorType*>(type);
             auto vecCount = int(GetIntVal(vecType->getElementCount()));
-            const IROp elemType = vecType->getElementType()->op;
+            auto elemType = vecType->getElementType();
 
-            if (target == CodeGenTarget::CPPSource)
+            if (target == CodeGenTarget::CPPSource || target == CodeGenTarget::CUDASource)
             {
-                out << "Vector<" << getBuiltinTypeName(elemType) << ", " << vecCount << ">";
+                out << "Vector<" << _getTypeName(elemType) << ", " << vecCount << ">";
             }
             else
             {             
                 out << "Vec";
-                UnownedStringSlice postFix = _getCTypeVecPostFix(elemType);
+                UnownedStringSlice postFix = _getCTypeVecPostFix(elemType->op);
 
                 out << postFix;
                 if (postFix.size() > 1)
@@ -431,9 +431,9 @@ SlangResult CPPSourceEmitter::calcTypeName(IRType* type, CodeGenTarget target, S
             const auto rowCount = int(GetIntVal(matType->getRowCount()));
             const auto colCount = int(GetIntVal(matType->getColumnCount()));
 
-            if (target == CodeGenTarget::CPPSource)
+            if (target == CodeGenTarget::CPPSource || target == CodeGenTarget::CUDASource)
             {
-                out << "Matrix<" << getBuiltinTypeName(elementType->op) << ", " << rowCount << ", " << colCount << ">";
+                out << "Matrix<" << _getTypeName(elementType) << ", " << rowCount << ", " << colCount << ">";
             }
             else
             {
@@ -800,6 +800,8 @@ void CPPSourceEmitter::_emitSignature(const UnownedStringSlice& funcName, const
     const int paramsCount = int(funcType->getParamCount());
     IRType* retType = specOp->returnType;
 
+    emitSpecializedOperationDefinitionPreamble(specOp);
+
     SourceWriter* writer = getSourceWriter();
 
     emitType(retType);
@@ -900,9 +902,19 @@ void CPPSourceEmitter::_emitCrossDefinition(const UnownedStringSlice& funcName,
     writer->indent();
 
     writer->emit("return ");
-    emitType(specOp->returnType);
-    writer->emit("{ a.y * b.z - a.z * b.y, a.z * b.x - a.x * b.z, a.x * b.y - a.y * b.x }; \n");
+    if (m_target == CodeGenTarget::CUDASource)
+    {
+        m_writer->emit("make_");
+        emitType(specOp->returnType);
+        writer->emit("( a.y * b.z - a.z * b.y, a.z * b.x - a.x * b.z, a.x * b.y - a.y * b.x ); \n");
+    }
+    else
+    {
+        emitType(specOp->returnType);
+        writer->emit("{ a.y * b.z - a.z * b.y, a.z * b.x - a.x * b.z, a.x * b.y - a.y * b.x }; \n");
+    }
 
+    
     writer->dedent();
     writer->emit("}\n\n");
 }
@@ -912,7 +924,7 @@ UnownedStringSlice CPPSourceEmitter::_getAndEmitSpecializedOperationDefinition(H
     HLSLIntrinsic intrinsic;
     m_intrinsicSet.calcIntrinsic(op, retType, argTypes, argCount, intrinsic);
     auto specOp = m_intrinsicSet.add(intrinsic);
-    emitSpecializedOperationDefinition(specOp);
+    _maybeEmitSpecializedOperationDefinition(specOp);
     return  _getFuncName(specOp);
 }
 
@@ -1184,15 +1196,19 @@ void CPPSourceEmitter::_emitReflectDefinition(const UnownedStringSlice& funcName
     writer->emit("}\n\n");
 }
 
-void CPPSourceEmitter::emitSpecializedOperationDefinition(const HLSLIntrinsic* specOp)
+void CPPSourceEmitter::_maybeEmitSpecializedOperationDefinition(const HLSLIntrinsic* specOp)
 {
-    typedef HLSLIntrinsic::Op Op;
-
     // Check if it's been emitted already, if not add it.
     if (!m_intrinsicEmitted.Add(specOp))
     {
         return;
     }
+    emitSpecializedOperationDefinition(specOp);
+}
+
+void CPPSourceEmitter::emitSpecializedOperationDefinition(const HLSLIntrinsic* specOp)
+{
+    typedef HLSLIntrinsic::Op Op;
 
     switch (specOp->op)
     {
@@ -1318,8 +1334,8 @@ void CPPSourceEmitter::emitCall(const HLSLIntrinsic* specOp, IRInst* inst, const
                     if (IRBasicType::isaImpl(retType->op))
                     {
                         SLANG_ASSERT(numOperands == 1);
-
-                        writer->emit(getBuiltinTypeName(retType->op));
+                        
+                        writer->emit(_getTypeName(retType));
                         writer->emitChar('(');
 
                         emitOperand(operands[0].get(), getInfo(EmitOp::General));
@@ -1432,16 +1448,29 @@ HLSLIntrinsic* CPPSourceEmitter::_addIntrinsic(HLSLIntrinsic::Op op, IRType* ret
     return addedIntrinsic;
 }
 
-StringSlicePool::Handle CPPSourceEmitter::_calcScalarFuncName(HLSLIntrinsic::Op op, IRBasicType* type)
+SlangResult CPPSourceEmitter::calcScalarFuncName(HLSLIntrinsic::Op op, IRBasicType* type, StringBuilder& outBuilder)
 {
-    StringBuilder builder;
-    builder << _getTypePrefix(type->op) << "_" << HLSLIntrinsic::getInfo(op).funcName;
-    return m_slicePool.add(builder);
+    outBuilder << _getTypePrefix(type->op) << "_" << HLSLIntrinsic::getInfo(op).funcName;
+    return SLANG_OK;   
 }
 
 UnownedStringSlice CPPSourceEmitter::_getScalarFuncName(HLSLIntrinsic::Op op, IRBasicType* type)
 {
-    return m_slicePool.getSlice(_calcScalarFuncName(op, type));
+    /* TODO(JS): This is kind of fast and loose. That we don't know all the parameters that are taken or
+    what the return type is, so we can't add to the HLSLIntrinsic map - we just generate the scalar
+    function name and use it (whilst also adding to the slice pool, so that we can return an
+    unowned slice). */
+
+    StringBuilder builder;
+    if (SLANG_FAILED(calcScalarFuncName(op, type, builder)))
+    {
+        SLANG_ASSERT(!"Unable to create scalar function name");
+        return UnownedStringSlice();
+    }
+
+    // Add to the pool. 
+    auto handle = m_slicePool.add(builder);
+    return m_slicePool.getSlice(handle);
 }
 
 UnownedStringSlice CPPSourceEmitter::_getFuncName(const HLSLIntrinsic* specOp)
@@ -1452,14 +1481,22 @@ UnownedStringSlice CPPSourceEmitter::_getFuncName(const HLSLIntrinsic* specOp)
         return m_slicePool.getSlice(handle);
     }
 
-    handle = _calcFuncName(specOp);
+    StringBuilder builder;
+    if (SLANG_FAILED(calcFuncName(specOp, builder)))
+    {
+        SLANG_ASSERT(!"Unable to create function name");
+        // Return an empty slice, as an error...
+        return UnownedStringSlice();
+    }
+
+    handle = m_slicePool.add(builder);
     m_intrinsicNameMap.Add(specOp, handle);
 
     SLANG_ASSERT(handle != StringSlicePool::kNullHandle);
     return m_slicePool.getSlice(handle);
 }
 
-StringSlicePool::Handle CPPSourceEmitter::_calcFuncName(const HLSLIntrinsic* specOp)
+SlangResult CPPSourceEmitter::calcFuncName(const HLSLIntrinsic* specOp, StringBuilder& outBuilder)
 {
     typedef HLSLIntrinsic::Op Op;
 
@@ -1468,7 +1505,7 @@ StringSlicePool::Handle CPPSourceEmitter::_calcFuncName(const HLSLIntrinsic* spe
         IRType* paramType = specOp->signatureType->getParamType(0);
         IRBasicType* basicType = as<IRBasicType>(paramType);
         SLANG_ASSERT(basicType);
-        return _calcScalarFuncName(specOp->op, basicType);
+        return calcScalarFuncName(specOp->op, basicType, outBuilder);
     }
     else
     {
@@ -1483,14 +1520,10 @@ StringSlicePool::Handle CPPSourceEmitter::_calcFuncName(const HLSLIntrinsic* spe
                 IRType* dstType = signatureType->getParamType(0);
                 //IRType* srcType = signatureType->getParamType(1);
 
-                StringBuilder builder;
-                builder << "convert_";
+                outBuilder << "convert_";
                 // I need a function that is called that will construct this
-                if (SLANG_FAILED(calcTypeName(dstType, CodeGenTarget::CSource, builder)))
-                {
-                    return StringSlicePool::kNullHandle;
-                }
-                return m_slicePool.add(builder);
+                SLANG_RETURN_ON_FAIL(calcTypeName(dstType, CodeGenTarget::CSource, outBuilder));
+                return SLANG_OK;
             }
             case Op::ConstructFromScalar:
             {
@@ -1500,22 +1533,20 @@ StringSlicePool::Handle CPPSourceEmitter::_calcFuncName(const HLSLIntrinsic* spe
 
                 IRType* dstType = signatureType->getParamType(0);
                 
-                StringBuilder builder;
-                builder << "constructFromScalar_";
+                outBuilder << "constructFromScalar_";
                 // I need a function that is called that will construct this
-                if (SLANG_FAILED(calcTypeName(dstType, CodeGenTarget::CSource, builder)))
-                {
-                    return StringSlicePool::kNullHandle;
-                }
-                return m_slicePool.add(builder);
+                SLANG_RETURN_ON_FAIL(calcTypeName(dstType, CodeGenTarget::CSource, outBuilder));
+                return SLANG_OK;
             }
             case Op::GetAt:
             {
-                return m_slicePool.add(UnownedStringSlice::fromLiteral("getAt"));
+                outBuilder << "getAt";
+                return SLANG_OK;
             }
             case Op::SetAt:
             {
-                return m_slicePool.add(UnownedStringSlice::fromLiteral("setAt"));
+                outBuilder << "setAt";
+                return SLANG_OK;
             }
             default: break;
         }
@@ -1525,10 +1556,15 @@ StringSlicePool::Handle CPPSourceEmitter::_calcFuncName(const HLSLIntrinsic* spe
         {
             if (!_isOperator(info.funcName))
             {
-                return m_slicePool.add(info.funcName);
+                // If there is a standard default name, just use that
+                outBuilder << info.funcName;
+                return SLANG_OK;
             }
         }
-        return m_slicePool.add(info.name);
+
+        // Just use the name of the Op. This is probably wrong, but gives a pretty good idea of what the desired (presumably missing) op is.
+        outBuilder << info.name;
+        return SLANG_OK;
     }
 }
 
@@ -1993,7 +2029,7 @@ void CPPSourceEmitter::emitPreprocessorDirectivesImpl()
     // Emit all the intrinsics that were used
     for (const auto& keyValue : m_intrinsicNameMap)
     {
-        emitSpecializedOperationDefinition(keyValue.Key);
+        _maybeEmitSpecializedOperationDefinition(keyValue.Key);
     }
 }
 
diff --git a/source/slang/slang-emit-cpp.h b/source/slang/slang-emit-cpp.h
index df1dec380..12bc0939e 100644
--- a/source/slang/slang-emit-cpp.h
+++ b/source/slang/slang-emit-cpp.h
@@ -80,8 +80,12 @@ protected:
 
     // Replaceable for classes derived from CPPSourceEmitter
     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 emitIntrinsicCallExpr(
         IRCall*                         inst,
@@ -115,11 +119,9 @@ protected:
     static TypeDimension _getTypeDimension(IRType* type, bool vecSwap);
     static void _emitAccess(const UnownedStringSlice& name, const TypeDimension& dimension, int row, int col, SourceWriter* writer);
 
-    StringSlicePool::Handle _calcScalarFuncName(HLSLIntrinsic::Op, IRBasicType* type);
     UnownedStringSlice _getScalarFuncName(HLSLIntrinsic::Op operation, IRBasicType* scalarType);
 
     UnownedStringSlice _getFuncName(const HLSLIntrinsic* specOp);
-    StringSlicePool::Handle _calcFuncName(const HLSLIntrinsic* specOp);
 
     UnownedStringSlice _getTypeName(IRType* type);
     
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index 39a25aafa..c72b9125a 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -10,6 +10,23 @@
 
 namespace Slang {
 
+static bool _isSingleNameBasicType(IROp op)
+{
+    switch (op)
+    {
+        case kIROp_Int64Type:   
+        case kIROp_UInt8Type: 
+        case kIROp_UInt16Type:
+        case kIROp_UIntType: 
+        case kIROp_UInt64Type:
+        {
+            return false;
+        }
+        default: return true;
+
+    }
+}
+
 /* static */ UnownedStringSlice CUDASourceEmitter::getBuiltinTypeName(IROp op)
 {
     switch (op)
@@ -110,10 +127,93 @@ SlangResult CUDASourceEmitter::_calcCUDATextureTypeName(IRTextureTypeBase* texTy
     return SLANG_OK;
 }
 
-void CUDASourceEmitter::emitSpecializedOperationDefinition(const HLSLIntrinsic* specOp)
+
+SlangResult CUDASourceEmitter::calcScalarFuncName(HLSLIntrinsic::Op op, IRBasicType* type, StringBuilder& outBuilder)
 {
-    m_writer->emit("__device__ ");
-    Super::emitSpecializedOperationDefinition(specOp);
+    typedef HLSLIntrinsic::Op Op;
+
+    UnownedStringSlice funcName;
+    
+    switch (op)
+    {
+        case Op::Sin:
+        case Op::Cos:
+        case Op::Tan:
+        case Op::ArcSin:
+        case Op::ArcCos:
+        case Op::ArcTan:
+        case Op::ArcTan2:
+        case Op::Floor:
+        case Op::Ceil:
+        case Op::FMod:
+        case Op::Exp2:
+        case Op::Exp:
+        case Op::Log:
+        case Op::Log2:
+        case Op::Log10:
+        case Op::FRem:
+        case Op::Sqrt:
+        case Op::RecipSqrt:
+        case Op::Pow:
+        case Op::Trunc:
+        {
+            if (type->op == kIROp_FloatType || type->op == kIROp_DoubleType)
+            {
+                funcName = HLSLIntrinsic::getInfo(op).funcName;
+            }
+            break;
+        }
+        case Op::Max:
+        case Op::Min:
+        case Op::Abs:
+        {
+            // There are only floating point built in versions of these, prefixed with f
+            if (type->op == kIROp_FloatType || type->op == kIROp_DoubleType)
+            {
+                outBuilder << "f";
+                outBuilder << HLSLIntrinsic::getInfo(op).funcName;
+
+                if (type->op == kIROp_FloatType)
+                {
+                    outBuilder << "f";
+                }
+                return SLANG_OK;
+            }
+            break;
+        }
+
+        default: break;
+    }
+
+    if (funcName.size())
+    {
+        outBuilder << funcName;
+        if (type->op == kIROp_FloatType)
+        {
+            outBuilder << "f";
+        }
+        return SLANG_OK;
+    }
+
+    // Missing ones:
+    // 
+    // sincos - the built in uses pointer, so we'll just define in prelude
+    // rcp
+    // sign
+    // saturate
+    // frac
+    // smoothstep
+    // lerp
+    // clamp
+    // step
+    // 
+    // For integer types
+    // abs
+    // min
+    // max
+
+    // Defer to the supers impl
+    return Super::calcScalarFuncName(op, type, outBuilder);
 }
 
 SlangResult CUDASourceEmitter::calcTypeName(IRType* type, CodeGenTarget target, StringBuilder& out)
@@ -278,73 +378,73 @@ void CUDASourceEmitter::emitOperandImpl(IRInst* inst, EmitOpInfo const& outerPre
     Super::emitOperandImpl(inst, outerPrec);
 }
 
-bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec)
+void CUDASourceEmitter::emitCall(const HLSLIntrinsic* specOp, IRInst* inst, const IRUse* operands, int numOperands, const EmitOpInfo& inOuterPrec)
 {
-    switch (inst->op)
+    switch (specOp->op)
     {
-        case kIROp_Construct:
-        case kIROp_makeVector:
+        case HLSLIntrinsic::Op::Init:
         {
-            if (inst->getOperandCount() == 1)
-            {
-                EmitOpInfo outerPrec = inOuterPrec;
-                bool needClose = false;
+            // For CUDA vector types we construct with make_
 
-                auto prec = getInfo(EmitOp::Prefix);
-                needClose = maybeEmitParens(outerPrec, prec);
+            auto writer = m_writer;
 
-                // Need to emit as cast for HLSL
-                m_writer->emit("(");
-                emitType(inst->getDataType());
-                m_writer->emit(") ");
-                emitOperand(inst->getOperand(0), rightSide(outerPrec, prec));
+            IRType* retType = specOp->returnType;
 
-                maybeCloseParens(needClose);
-                // Handled
-                return true;
-            }
-            else
+            switch (retType->op)
             {
-                m_writer->emit("make_");
-                m_writer->emit(_getTypeName(inst->getDataType()));
-                emitArgs(inst);
-                return true;
+                case kIROp_VectorType:
+                {
+                    // Get the type name
+                    writer->emit("make_");
+                    emitType(retType);
+                    writer->emitChar('(');
+
+                    for (int i = 0; i < numOperands; ++i)
+                    {
+                        if (i > 0)
+                        {
+                            writer->emit(", ");
+                        }
+                        emitOperand(operands[i].get(), getInfo(EmitOp::General));
+                    }
+
+                    writer->emitChar(')');
+                    return;
+                }
+                default: break; 
             }
             break;
         }
-        case kIROp_MakeMatrix:
-        {
-            return false;
-        }
-        case kIROp_BitCast:
+        default: break;
+    }
+
+    return Super::emitCall(specOp, inst, operands, numOperands, inOuterPrec);
+}
+
+bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec)
+{
+    switch(inst->op)
+    {
+        case kIROp_Construct:
         {
-            auto toType = extractBaseType(inst->getDataType());
-            switch (toType)
+            // Simple constructor call
+            // On CUDA some of the built in types can't be used as constructors directly
+
+            IRType* type = inst->getDataType();
+            if (auto basicType = as<IRBasicType>(type) && !_isSingleNameBasicType(type->op))
             {
-                default:
-                    m_writer->emit("/* unhandled */");
-                    break;
-                case BaseType::UInt:
-                    break;
-                case BaseType::Int:
-                    m_writer->emit("(");
-                    emitType(inst->getDataType());
-                    m_writer->emit(")");
-                    break;
-                case BaseType::Float:
-                    m_writer->emit("asfloat");
-                    break;
+                m_writer->emit("(");
+                emitType(inst->getDataType());
+                m_writer->emit(")");
+                emitArgs(inst);
+                return true;
             }
-
-            m_writer->emit("(");
-            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-            m_writer->emit(")");
-            return true;
+            break;
         }
         default: break;
     }
-    // Not handled
-    return false;
+
+    return Super::tryEmitInstExprImpl(inst, inOuterPrec);
 }
 
 void CUDASourceEmitter::emitLayoutDirectivesImpl(TargetRequest* targetReq)
@@ -398,7 +498,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 separtely to call the __device_ implementation
+    // 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__ ");    
 
     CLikeSourceEmitter::emitSimpleFuncImpl(func);
@@ -444,7 +544,7 @@ void CUDASourceEmitter::emitPreprocessorDirectivesImpl()
     // Emit all the intrinsics that were used
     for (const auto& keyValue : m_intrinsicNameMap)
     {
-        emitSpecializedOperationDefinition(keyValue.Key);
+        _maybeEmitSpecializedOperationDefinition(keyValue.Key);
     }
 }
 
diff --git a/source/slang/slang-emit-cuda.h b/source/slang/slang-emit-cuda.h
index c3c88e156..e75eb4e88 100644
--- a/source/slang/slang-emit-cuda.h
+++ b/source/slang/slang-emit-cuda.h
@@ -46,6 +46,7 @@ protected:
     virtual void emitVarDecorationsImpl(IRInst* varDecl) SLANG_OVERRIDE;
     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 bool tryEmitGlobalParamImpl(IRGlobalParam* varDecl, IRType* varType) SLANG_OVERRIDE;
     virtual bool tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOuterPrec) SLANG_OVERRIDE;
@@ -56,7 +57,8 @@ protected:
 
     // CPPSourceEmitter overrides 
     virtual SlangResult calcTypeName(IRType* type, CodeGenTarget target, StringBuilder& out) SLANG_OVERRIDE;
-    virtual void emitSpecializedOperationDefinition(const HLSLIntrinsic* specOp) 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);
 };
diff --git a/source/slang/slang-hlsl-intrinsic-set.h b/source/slang/slang-hlsl-intrinsic-set.h
index 5e01c0599..ee17dd571 100644
--- a/source/slang/slang-hlsl-intrinsic-set.h
+++ b/source/slang/slang-hlsl-intrinsic-set.h
@@ -96,6 +96,10 @@ just constructXXXFromScalar. Would be good if there was a suitable name to encom
         x(Exp2, "exp2", 1) \
         x(Exp, "exp", 1) \
         \
+        x(Log, "log", 1) \
+        x(Log2, "log2", 1) \
+        x(Log10, "log10", 1) \
+        \
         x(Abs, "abs", 1) \
         \
         x(Min, "min", 2) \