3 files changed, 9 insertions, 1 deletions
diff --git a/source/core/slang-render-api-util.cpp b/source/core/slang-render-api-util.cpp
index a9339c14e..960537a0b 100644
--- a/source/core/slang-render-api-util.cpp
+++ b/source/core/slang-render-api-util.cpp
@@ -18,6 +18,7 @@ namespace Slang {
     { RenderApiType::D3D12,  "dx12,d3d12",      ""},
     { RenderApiType::D3D11,  "dx11,d3d11",      "hlsl,hlsl-rewrite,slang"},
     { RenderApiType::CPU,    "cpu",             ""},
+    { RenderApiType::CUDA,   "cuda",            "cuda,ptx"},
 };
 
 static int _calcAvailableApis()
@@ -268,6 +269,11 @@ static bool _canLoadSharedLibrary(const char* libName)
         case RenderApiType::D3D11:     return _canLoadSharedLibrary("d3d11"); 
         case RenderApiType::D3D12:     return _canLoadSharedLibrary("d3d12");
         case RenderApiType::CPU:       return true;
+        case RenderApiType::CUDA:
+        {
+            // We'll assume it's available, and if not trying to create it will detect it
+            return true;
+        }
         default: break; 
     }
 #elif SLANG_UNIX_FAMILY
diff --git a/source/core/slang-render-api-util.h b/source/core/slang-render-api-util.h
index 48b599653..b028d3996 100644
--- a/source/core/slang-render-api-util.h
+++ b/source/core/slang-render-api-util.h
@@ -16,6 +16,7 @@ enum class RenderApiType
     D3D12,
     D3D11,
     CPU,
+    CUDA,
     CountOf,
 };
 
@@ -29,6 +30,7 @@ struct RenderApiFlag
         D3D12 = 1 << int(RenderApiType::D3D12),
         D3D11 = 1 << int(RenderApiType::D3D11),
         CPU   = 1 << int(RenderApiType::CPU),
+        CUDA  = 1 << int(RenderApiType::CUDA),
         AllOf = (1 << int(RenderApiType::CountOf)) - 1                   ///< All bits set
     };
 };
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index f2c9a1e80..980e94a29 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -270,7 +270,7 @@ void CUDASourceEmitter::emitEntryPointAttributesImpl(IRFunc* irFunc, IREntryPoin
             m_writer->emit(")]\n");
 #endif
 
-            m_writer->emit("__global__ ");
+            m_writer->emit("extern \"C\" __global__ ");
             break;
         }