WIP: Emit header files for cpp, cuda and metal targets

docs/cpu-target.md

@@ -47,13 +47,15 @@ In the API the `SlangCompileTarget`s are
 ```
 SLANG_C_SOURCE             ///< The C language
 SLANG_CPP_SOURCE           ///< The C++ language
+SLANG_CPP_HEADER           ///< The C++ language (header)
 SLANG_HOST_CPP_SOURCE,     ///< C++ code for `host` style 
 ```        
 
 Using the `-target` command line option
 
 * C_SOURCE: c
 * CPP_SOURCE: cpp,c++,cxx
+* CPP_HEADER: hpp,h++,hxx
 * HOST_CPP_SOURCE: host-cpp,host-c++,host-cxx
 
 Note! Output of C source is not currently supported.

include/slang.h

@@ -589,17 +589,20 @@ extern "C"
         SLANG_DXIL_ASM,
         SLANG_C_SOURCE,                 ///< The C language
         SLANG_CPP_SOURCE,               ///< C++ code for shader kernels.
+        SLANG_CPP_HEADER,               ///< C++ header for shader kernels.
         SLANG_HOST_EXECUTABLE,          ///< Standalone binary executable (for hosting CPU/OS)
         SLANG_SHADER_SHARED_LIBRARY,    ///< A shared library/Dll for shader kernels (for hosting CPU/OS)
         SLANG_SHADER_HOST_CALLABLE,     ///< A CPU target that makes the compiled shader code available to be run immediately
         SLANG_CUDA_SOURCE,              ///< Cuda source
+        SLANG_CUDA_HEADER,              ///< Cuda header
         SLANG_PTX,                      ///< PTX
         SLANG_CUDA_OBJECT_CODE,         ///< Object code that contains CUDA functions.
         SLANG_OBJECT_CODE,              ///< Object code that can be used for later linking
         SLANG_HOST_CPP_SOURCE,          ///< C++ code for host library or executable.
         SLANG_HOST_HOST_CALLABLE,       ///< Host callable host code (ie non kernel/shader) 
         SLANG_CPP_PYTORCH_BINDING,      ///< C++ PyTorch binding code.
         SLANG_METAL,                    ///< Metal shading language
+        SLANG_METAL_HEADER,             ///< Metal shading language header
         SLANG_METAL_LIB,                ///< Metal library
         SLANG_METAL_LIB_ASM,            ///< Metal library assembly
         SLANG_HOST_SHARED_LIBRARY,      ///< A shared library/Dll for host code (for hosting CPU/OS)

prelude/slang-cuda-prelude.h

@@ -1,3 +1,6 @@
+#ifndef SLANG_CUDA_PRELUDE_H
+#define SLANG_CUDA_PRELUDE_H
+
 #define SLANG_PRELUDE_EXPORT
 
 #ifdef __CUDACC_RTC__
@@ -2109,12 +2112,12 @@ __inline__ __device__ uint4 _waveMatchMultiple(WarpMask mask, const T& inVal)
     return make_uint4(matchBits, 0, 0, 0);
 }
 
-__device__ uint getAt(dim3 a,  int b)
+__inline__ __device__ uint getAt(dim3 a,  int b)
 {
     SLANG_PRELUDE_ASSERT(b >= 0 && b < 3);
     return (&a.x)[b];
 }
-__device__ uint3 operator*(uint3 a, dim3 b)
+__inline__ __device__ uint3 operator*(uint3 a, dim3 b)
 {
     uint3 r;
     r.x = a.x * b.x;
@@ -2364,3 +2367,5 @@ struct TensorView
         *reinterpret_cast<T*>(data + offset) = val;
     }
 };
+
+#endif

source/compiler-core/slang-artifact-desc-util.cpp

@@ -271,13 +271,15 @@ SLANG_HIERARCHICAL_ENUM(ArtifactStyle, SLANG_ARTIFACT_STYLE, SLANG_ARTIFACT_STYL
         case SLANG_DXIL_ASM:                return Desc::make(Kind::Assembly, Payload::DXIL, Style::Kernel, 0);
         case SLANG_C_SOURCE:                return Desc::make(Kind::Source, Payload::C, Style::Kernel, 0);
         case SLANG_CPP_SOURCE:              return Desc::make(Kind::Source, Payload::Cpp, Style::Kernel, 0);
+        case SLANG_CPP_HEADER:              return Desc::make(Kind::Source, Payload::Cpp, Style::Kernel, 0);
         case SLANG_HOST_CPP_SOURCE:         return Desc::make(Kind::Source, Payload::Cpp, Style::Host, 0);
         case SLANG_CPP_PYTORCH_BINDING:     return Desc::make(Kind::Source, Payload::Cpp, Style::Host, 0);
         case SLANG_HOST_EXECUTABLE:         return Desc::make(Kind::Executable, Payload::HostCPU, Style::Host, 0);
         case SLANG_HOST_SHARED_LIBRARY:     return Desc::make(Kind::SharedLibrary, Payload::HostCPU, Style::Host, 0);
         case SLANG_SHADER_SHARED_LIBRARY:   return Desc::make(Kind::SharedLibrary, Payload::HostCPU, Style::Kernel, 0);
         case SLANG_SHADER_HOST_CALLABLE:    return Desc::make(Kind::HostCallable, Payload::HostCPU, Style::Kernel, 0);
         case SLANG_CUDA_SOURCE:             return Desc::make(Kind::Source, Payload::CUDA, Style::Kernel, 0);
+        case SLANG_CUDA_HEADER:             return Desc::make(Kind::Source, Payload::CUDA, Style::Kernel, 0);
             // TODO(JS):
             // Not entirely clear how best to represent PTX here. We could mark as 'Assembly'. Saying it is 
             // 'Executable' implies it is Binary (which PTX isn't). Executable also implies 'complete for executation', 
@@ -286,6 +288,7 @@ SLANG_HIERARCHICAL_ENUM(ArtifactStyle, SLANG_ARTIFACT_STYLE, SLANG_ARTIFACT_STYL
         case SLANG_OBJECT_CODE:             return Desc::make(Kind::ObjectCode, Payload::HostCPU, Style::Kernel, 0);
         case SLANG_HOST_HOST_CALLABLE:      return Desc::make(Kind::HostCallable, Payload::HostCPU, Style::Host, 0);
         case SLANG_METAL:                   return Desc::make(Kind::Source, Payload::Metal, Style::Kernel, 0);
+        case SLANG_METAL_HEADER:            return Desc::make(Kind::Source, Payload::Metal, Style::Kernel, 0);
         case SLANG_METAL_LIB:               return Desc::make(Kind::Executable, Payload::MetalAIR, Style::Kernel, 0);
         case SLANG_METAL_LIB_ASM:           return Desc::make(Kind::Assembly, Payload::MetalAIR, Style::Kernel, 0);
         default: break;

source/core/slang-type-text-util.cpp

@@ -49,18 +49,21 @@ static const TypeTextUtil::CompileTargetInfo s_compileTargetInfos[] =
     { SLANG_SPIRV_ASM,      "spv-asm",                                          "spirv-asm,spirv-assembly", "SPIR-V assembly" },
     { SLANG_C_SOURCE,       "c",                                                "c",                        "C source code" },
     { SLANG_CPP_SOURCE,     "cpp,c++,cxx",                                      "cpp,c++,cxx",              "C++ source code" },
+    { SLANG_CPP_HEADER,     "hpp,h++,hxx",                                      "hpp,h++,hxx",              "C++ source header" },
     { SLANG_CPP_PYTORCH_BINDING, "cpp,c++,cxx",                                 "torch,torch-binding,torch-cpp,torch-cpp-binding", "C++ for pytorch binding" },
     { SLANG_HOST_CPP_SOURCE, "cpp,c++,cxx",                                     "host-cpp,host-c++,host-cxx", "C++ source for host execution"},
     { SLANG_HOST_EXECUTABLE,"exe",                                              "exe,executable",           "Executable binary" },
     { SLANG_SHADER_SHARED_LIBRARY, "shader-dll,shader-so",                      "shader-sharedlib,shader-sharedlibrary,shader-dll", "Shared library/Dll for shader kernel" },
     { SLANG_HOST_SHARED_LIBRARY, "dll,so",                                      "sharedlib,sharedlibrary,dll", "Shared library/Dll for host execution" },
     { SLANG_CUDA_SOURCE,    "cu",                                               "cuda,cu",                  "CUDA source code"  },
+    { SLANG_CUDA_HEADER,    "cuh",                                              "cuh",                      "CUDA source header"  },
     { SLANG_PTX,            "ptx",                                              "ptx",                      "PTX assembly" },
     { SLANG_CUDA_OBJECT_CODE, "obj,o",                                          "cuobj,cubin",              "CUDA binary" },
     { SLANG_SHADER_HOST_CALLABLE,  "",                                          "host-callable,callable",   "Host callable" },
     { SLANG_OBJECT_CODE,    "obj,o",                                            "object-code",              "Object code" },
     { SLANG_HOST_HOST_CALLABLE, "",                                             "host-host-callable",       "Host callable for host execution" },
     { SLANG_METAL,          "metal",                                            "metal",                    "Metal shader source" },
+    { SLANG_METAL_HEADER,   "metal.hpp",                                        "metal.hpp",                "Metal shader header" },
     { SLANG_METAL_LIB,      "metallib",                                         "metallib",                 "Metal Library Bytecode" },
     { SLANG_METAL_LIB_ASM,  "metallib-asm"                                      "metallib-asm",             "Metal Library Bytecode assembly" },
 };

source/slang/slang-compiler.cpp

@@ -538,11 +538,14 @@ namespace Slang
             case CodeGenTarget::GLSL:
             case CodeGenTarget::HLSL:
             case CodeGenTarget::CUDASource:
+            case CodeGenTarget::CUDAHeader:
             case CodeGenTarget::CPPSource:
+            case CodeGenTarget::CPPHeader:
             case CodeGenTarget::HostCPPSource:
             case CodeGenTarget::PyTorchCppBinding:
             case CodeGenTarget::CSource:
             case CodeGenTarget::Metal:
+            case CodeGenTarget::MetalHeader:
             {
                 return PassThroughMode::None;
             }
@@ -973,6 +976,7 @@ namespace Slang
         {
             case CodeGenTarget::PTX:
             case CodeGenTarget::CUDASource:
+            case CodeGenTarget::CUDAHeader:
             {
                 return new CUDAExtensionTracker;
             }
@@ -1716,11 +1720,14 @@ namespace Slang
         case CodeGenTarget::GLSL:
         case CodeGenTarget::HLSL:
         case CodeGenTarget::CUDASource:
+        case CodeGenTarget::CUDAHeader:
         case CodeGenTarget::CPPSource:
+        case CodeGenTarget::CPPHeader:
         case CodeGenTarget::HostCPPSource:
         case CodeGenTarget::PyTorchCppBinding:
         case CodeGenTarget::CSource:
         case CodeGenTarget::Metal:
+        case CodeGenTarget::MetalHeader:
             {
                 RefPtr<ExtensionTracker> extensionTracker = _newExtensionTracker(target);
 

source/slang/slang-compiler.h

@@ -80,18 +80,21 @@ namespace Slang
         DXILAssembly        = SLANG_DXIL_ASM,
         CSource             = SLANG_C_SOURCE,
         CPPSource           = SLANG_CPP_SOURCE,
+        CPPHeader           = SLANG_CPP_HEADER,
         PyTorchCppBinding   = SLANG_CPP_PYTORCH_BINDING,
         HostCPPSource       = SLANG_HOST_CPP_SOURCE,
         HostExecutable      = SLANG_HOST_EXECUTABLE,
         HostSharedLibrary   = SLANG_HOST_SHARED_LIBRARY,
         ShaderSharedLibrary = SLANG_SHADER_SHARED_LIBRARY,
         ShaderHostCallable  = SLANG_SHADER_HOST_CALLABLE,
         CUDASource          = SLANG_CUDA_SOURCE,
+        CUDAHeader          = SLANG_CUDA_HEADER,
         PTX                 = SLANG_PTX,
         CUDAObjectCode      = SLANG_CUDA_OBJECT_CODE,
         ObjectCode          = SLANG_OBJECT_CODE,
         HostHostCallable    = SLANG_HOST_HOST_CALLABLE,
         Metal               = SLANG_METAL,
+        MetalHeader         = SLANG_METAL_HEADER,
         MetalLib            = SLANG_METAL_LIB,
         MetalLibAssembly    = SLANG_METAL_LIB_ASM,
         CountOf             = SLANG_TARGET_COUNT_OF,

source/slang/slang-emit-c-like.cpp

@@ -82,16 +82,19 @@ struct CLikeSourceEmitter::ComputeEmitActionsContext
             return SourceLanguage::C;
         }
         case CodeGenTarget::CPPSource:
+        case CodeGenTarget::CPPHeader:
         case CodeGenTarget::HostCPPSource:
         case CodeGenTarget::PyTorchCppBinding:
         {
             return SourceLanguage::CPP;
         }
         case CodeGenTarget::CUDASource:
+        case CodeGenTarget::CUDAHeader:
         {
             return SourceLanguage::CUDA;
         }
         case CodeGenTarget::Metal:
+        case CodeGenTarget::MetalHeader:
         {
             return SourceLanguage::Metal;
         }
@@ -1150,7 +1153,20 @@ String CLikeSourceEmitter::getName(IRInst* inst)
     String name;
     if(!m_mapInstToName.tryGetValue(inst, name))
     {
-        name = generateName(inst);
+        // unmangle names, when emitting header
+        if (shouldEmitOnlyHeader())
+        {
+            if (auto nameHintDecor = inst->findDecoration<IRNameHintDecoration>())
+            {
+                StringBuilder sb;
+                appendScrubbedName(nameHintDecor->getName(), sb);
+                name = sb.produceString();
+            }
+        }
+        else
+        {
+            name = generateName(inst);
+        }
         m_mapInstToName.add(inst, name);
     }
     return name;
@@ -3476,7 +3492,7 @@ void CLikeSourceEmitter::emitSimpleFuncImpl(IRFunc* func)
     emitSemantics(func);
 
     // TODO: encode declaration vs. definition
-    if(isDefinition(func))
+    if(!shouldEmitOnlyHeader() && isDefinition(func))
     {
         m_writer->emit("\n{\n");
         m_writer->indent();
@@ -3622,6 +3638,11 @@ bool shouldWrapInExternCBlock(IRFunc* func)
 
 void CLikeSourceEmitter::emitFunc(IRFunc* func)
 {
+    if (shouldEmitOnlyHeader() && !func->findDecoration<IRExternCppDecoration>())
+    {
+        return;
+    }
+
     // Target-intrinsic functions should never be emitted
     // even if they happen to have a body.
     //
@@ -3669,6 +3690,11 @@ void CLikeSourceEmitter::emitFuncDecorationsImpl(IRFunc* func)
 
 void CLikeSourceEmitter::emitStruct(IRStructType* structType)
 {
+    if (shouldEmitOnlyHeader() && !structType->findDecoration<IRExternCppDecoration>())
+    {
+        return;
+    }
+
     ensureTypePrelude(structType);
 
     // If the selected `struct` type is actually an intrinsic
@@ -4064,11 +4090,20 @@ void CLikeSourceEmitter::emitGlobalVar(IRGlobalVar* varDecl)
 
             m_writer->emit("\n");
             emitType(varType, initFuncName);
-            m_writer->emit("()\n{\n");
-            m_writer->indent();
-            emitFunctionBody(varDecl);
-            m_writer->dedent();
-            m_writer->emit("}\n");
+
+            // When emiting header, emit only declaration.
+            if (shouldEmitOnlyHeader())
+            {
+                m_writer->emit(";\n");
+            }
+            else
+            {
+                m_writer->emit("()\n{\n");
+                m_writer->indent();
+                emitFunctionBody(varDecl);
+                m_writer->dedent();
+                m_writer->emit("}\n");
+            }
         }
     }
 

source/slang/slang-emit-c-like.h

@@ -469,6 +469,8 @@ class CLikeSourceEmitter: public SourceEmitterBase
     protected:
 
 
+
+    virtual bool shouldEmitOnlyHeader() { return false; }
     virtual void emitPostDeclarationAttributesForType(IRInst* type) { SLANG_UNUSED(type); }
     virtual bool doesTargetSupportPtrTypes() { return false; }
     virtual void emitLayoutSemanticsImpl(IRInst* inst, char const* uniformSemanticSpelling, EmitLayoutSemanticOption layoutSemanticOption) { SLANG_UNUSED(inst); SLANG_UNUSED(uniformSemanticSpelling); SLANG_UNUSED(layoutSemanticOption); }

source/slang/slang-emit-cpp.cpp

@@ -950,7 +950,7 @@ void CPPSourceEmitter::emitSimpleFuncImpl(IRFunc* func)
     emitSemantics(func);
 
     // TODO: encode declaration vs. definition
-    if (isDefinition(func))
+    if (!shouldEmitOnlyHeader() && isDefinition(func))
     {
         m_writer->emit("\n{\n");
         m_writer->indent();
@@ -1713,7 +1713,7 @@ bool CPPSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOut
 
 void CPPSourceEmitter::emitPreModuleImpl()
 {
-    if (m_target == CodeGenTarget::CPPSource)
+    if (m_target == CodeGenTarget::CPPSource || m_target == CodeGenTarget::CPPHeader)
     {
         // TODO(JS): Previously this opened an anonymous scope for all generated functions
         // Unfortunately this is a problem if we are just emitting code that is externally available
@@ -1737,6 +1737,13 @@ void CPPSourceEmitter::emitPreModuleImpl()
     Super::emitPreModuleImpl();
 }
 
+void CPPSourceEmitter::emitFrontMatterImpl(TargetRequest*)
+{
+    if (shouldEmitOnlyHeader())
+    {
+        m_writer->emit("#pragma once\n\n");
+    }
+}
 
 void CPPSourceEmitter::emitGlobalInstImpl(IRInst* inst)
 {
@@ -2145,7 +2152,7 @@ void CPPSourceEmitter::emitModuleImpl(IRModule* module, DiagnosticSink* sink)
     // Now that we can have any function available externally (not just entry points)
     // this doesn't work. 
 
-    //if (m_target == CodeGenTarget::CPPSource)
+    //if (m_target == CodeGenTarget::CPPSource || m_target == CodeGenTarget::CPPHeader)
     //{
         // Need to close the anonymous namespace when outputting for C++ kernel.
         //m_writer->emit("} // anonymous\n\n");

source/slang/slang-emit-cpp.h

@@ -46,9 +46,11 @@ class CPPSourceEmitter: public CLikeSourceEmitter
 protected:
 
     // Implement CLikeSourceEmitter interface
+    virtual bool shouldEmitOnlyHeader() SLANG_OVERRIDE { return m_target == CodeGenTarget::CPPHeader; }
     virtual bool doesTargetSupportPtrTypes() SLANG_OVERRIDE { return true; }
     virtual void emitParameterGroupImpl(IRGlobalParam* varDecl, IRUniformParameterGroupType* type) SLANG_OVERRIDE;
     virtual void emitEntryPointAttributesImpl(IRFunc* irFunc, IREntryPointDecoration* entryPointDecor) SLANG_OVERRIDE;
+    virtual void emitFrontMatterImpl(TargetRequest* targetReq) SLANG_OVERRIDE;
     virtual void emitSimpleTypeImpl(IRType* type) SLANG_OVERRIDE;
     virtual void _emitType(IRType* type, DeclaratorInfo* declarator) SLANG_OVERRIDE;
     virtual void emitVectorTypeNameImpl(IRType* elementType, IRIntegerValue elementCount) SLANG_OVERRIDE;

source/slang/slang-emit-cuda.cpp

@@ -142,7 +142,7 @@ SlangResult CUDASourceEmitter::calcTypeName(IRType* type, CodeGenTarget target,
     SLANG_UNUSED(target);
 
     // The names CUDA produces are all compatible with 'C' (ie they aren't templated types)
-    SLANG_ASSERT(target == CodeGenTarget::CUDASource || target == CodeGenTarget::CSource);
+    SLANG_ASSERT(target == CodeGenTarget::CUDASource || target == CodeGenTarget::CUDAHeader || target == CodeGenTarget::CSource);
 
     switch (type->getOp())
     {
@@ -794,6 +794,14 @@ void CUDASourceEmitter::emitSimpleValueImpl(IRInst* inst)
     Super::emitSimpleValueImpl(inst);
 }
 
+void CUDASourceEmitter::emitFrontMatterImpl(TargetRequest*)
+{
+    if (shouldEmitOnlyHeader())
+    {
+        m_writer->emit("#pragma once\n\n");
+    }
+}
+
 
 void CUDASourceEmitter::emitSemanticsImpl(IRInst* inst, bool allowOffsetLayout)
 {

source/slang/slang-emit-cuda.h

@@ -63,9 +63,11 @@ class CUDASourceEmitter : public CPPSourceEmitter
 
 protected:
 
+    virtual bool shouldEmitOnlyHeader() SLANG_OVERRIDE { return m_target == CodeGenTarget::CUDAHeader; }
     virtual void emitLayoutSemanticsImpl(IRInst* inst, char const* uniformSemanticSpelling, EmitLayoutSemanticOption layoutSemanticOption) SLANG_OVERRIDE;
     virtual void emitParameterGroupImpl(IRGlobalParam* varDecl, IRUniformParameterGroupType* type) SLANG_OVERRIDE;
     virtual void emitEntryPointAttributesImpl(IRFunc* irFunc, IREntryPointDecoration* entryPointDecor) SLANG_OVERRIDE;
+    virtual void emitFrontMatterImpl(TargetRequest* targetReq) SLANG_OVERRIDE;
 
     virtual void emitRateQualifiersAndAddressSpaceImpl(IRRate* rate, AddressSpace addressSpace) SLANG_OVERRIDE;
     virtual void emitSemanticsImpl(IRInst* inst, bool allowOffsetLayout) SLANG_OVERRIDE;