Initial work to support OptiX output for ray tracing shaders (#1307)

* Initial work to support OptiX output for ray tracing shaders

This change represents in-progress work toward allowing Slang/HLSL ray-tracing shaders to be cross-compiled for execution on top of OptiX. The work as it exists here is incomplete, but the changes are incremental and should not disturb existing supported use cases.

One major unresolved issue in this work is that the OptiX SDK does not appear to set an environment variable

Changes include:

* Modified the premake script to support new options for adding OptiX to the build. Right now the default path to the OptiX SDK is hard-coded because the installer doesn't seem to set an environment variable. We will want to update that to have a reasonable default path for both Windows and Unix-y platforms in a later chance.

* I ran the premake generator on the project since I added new options, which resulted in a bunch of diffs to the Visual Studio project files that are unrelated to this change. Many of the diffs come from previous edits that added files using only the Visual Studio IDE rather than by re-running premake, so it is arguably better to have the checked-in project files more accurately reflect the generated files used for CI builds.

* The "downstream compiler" abstraction was extended to have an explicit notion of the kind of pipeline that shaders are being compiled for (e.g., compute vs. rasterization vs. ray tracing). This option is used to tell the NVRTC case when it needs to include the OptiX SDK headers in the search path for shader compilation (and also when it should add a `#define` to make the prelude pull in OptiX). This code again uses a hard-coded default path for the OptiX SDK; we will need to modify that to have a better discovery approach and also to support an API or command-line override.

  * One note for the future is that instead of passing down a "pipeline type" we could instead pass down the list/set of stages for the kernels being compiled, and the OptiX support could be enabled whenever there is *any* ray tracing entry point present in a module. That approach would allow mixing RT and compute kernels during downstream compilation. We will need to revisit these choices when we start supporting code generation for multiple entry points at a time.

* The CUDA emit logic is currently mostly unchanged. The biggest difference is that when emitting a ray-tracing entry point we prefix the name of the generated `__global__` function with a marker for its stage type, as required by the OptiX runtime (e.g., a `__raygen__` prefix is required on all ray-generation entry points).

* The `Renderer` abstraction had a bare minimum of changes made to be able to understand that ray-tracing pipelines exist, and also that some APIs will require the name of each entry point along with its binary data in order to create a program.

* The `ShaderCompileRequest` type was updated so that only a single "source" is supported (rather than distinct source for each entry point), and also the entry points have been turned into a single list where each entry identifies its stage instead of a fixed list of fields for the supported entry-point types.

* The CUDA compute path had a lot of code added to support execution for the new ray-tracing pipeline type. The logic is mostly derived from the `optixHello` example in the OptiX SDK, and at present only supports running a single ray-generation shader with no parameters. The code here is not intended to be ready for use, but represents a signficiant amount of learning-by-doing.

* The `slang-support.cpp` file in `render-test` was updated so that instead of having separate compilation logic for compute vs. rasterization shaders (which would mean adding a third path for ray tracing), there is now a single flow to the code that works for all pipeline types and any kind of entry points.

  * Implicit in the new code is dropping support for the way GLSL was being compiled for pass-through render tests, which means pass-through GLSL render tests will no longer work. It seems like we didn't have any of those to begin with, though, so it is no great loss.

  * Also implicit are some new invariants about how shaders without known/default entry points need to be handled. For example, the ray tracing case intentionally does not fill in entry points on the `ShaderCompileRequest` and instead fully relies on the Slang compiler's support for discovering and enumerating entry points via reflection. As a consequence of those edits the `-no-default-entry-point` flag on `render-test` is probably not working, but it seems like we don't have any test cases that use that flag anyway.

Given the seemingly breaking changes in those last two bullets, I was surprised to find that all our current tests seem to pass with this change. If there are things that I'm missing, I hope they will come up in review.

* fixup: issues from review and CI

* Some issues noted during the review process (e.g., a missing `break`)

* Fix logic for render tests with `-no-default-entry-point`. I had somehow missed that we had tests reliant on that flag. This required a bit of refactoring to pass down the relevant flag (luckily the function in question was already being passed most of what was in `Options`, so that just passing that in directly actually simplifies the call sites a bit.

* There was a missing line of code to actually add the default compute entry points to the compile request. I think this was a problem that slipped in as part of some pre-PR refactoring/cleanup changes that I failed to re-test.

19 files changed, 784 insertions, 285 deletions

diff --git a/examples/gpu-printing/gpu-printing.vcxproj b/examples/gpu-printing/gpu-printing.vcxproj
index d97c97914..1cb664892 100644
--- a/examples/gpu-printing/gpu-printing.vcxproj
+++ b/examples/gpu-printing/gpu-printing.vcxproj
@@ -162,6 +162,10 @@
     </Link>
   </ItemDefinitionGroup>
   <ItemGroup>
+    <ClInclude Include="gpu-printing-ops.h" />
+    <ClInclude Include="gpu-printing.h" />
+  </ItemGroup>
+  <ItemGroup>
     <ClCompile Include="gpu-printing.cpp" />
     <ClCompile Include="main.cpp" />
   </ItemGroup>
@@ -180,10 +184,6 @@
       <Project>{222F7498-B40C-4F3F-A704-DDEB91A4484A}</Project>
     </ProjectReference>
   </ItemGroup>
-  <ItemGroup>
-    <ClInclude Include="gpu-printing-ops.h" />
-    <ClInclude Include="gpu-printing.h" />
-  </ItemGroup>
   <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
   <ImportGroup Label="ExtensionTargets">
   </ImportGroup>
diff --git a/examples/gpu-printing/gpu-printing.vcxproj.filters b/examples/gpu-printing/gpu-printing.vcxproj.filters
index c539443c3..7b1743f46 100644
--- a/examples/gpu-printing/gpu-printing.vcxproj.filters
+++ b/examples/gpu-printing/gpu-printing.vcxproj.filters
@@ -1,15 +1,26 @@
-<?xml version="1.0" encoding="utf-8"?>
+<?xml version="1.0" encoding="utf-8"?>
 <Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
   <ItemGroup>
+    <Filter Include="Header Files">
+      <UniqueIdentifier>{21EB8090-0D4E-1035-B6D3-48EBA215DCB7}</UniqueIdentifier>
+    </Filter>
     <Filter Include="Source Files">
       <UniqueIdentifier>{E9C7FDCE-D52A-8D73-7EB0-C5296AF258F6}</UniqueIdentifier>
     </Filter>
   </ItemGroup>
   <ItemGroup>
-    <ClCompile Include="main.cpp">
+    <ClInclude Include="gpu-printing-ops.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="gpu-printing.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+  </ItemGroup>
+  <ItemGroup>
+    <ClCompile Include="gpu-printing.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
-    <ClCompile Include="gpu-printing.cpp">
+    <ClCompile Include="main.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
   </ItemGroup>
@@ -21,12 +32,4 @@
       <Filter>Source Files</Filter>
     </None>
   </ItemGroup>
-  <ItemGroup>
-    <ClInclude Include="gpu-printing-ops.h">
-      <Filter>Source Files</Filter>
-    </ClInclude>
-    <ClInclude Include="gpu-printing.h">
-      <Filter>Source Files</Filter>
-    </ClInclude>
-  </ItemGroup>
 </Project>
\ No newline at end of file
diff --git a/prelude/slang-cuda-prelude.h b/prelude/slang-cuda-prelude.h
index b5d8b3788..c23189320 100644
--- a/prelude/slang-cuda-prelude.h
+++ b/prelude/slang-cuda-prelude.h
@@ -1,3 +1,7 @@
+#ifdef SLANG_CUDA_ENABLE_OPTIX
+#include <optix.h>
+#endif
+
 
 // Must be large enough to cause overflow and therefore infinity
 #ifndef SLANG_INFINITY
diff --git a/premake5.lua b/premake5.lua
index bad578ef5..f0fb7b341 100644
--- a/premake5.lua
+++ b/premake5.lua
@@ -98,6 +98,21 @@ newoption {
 }
 
 newoption {
+   trigger     = "enable-optix",
+   description = "(Optional) If true will enable OptiX build/ tests (also implicitly enables CUDA)",
+   value       = "bool",
+   default     = "false",
+   allowed     = { { "true", "True"}, { "false", "False" } }
+}
+
+newoption {
+   trigger     = "optix-sdk-path",
+   description = "(Optional) Path to the root of OptiX SDK. (Implicitly enabled OptiX and CUDA)",
+   value       = "path"
+}
+
+
+newoption {
    trigger     = "enable-profile",
    description = "(Optional) If true will enable slang-profile tool - suitable for gprof usage on linux",
    value       = "bool",
@@ -111,13 +126,21 @@ targetDetail = _OPTIONS["target-detail"]
 buildGlslang = (_OPTIONS["build-glslang"] == "true")
 enableCuda = not not (_OPTIONS["enable-cuda"] == "true" or _OPTIONS["cuda-sdk-path"])
 enableProfile = (_OPTIONS["enable-profile"] == "true")
+optixPath = _OPTIONS["optix-sdk-path"]
+enableOptix = not not (_OPTIONS["enable-optix"] == "true" or optixPath)
+enableProfile = (_OPTIONS["enable-profile"] == "true")
+
+if enableOptix then
+    optixPath = optixPath or "C:/ProgramData/NVIDIA Corporation/OptiX SDK 7.0.0/"
+    enableCuda = true;
+end
 
 -- cudaPath is only set if cuda is enabled, and CUDA_PATH enviromental variable is set
 cudaPath = nil
 if enableCuda then
     -- Get the CUDA path. Use the value set on cuda-sdk-path by default, if not set use the environment variable. 
     cudaPath = (_OPTIONS["cuda-sdk-path"] or os.getenv("CUDA_PATH"))
-end    
+end
 
 -- Is true when the target is really windows (ie not something on top of windows like cygwin)
 local isTargetWindows = (os.target() == "windows") and not (targetDetail == "mingw" or targetDetail == "cygwin")
@@ -568,6 +591,11 @@ toolSharedLibrary "render-test"
         defines { "RENDER_TEST_CUDA" }
         includedirs { cudaPath .. "/include" }
         includedirs { cudaPath .. "/include", cudaPath .. "/common/inc" }
+
+        if optixPath then
+            defines { "RENDER_TEST_OPTIX" }
+            includedirs { optixPath .. "include/" }
+        end
         
         links { "cuda", "cudart" }   
         
@@ -575,8 +603,7 @@ toolSharedLibrary "render-test"
             libdirs { cudaPath .. "/lib/Win32/" }
            
         filter { "platforms:x64" }
-            libdirs { cudaPath .. "/lib/x64/" }   
-           
+            libdirs { cudaPath .. "/lib/x64/" }       
     end
   
 --
diff --git a/slang.sln b/slang.sln
index 2160d7e3b..cd70d88e1 100644
--- a/slang.sln
+++ b/slang.sln
@@ -1,38 +1,38 @@
 
 Microsoft Visual Studio Solution File, Format Version 12.00
 # Visual Studio 14
-Project("{2150E333-8FDC-42A3-9474-1A3956D46DE8}") = "examples", "examples", "{EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}"
-EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "hello-world", "examples\hello-world\hello-world.vcxproj", "{010BE414-ED5B-CF56-16C0-BD18027062C0}"
-EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "model-viewer", "examples\model-viewer\model-viewer.vcxproj", "{2F8724C6-1BC3-2730-84D5-3F277030D04A}"
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "core", "source\core\core.vcxproj", "{F9BE7957-8399-899E-0C49-E714FDDD4B65}"
 EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "gpu-printing", "examples\gpu-printing\gpu-printing.vcxproj", "{57C81DD3-4304-213D-AC16-39349871C957}"
+Project("{2150E333-8FDC-42A3-9474-1A3956D46DE8}") = "examples", "examples", "{EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}"
 EndProject
 Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "cpu-hello-world", "examples\cpu-hello-world\cpu-hello-world.vcxproj", "{4B47A364-37C4-96A7-6041-97BB4C1D333B}"
 EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "core", "source\core\core.vcxproj", "{F9BE7957-8399-899E-0C49-E714FDDD4B65}"
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "gpu-printing", "examples\gpu-printing\gpu-printing.vcxproj", "{57C81DD3-4304-213D-AC16-39349871C957}"
 EndProject
-Project("{2150E333-8FDC-42A3-9474-1A3956D46DE8}") = "tools", "tools", "{FD47AE19-69FD-260F-F2F1-20E65EA61D13}"
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "hello-world", "examples\hello-world\hello-world.vcxproj", "{010BE414-ED5B-CF56-16C0-BD18027062C0}"
 EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "slang-generate", "tools\slang-generate\slang-generate.vcxproj", "{66174227-8541-41FC-A6DF-4764FC66F78E}"
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "model-viewer", "examples\model-viewer\model-viewer.vcxproj", "{2F8724C6-1BC3-2730-84D5-3F277030D04A}"
 EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "slang-test", "tools\slang-test\slang-test.vcxproj", "{0C768A18-1D25-4000-9F37-DA5FE99E3B64}"
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "slang", "source\slang\slang.vcxproj", "{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}"
+	ProjectSection(ProjectDependencies) = postProject
+		{66174227-8541-41FC-A6DF-4764FC66F78E} = {66174227-8541-41FC-A6DF-4764FC66F78E}
+	EndProjectSection
 EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "gfx", "tools\gfx\gfx.vcxproj", "{222F7498-B40C-4F3F-A704-DDEB91A4484A}"
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "slangc", "source\slangc\slangc.vcxproj", "{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}"
 EndProject
 Project("{2150E333-8FDC-42A3-9474-1A3956D46DE8}") = "test-tool", "test-tool", "{57B5AA5E-C340-1823-CC51-9B17385C7423}"
 EndProject
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "render-test-tool", "tools\render-test\render-test-tool.vcxproj", "{61F7EB00-7281-4BF3-9470-7C2EA92620C3}"
+EndProject
 Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "slang-reflection-test-tool", "tools\slang-reflection-test\slang-reflection-test-tool.vcxproj", "{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}"
 EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "render-test-tool", "tools\render-test\render-test-tool.vcxproj", "{61F7EB00-7281-4BF3-9470-7C2EA92620C3}"
+Project("{2150E333-8FDC-42A3-9474-1A3956D46DE8}") = "tools", "tools", "{FD47AE19-69FD-260F-F2F1-20E65EA61D13}"
 EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "slangc", "source\slangc\slangc.vcxproj", "{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}"
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "gfx", "tools\gfx\gfx.vcxproj", "{222F7498-B40C-4F3F-A704-DDEB91A4484A}"
 EndProject
-Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "slang", "source\slang\slang.vcxproj", "{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}"
-	ProjectSection(ProjectDependencies) = postProject
-		{66174227-8541-41FC-A6DF-4764FC66F78E} = {66174227-8541-41FC-A6DF-4764FC66F78E}
-	EndProjectSection
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "slang-generate", "tools\slang-generate\slang-generate.vcxproj", "{66174227-8541-41FC-A6DF-4764FC66F78E}"
+EndProject
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "slang-test", "tools\slang-test\slang-test.vcxproj", "{0C768A18-1D25-4000-9F37-DA5FE99E3B64}"
 EndProject
 Global
 	GlobalSection(SolutionConfigurationPlatforms) = preSolution
@@ -42,6 +42,30 @@ Global
 		Release|x64 = Release|x64
 	EndGlobalSection
 	GlobalSection(ProjectConfigurationPlatforms) = postSolution
+		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Debug|Win32.ActiveCfg = Debug|Win32
+		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Debug|Win32.Build.0 = Debug|Win32
+		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Debug|x64.ActiveCfg = Debug|x64
+		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Debug|x64.Build.0 = Debug|x64
+		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|Win32.ActiveCfg = Release|Win32
+		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|Win32.Build.0 = Release|Win32
+		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|x64.ActiveCfg = Release|x64
+		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|x64.Build.0 = Release|x64
+		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Debug|Win32.ActiveCfg = Debug|Win32
+		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Debug|Win32.Build.0 = Debug|Win32
+		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Debug|x64.ActiveCfg = Debug|x64
+		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Debug|x64.Build.0 = Debug|x64
+		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Release|Win32.ActiveCfg = Release|Win32
+		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Release|Win32.Build.0 = Release|Win32
+		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Release|x64.ActiveCfg = Release|x64
+		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Release|x64.Build.0 = Release|x64
+		{57C81DD3-4304-213D-AC16-39349871C957}.Debug|Win32.ActiveCfg = Debug|Win32
+		{57C81DD3-4304-213D-AC16-39349871C957}.Debug|Win32.Build.0 = Debug|Win32
+		{57C81DD3-4304-213D-AC16-39349871C957}.Debug|x64.ActiveCfg = Debug|x64
+		{57C81DD3-4304-213D-AC16-39349871C957}.Debug|x64.Build.0 = Debug|x64
+		{57C81DD3-4304-213D-AC16-39349871C957}.Release|Win32.ActiveCfg = Release|Win32
+		{57C81DD3-4304-213D-AC16-39349871C957}.Release|Win32.Build.0 = Release|Win32
+		{57C81DD3-4304-213D-AC16-39349871C957}.Release|x64.ActiveCfg = Release|x64
+		{57C81DD3-4304-213D-AC16-39349871C957}.Release|x64.Build.0 = Release|x64
 		{010BE414-ED5B-CF56-16C0-BD18027062C0}.Debug|Win32.ActiveCfg = Debug|Win32
 		{010BE414-ED5B-CF56-16C0-BD18027062C0}.Debug|Win32.Build.0 = Debug|Win32
 		{010BE414-ED5B-CF56-16C0-BD18027062C0}.Debug|x64.ActiveCfg = Debug|x64
@@ -58,30 +82,46 @@ Global
 		{2F8724C6-1BC3-2730-84D5-3F277030D04A}.Release|Win32.Build.0 = Release|Win32
 		{2F8724C6-1BC3-2730-84D5-3F277030D04A}.Release|x64.ActiveCfg = Release|x64
 		{2F8724C6-1BC3-2730-84D5-3F277030D04A}.Release|x64.Build.0 = Release|x64
-		{57C81DD3-4304-213D-AC16-39349871C957}.Debug|Win32.ActiveCfg = Debug|Win32
-		{57C81DD3-4304-213D-AC16-39349871C957}.Debug|Win32.Build.0 = Debug|Win32
-		{57C81DD3-4304-213D-AC16-39349871C957}.Debug|x64.ActiveCfg = Debug|x64
-		{57C81DD3-4304-213D-AC16-39349871C957}.Debug|x64.Build.0 = Debug|x64
-		{57C81DD3-4304-213D-AC16-39349871C957}.Release|Win32.ActiveCfg = Release|Win32
-		{57C81DD3-4304-213D-AC16-39349871C957}.Release|Win32.Build.0 = Release|Win32
-		{57C81DD3-4304-213D-AC16-39349871C957}.Release|x64.ActiveCfg = Release|x64
-		{57C81DD3-4304-213D-AC16-39349871C957}.Release|x64.Build.0 = Release|x64
-		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Debug|Win32.ActiveCfg = Debug|Win32
-		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Debug|Win32.Build.0 = Debug|Win32
-		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Debug|x64.ActiveCfg = Debug|x64
-		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Debug|x64.Build.0 = Debug|x64
-		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Release|Win32.ActiveCfg = Release|Win32
-		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Release|Win32.Build.0 = Release|Win32
-		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Release|x64.ActiveCfg = Release|x64
-		{4B47A364-37C4-96A7-6041-97BB4C1D333B}.Release|x64.Build.0 = Release|x64
-		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Debug|Win32.ActiveCfg = Debug|Win32
-		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Debug|Win32.Build.0 = Debug|Win32
-		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Debug|x64.ActiveCfg = Debug|x64
-		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Debug|x64.Build.0 = Debug|x64
-		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|Win32.ActiveCfg = Release|Win32
-		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|Win32.Build.0 = Release|Win32
-		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|x64.ActiveCfg = Release|x64
-		{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|x64.Build.0 = Release|x64
+		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Debug|Win32.ActiveCfg = Debug|Win32
+		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Debug|Win32.Build.0 = Debug|Win32
+		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Debug|x64.ActiveCfg = Debug|x64
+		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Debug|x64.Build.0 = Debug|x64
+		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Release|Win32.ActiveCfg = Release|Win32
+		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Release|Win32.Build.0 = Release|Win32
+		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Release|x64.ActiveCfg = Release|x64
+		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Release|x64.Build.0 = Release|x64
+		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Debug|Win32.ActiveCfg = Debug|Win32
+		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Debug|Win32.Build.0 = Debug|Win32
+		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Debug|x64.ActiveCfg = Debug|x64
+		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Debug|x64.Build.0 = Debug|x64
+		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Release|Win32.ActiveCfg = Release|Win32
+		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Release|Win32.Build.0 = Release|Win32
+		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Release|x64.ActiveCfg = Release|x64
+		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Release|x64.Build.0 = Release|x64
+		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Debug|Win32.ActiveCfg = Debug|Win32
+		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Debug|Win32.Build.0 = Debug|Win32
+		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Debug|x64.ActiveCfg = Debug|x64
+		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Debug|x64.Build.0 = Debug|x64
+		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Release|Win32.ActiveCfg = Release|Win32
+		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Release|Win32.Build.0 = Release|Win32
+		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Release|x64.ActiveCfg = Release|x64
+		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Release|x64.Build.0 = Release|x64
+		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Debug|Win32.ActiveCfg = Debug|Win32
+		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Debug|Win32.Build.0 = Debug|Win32
+		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Debug|x64.ActiveCfg = Debug|x64
+		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Debug|x64.Build.0 = Debug|x64
+		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Release|Win32.ActiveCfg = Release|Win32
+		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Release|Win32.Build.0 = Release|Win32
+		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Release|x64.ActiveCfg = Release|x64
+		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Release|x64.Build.0 = Release|x64
+		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Debug|Win32.ActiveCfg = Debug|Win32
+		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Debug|Win32.Build.0 = Debug|Win32
+		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Debug|x64.ActiveCfg = Debug|x64
+		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Debug|x64.Build.0 = Debug|x64
+		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Release|Win32.ActiveCfg = Release|Win32
+		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Release|Win32.Build.0 = Release|Win32
+		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Release|x64.ActiveCfg = Release|x64
+		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Release|x64.Build.0 = Release|x64
 		{66174227-8541-41FC-A6DF-4764FC66F78E}.Debug|Win32.ActiveCfg = Debug|Win32
 		{66174227-8541-41FC-A6DF-4764FC66F78E}.Debug|Win32.Build.0 = Debug|Win32
 		{66174227-8541-41FC-A6DF-4764FC66F78E}.Debug|x64.ActiveCfg = Debug|x64
@@ -98,59 +138,19 @@ Global
 		{0C768A18-1D25-4000-9F37-DA5FE99E3B64}.Release|Win32.Build.0 = Release|Win32
 		{0C768A18-1D25-4000-9F37-DA5FE99E3B64}.Release|x64.ActiveCfg = Release|x64
 		{0C768A18-1D25-4000-9F37-DA5FE99E3B64}.Release|x64.Build.0 = Release|x64
-		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Debug|Win32.ActiveCfg = Debug|Win32
-		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Debug|Win32.Build.0 = Debug|Win32
-		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Debug|x64.ActiveCfg = Debug|x64
-		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Debug|x64.Build.0 = Debug|x64
-		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Release|Win32.ActiveCfg = Release|Win32
-		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Release|Win32.Build.0 = Release|Win32
-		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Release|x64.ActiveCfg = Release|x64
-		{222F7498-B40C-4F3F-A704-DDEB91A4484A}.Release|x64.Build.0 = Release|x64
-		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Debug|Win32.ActiveCfg = Debug|Win32
-		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Debug|Win32.Build.0 = Debug|Win32
-		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Debug|x64.ActiveCfg = Debug|x64
-		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Debug|x64.Build.0 = Debug|x64
-		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Release|Win32.ActiveCfg = Release|Win32
-		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Release|Win32.Build.0 = Release|Win32
-		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Release|x64.ActiveCfg = Release|x64
-		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F}.Release|x64.Build.0 = Release|x64
-		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Debug|Win32.ActiveCfg = Debug|Win32
-		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Debug|Win32.Build.0 = Debug|Win32
-		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Debug|x64.ActiveCfg = Debug|x64
-		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Debug|x64.Build.0 = Debug|x64
-		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Release|Win32.ActiveCfg = Release|Win32
-		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Release|Win32.Build.0 = Release|Win32
-		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Release|x64.ActiveCfg = Release|x64
-		{61F7EB00-7281-4BF3-9470-7C2EA92620C3}.Release|x64.Build.0 = Release|x64
-		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Debug|Win32.ActiveCfg = Debug|Win32
-		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Debug|Win32.Build.0 = Debug|Win32
-		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Debug|x64.ActiveCfg = Debug|x64
-		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Debug|x64.Build.0 = Debug|x64
-		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Release|Win32.ActiveCfg = Release|Win32
-		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Release|Win32.Build.0 = Release|Win32
-		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Release|x64.ActiveCfg = Release|x64
-		{D56CBCEB-1EB5-4CA8-AEC4-48EA35ED61C7}.Release|x64.Build.0 = Release|x64
-		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Debug|Win32.ActiveCfg = Debug|Win32
-		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Debug|Win32.Build.0 = Debug|Win32
-		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Debug|x64.ActiveCfg = Debug|x64
-		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Debug|x64.Build.0 = Debug|x64
-		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Release|Win32.ActiveCfg = Release|Win32
-		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Release|Win32.Build.0 = Release|Win32
-		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Release|x64.ActiveCfg = Release|x64
-		{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}.Release|x64.Build.0 = Release|x64
 	EndGlobalSection
 	GlobalSection(SolutionProperties) = preSolution
 		HideSolutionNode = FALSE
 	EndGlobalSection
 	GlobalSection(NestedProjects) = preSolution
+		{4B47A364-37C4-96A7-6041-97BB4C1D333B} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}
+		{57C81DD3-4304-213D-AC16-39349871C957} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}
 		{010BE414-ED5B-CF56-16C0-BD18027062C0} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}
 		{2F8724C6-1BC3-2730-84D5-3F277030D04A} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}
-		{57C81DD3-4304-213D-AC16-39349871C957} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}
-		{4B47A364-37C4-96A7-6041-97BB4C1D333B} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}
+		{61F7EB00-7281-4BF3-9470-7C2EA92620C3} = {57B5AA5E-C340-1823-CC51-9B17385C7423}
+		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F} = {57B5AA5E-C340-1823-CC51-9B17385C7423}
+		{222F7498-B40C-4F3F-A704-DDEB91A4484A} = {FD47AE19-69FD-260F-F2F1-20E65EA61D13}
 		{66174227-8541-41FC-A6DF-4764FC66F78E} = {FD47AE19-69FD-260F-F2F1-20E65EA61D13}
 		{0C768A18-1D25-4000-9F37-DA5FE99E3B64} = {FD47AE19-69FD-260F-F2F1-20E65EA61D13}
-		{222F7498-B40C-4F3F-A704-DDEB91A4484A} = {FD47AE19-69FD-260F-F2F1-20E65EA61D13}
-		{C5ACCA6E-C04D-4B36-8516-3752B3C13C2F} = {57B5AA5E-C340-1823-CC51-9B17385C7423}
-		{61F7EB00-7281-4BF3-9470-7C2EA92620C3} = {57B5AA5E-C340-1823-CC51-9B17385C7423}
 	EndGlobalSection
 EndGlobal
diff --git a/source/core/slang-downstream-compiler.h b/source/core/slang-downstream-compiler.h
index 3ffa32097..bcf53d065 100644
--- a/source/core/slang-downstream-compiler.h
+++ b/source/core/slang-downstream-compiler.h
@@ -202,6 +202,14 @@ public:
         Object,             ///< Produce an object file
     };
 
+    enum PipelineType
+    {
+        Unknown,
+        Compute,
+        Rasterization,
+        RayTracing,
+    };
+
     struct Define
     {
         String nameWithSig;             ///< If macro takes parameters include in brackets
@@ -236,6 +244,7 @@ public:
         TargetType targetType = TargetType::Executable;
         SlangSourceLanguage sourceLanguage = SLANG_SOURCE_LANGUAGE_CPP;
         FloatingPointMode floatingPointMode = FloatingPointMode::Default;
+        PipelineType pipelineType = PipelineType::Unknown;
 
         Flags flags = Flag::EnableExceptionHandling;
 
diff --git a/source/core/slang-nvrtc-compiler.cpp b/source/core/slang-nvrtc-compiler.cpp
index 0e167bf80..6b01ef070 100644
--- a/source/core/slang-nvrtc-compiler.cpp
+++ b/source/core/slang-nvrtc-compiler.cpp
@@ -334,8 +334,63 @@ SlangResult NVRTCDownstreamCompiler::compile(const CompileOptions& options, RefP
         cmdLine.addArg(builder);
     }
 
+    List<const char*> headers;
+    List<const char*> headerIncludeNames;
+
+    // If compiling for OptiX, we need to add the appropriate search paths to the command line.
+    //
+    if(options.pipelineType == PipelineType::RayTracing)
+    {
+        // The device-side OptiX API is accessed through a constellation
+        // of headers provided by the OptiX SDK, so we need to set an
+        // include path for the compile that makes those visible.
+        //
+        // TODO: The OptiX SDK installer doesn't set any kind of environment
+        // variable to indicate where the SDK was installed, so we seemingly
+        // need to probe paths instead. The form of the path will differ
+        // betwene Windows and Unix-y platforms, and we will need some kind
+        // of approach to probe multiple versiosn and use the latest.
+        //
+        // HACK: For now I'm using the fixed path for the most recent SDK
+        // release on Windows. This means that OptiX cross-compilation will
+        // only "work" on a subset of platforms, but that doesn't matter
+        // for now since it doesn't really "work" at all.
+        //
+        cmdLine.addArg("-I");
+        cmdLine.addArg("C:/ProgramData/NVIDIA Corporation/OptiX SDK 7.0.0/include/");
+
+        // The OptiX headers in turn `#include <stddef.h>` and expect that
+        // to work. We could try to also add in an include path from the CUDA
+        // SDK (which seems to provide a `stddef.h` in the most recent version),
+        // but using that version doesn't seem to work (and also bakes in a
+        // requirement that the user have the CUDA SDK installed in addition
+        // to the OptiX SDK).
+        //
+        // Instead, we will rely on the NVRTC feature that lets us set up
+        // memory buffers to be used as include files by the we compile.
+        // We will define a dummy `stddef.h` that includes the bare minimum
+        // lines required to get the OptiX headers to compile without complaint.
+        //
+        // TODO: Confirm that the `LP64` definition herei s actually needed.
+        //
+        headerIncludeNames.add("stddef.h");
+        headers.add("#pragma once\n" "#define LP64\n");
+
+        // Finally, we want the CUDA prelude to be able to react to whether
+        // or not OptiX is required (most notably by `#include`ing the appropriate
+        // header(s)), so we will insert a preprocessor define to indicate
+        // the requirement.
+        //
+        cmdLine.addArg("-DSLANG_CUDA_ENABLE_OPTIX");
+    }
+
+    SLANG_ASSERT(headers.getCount() == headerIncludeNames.getCount());
+
     nvrtcProgram program = nullptr;
-    nvrtcResult res = m_nvrtcCreateProgram(&program, options.sourceContents.getBuffer(), options.sourceContentsPath.getBuffer(), 0, nullptr, nullptr);
+    nvrtcResult res = m_nvrtcCreateProgram(&program, options.sourceContents.getBuffer(), options.sourceContentsPath.getBuffer(),
+        (int) headers.getCount(),
+        headers.getBuffer(),
+        headerIncludeNames.getBuffer());
     if (res != NVRTC_SUCCESS)
     {
         return _asResult(res);
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index 483c2299f..4279e4a4e 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -3162,6 +3162,7 @@ void AcceptHitAndEndSearch();
 // 10.4.1 - Ray Dispatch System Values
 
 __target_intrinsic(glsl, "(gl_LaunchIDNV)")
+__target_intrinsic(cuda, "optixGetLaunchIndex")
 uint3 DispatchRaysIndex();
 
 __target_intrinsic(glsl, "(gl_LaunchSizeNV)")
diff --git a/source/slang/slang-compiler.cpp b/source/slang/slang-compiler.cpp
index 53a028483..f81a63275 100644
--- a/source/slang/slang-compiler.cpp
+++ b/source/slang/slang-compiler.cpp
@@ -1373,6 +1373,56 @@ SlangResult dissassembleDXILUsingDXC(
                 default: SLANG_ASSERT(!"Unhandled floating point mode");
             }
 
+            {
+                // We need to look at the stage of the entry point(s) we are
+                // being asked to compile, since this will determine the
+                // "pipeline" that the result should be compiled for (e.g.,
+                // compute vs. ray tracing).
+                //
+                // TODO: This logic is kind of messy in that it assumes
+                // a program to be compiled will only contain kernels for
+                // a single pipeline type, but that invariant isn't expressed
+                // at all in the front-end today. It also has no error
+                // checking for the case where there are conflicts.
+                //
+                // HACK: Right now none of the above concerns matter
+                // because we always perform code generation on a single
+                // entry point at a time.
+                //
+                Index entryPointCount = slangRequest->getProgram()->getEntryPointCount();
+                for(Index ee = 0; ee < entryPointCount; ++ee)
+                {
+                    auto stage = slangRequest->getProgram()->getEntryPoint(ee)->getStage();
+                    switch(stage)
+                    {
+                    default:
+                        break;
+
+                    case Stage::Compute:
+                        options.pipelineType = DownstreamCompiler::PipelineType::Compute;
+                        break;
+
+                    case Stage::Vertex:
+                    case Stage::Hull:
+                    case Stage::Domain:
+                    case Stage::Geometry:
+                    case Stage::Fragment:
+                        options.pipelineType = DownstreamCompiler::PipelineType::Rasterization;
+                        break;
+
+                    case Stage::RayGeneration:
+                    case Stage::Intersection:
+                    case Stage::AnyHit:
+                    case Stage::ClosestHit:
+                    case Stage::Miss:
+                    case Stage::Callable:
+                        options.pipelineType = DownstreamCompiler::PipelineType::RayTracing;
+                        break;
+                    }
+                }
+                
+            }
+
             // Add all the search paths (as calculated earlier - they will only be set if this is a pass through else will be empty)
             options.includePaths = includePaths;
 
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index a728df755..702543fc8 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -606,21 +606,57 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
 
             IREntryPointDecoration* entryPointDecor = func->findDecoration<IREntryPointDecoration>();
 
-            if (entryPointDecor && entryPointDecor->getProfile().GetStage() == Stage::Compute)
+            if (entryPointDecor)
             {
-                Int sizeAlongAxis[kThreadGroupAxisCount];
-                getComputeThreadGroupSize(func, sizeAlongAxis);
-
-                // 
-                m_writer->emit("// [numthreads(");
-                for (int ii = 0; ii < kThreadGroupAxisCount; ++ii)
+                // We have an entry-point function in the IR module, which we
+                // will want to emit as a `__global__` function in the generated
+                // CUDA C++.
+                //
+                // The most common case will be a compute kernel, in which case
+                // we will emit the function more or less as-is, including
+                // usingits original name as the name of the global symbol.
+                //
+                String funcName = getName(func);
+                String globalSymbolName = funcName;
+
+                // We also suport emitting ray tracing kernels for use with
+                // OptiX, and in that case the name of the global symbol
+                // must be prefixed to indicate to the OptiX runtime what
+                // stage it is to be compiled for.
+                //
+                auto stage = entryPointDecor->getProfile().GetStage();
+                switch( stage )
                 {
-                    if (ii != 0) m_writer->emit(", ");
-                    m_writer->emit(sizeAlongAxis[ii]);
+                default:
+                    break;
+
+            #define CASE(STAGE, PREFIX) \
+                case Stage::STAGE: globalSymbolName = #PREFIX + funcName; break
+
+                CASE(RayGeneration, __raygen__);
+                // TODO: Add the other ray tracing shader stages here.
+            #undef CASE
                 }
-                m_writer->emit(")]\n");
 
-                String funcName = getName(func);
+                // As a convenience for anybody reading the generated
+                // CUDA C++ code, we will prefix a compute kernel
+                // with the information from the `[numthreads(...)]`
+                // attribute in the source.
+                //
+                if(stage == Stage::Compute)
+                {
+                    Int sizeAlongAxis[kThreadGroupAxisCount];
+                    getComputeThreadGroupSize(func, sizeAlongAxis);
+
+                    // 
+                    m_writer->emit("// [numthreads(");
+                    for (int ii = 0; ii < kThreadGroupAxisCount; ++ii)
+                    {
+                        if (ii != 0) m_writer->emit(", ");
+                        m_writer->emit(sizeAlongAxis[ii]);
+                    }
+                    m_writer->emit(")]\n");
+                }
 
                 m_writer->emit("extern \"C\" __global__  ");
                
@@ -628,7 +664,7 @@ void CUDASourceEmitter::emitModuleImpl(IRModule* module)
 
                 // Emit the actual function
                 emitEntryPointAttributes(func, entryPointDecor);
-                emitType(resultType, funcName);
+                emitType(resultType, globalSymbolName);
 
                 m_writer->emit("(UniformEntryPointParams* params, UniformState* uniformState)");
                 emitSemantics(func);
diff --git a/source/slang/slang.vcxproj b/source/slang/slang.vcxproj
index 027c43a95..766893da3 100644
--- a/source/slang/slang.vcxproj
+++ b/source/slang/slang.vcxproj
@@ -188,7 +188,6 @@
   <ItemGroup>
     <ClInclude Include="..\..\slang.h" />
     <ClInclude Include="core.meta.slang.h" />
-    <ClInclude Include="glsl.meta.slang.h" />
     <ClInclude Include="hlsl.meta.slang.h" />
     <ClInclude Include="slang-check-impl.h" />
     <ClInclude Include="slang-check.h" />
@@ -345,17 +344,13 @@
     <ClCompile Include="slang.cpp" />
   </ItemGroup>
   <ItemGroup>
-    <None Include="..\core\core.natvis" />
-    <None Include="slang.natvis" />
-  </ItemGroup>
-  <ItemGroup>
     <CustomBuild Include="core.meta.slang">
       <FileType>Document</FileType>
       <Command Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">"../../bin/windows-x86/debug/slang-generate" %(Identity)</Command>
       <Command Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">"../../bin/windows-x64/debug/slang-generate" %(Identity)</Command>
       <Command Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">"../../bin/windows-x86/release/slang-generate" %(Identity)</Command>
       <Command Condition="'$(Configuration)|$(Platform)'=='Release|x64'">"../../bin/windows-x64/release/slang-generate" %(Identity)</Command>
-      <Outputs>%(Identity).h</Outputs>
+      <Outputs>../../core.meta.slang.h</Outputs>
       <Message>slang-generate %(Identity)</Message>
       <AdditionalInputs Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">../../bin/windows-x86/debug/slang-generate.exe</AdditionalInputs>
       <AdditionalInputs Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">../../bin/windows-x64/debug/slang-generate.exe</AdditionalInputs>
@@ -368,7 +363,7 @@
       <Command Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">"../../bin/windows-x64/debug/slang-generate" %(Identity)</Command>
       <Command Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">"../../bin/windows-x86/release/slang-generate" %(Identity)</Command>
       <Command Condition="'$(Configuration)|$(Platform)'=='Release|x64'">"../../bin/windows-x64/release/slang-generate" %(Identity)</Command>
-      <Outputs>%(Identity).h</Outputs>
+      <Outputs>../../hlsl.meta.slang.h</Outputs>
       <Message>slang-generate %(Identity)</Message>
       <AdditionalInputs Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">../../bin/windows-x86/debug/slang-generate.exe</AdditionalInputs>
       <AdditionalInputs Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">../../bin/windows-x64/debug/slang-generate.exe</AdditionalInputs>
@@ -377,6 +372,10 @@
     </CustomBuild>
   </ItemGroup>
   <ItemGroup>
+    <Natvis Include="..\core\core.natvis" />
+    <Natvis Include="slang.natvis" />
+  </ItemGroup>
+  <ItemGroup>
     <ProjectReference Include="..\core\core.vcxproj">
       <Project>{F9BE7957-8399-899E-0C49-E714FDDD4B65}</Project>
     </ProjectReference>
diff --git a/source/slang/slang.vcxproj.filters b/source/slang/slang.vcxproj.filters
index 9be567654..442f545c6 100644
--- a/source/slang/slang.vcxproj.filters
+++ b/source/slang/slang.vcxproj.filters
@@ -1,4 +1,4 @@
-<?xml version="1.0" encoding="utf-8"?>
+<?xml version="1.0" encoding="utf-8"?>
 <Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
   <ItemGroup>
     <Filter Include="Header Files">
@@ -15,9 +15,6 @@
     <ClInclude Include="core.meta.slang.h">
       <Filter>Header Files</Filter>
     </ClInclude>
-    <ClInclude Include="glsl.meta.slang.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
     <ClInclude Include="hlsl.meta.slang.h">
       <Filter>Header Files</Filter>
     </ClInclude>
@@ -99,6 +96,9 @@
     <ClInclude Include="slang-ir-glsl-legalize.h">
       <Filter>Header Files</Filter>
     </ClInclude>
+    <ClInclude Include="slang-ir-inline.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
     <ClInclude Include="slang-ir-inst-defs.h">
       <Filter>Header Files</Filter>
     </ClInclude>
@@ -138,6 +138,9 @@
     <ClInclude Include="slang-ir-string-hash.h">
       <Filter>Header Files</Filter>
     </ClInclude>
+    <ClInclude Include="slang-ir-strip-witness-tables.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
     <ClInclude Include="slang-ir-strip.h">
       <Filter>Header Files</Filter>
     </ClInclude>
@@ -246,12 +249,6 @@
     <ClInclude Include="slang-visitor.h">
       <Filter>Header Files</Filter>
     </ClInclude>
-    <ClInclude Include="slang-ir-strip-witness-tables.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
-    <ClInclude Include="slang-ir-inline.h">
-      <Filter>Header Files</Filter>
-    </ClInclude>
   </ItemGroup>
   <ItemGroup>
     <ClCompile Include="slang-check-conformance.cpp">
@@ -353,6 +350,9 @@
     <ClCompile Include="slang-ir-glsl-legalize.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
+    <ClCompile Include="slang-ir-inline.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
     <ClCompile Include="slang-ir-legalize-types.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
@@ -389,6 +389,9 @@
     <ClCompile Include="slang-ir-string-hash.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
+    <ClCompile Include="slang-ir-strip-witness-tables.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
     <ClCompile Include="slang-ir-strip.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
@@ -470,20 +473,6 @@
     <ClCompile Include="slang.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
-    <ClCompile Include="slang-ir-strip-witness-tables.cpp">
-      <Filter>Source Files</Filter>
-    </ClCompile>
-    <ClCompile Include="slang-ir-inline.cpp">
-      <Filter>Source Files</Filter>
-    </ClCompile>
-  </ItemGroup>
-  <ItemGroup>
-    <None Include="..\core\core.natvis">
-      <Filter>Source Files</Filter>
-    </None>
-    <None Include="slang.natvis">
-      <Filter>Source Files</Filter>
-    </None>
   </ItemGroup>
   <ItemGroup>
     <CustomBuild Include="core.meta.slang">
@@ -493,4 +482,12 @@
       <Filter>Source Files</Filter>
     </CustomBuild>
   </ItemGroup>
+  <ItemGroup>
+    <Natvis Include="..\core\core.natvis">
+      <Filter>Source Files</Filter>
+    </Natvis>
+    <Natvis Include="slang.natvis">
+      <Filter>Source Files</Filter>
+    </Natvis>
+  </ItemGroup>
 </Project>
\ No newline at end of file
diff --git a/tools/gfx/render.h b/tools/gfx/render.h
index a4d042a9a..423820a0b 100644
--- a/tools/gfx/render.h
+++ b/tools/gfx/render.h
@@ -45,6 +45,7 @@ enum class PipelineType
     Unknown,
     Graphics,
     Compute,
+    RayTracing,
     CountOf,
 };
 
@@ -57,6 +58,12 @@ enum class StageType
     Geometry,
     Fragment,
     Compute,
+    RayGeneration,
+    Intersection,
+    AnyHit,
+    ClosestHit,
+    Miss,
+    Callable,
     CountOf,
 };
 
@@ -102,6 +109,7 @@ public:
         StageType   stage;
         void const* codeBegin;
         void const* codeEnd;
+        char const* entryPointName;
 
         UInt getCodeSize() const { return (char const*)codeEnd - (char const*)codeBegin; }
     };
@@ -141,13 +149,12 @@ struct ShaderCompileRequest
     struct EntryPoint
     {
         char const* name = nullptr;
-        SourceInfo  source;
+        SlangStage  slangStage;
     };
 
     SourceInfo source;
-    EntryPoint vertexShader;
-    EntryPoint fragmentShader;
-    EntryPoint computeShader;
+    Slang::List<EntryPoint> entryPoints;
+
     Slang::List<Slang::String> globalSpecializationArgs;
     Slang::List<Slang::String> entryPointSpecializationArgs;
 
diff --git a/tools/render-test/cuda/cuda-compute-util.cpp b/tools/render-test/cuda/cuda-compute-util.cpp
index 48d73fa93..b2006a7e8 100644
--- a/tools/render-test/cuda/cuda-compute-util.cpp
+++ b/tools/render-test/cuda/cuda-compute-util.cpp
@@ -13,6 +13,14 @@
 
 #include <cuda_runtime_api.h>
 
+// TODO: should conditionalize this on OptiX being present
+#ifdef RENDER_TEST_OPTIX
+#include <optix.h>
+#include <optix_function_table_definition.h>
+#define _CRT_SECURE_NO_WARNINGS 1
+#include <optix_stubs.h>
+#endif
+
 namespace renderer_test {
 using namespace Slang;
 
@@ -111,6 +119,36 @@ static SlangResult _handleCUDAError(cudaError_t error)
 
 #define SLANG_CUDA_ASSERT_ON_FAIL(x) { auto _res = x; if (_isError(_res)) { SLANG_ASSERT(!"Failed CUDA call"); }; }
 
+#ifdef RENDER_TEST_OPTIX
+
+static bool _isError(OptixResult result) { return result != OPTIX_SUCCESS; }
+
+#if 1
+static SlangResult _handleOptixError(OptixResult result, char const* file, int line)
+{
+    fprintf(stderr, "%s(%d): optix: %s (%s)\n",
+        file,
+        line,
+        optixGetErrorString(result),
+        optixGetErrorName(result));
+    return SLANG_FAIL;
+}
+#define SLANG_OPTIX_HANDLE_ERROR(RESULT) _handleOptixError(RESULT, __FILE__, __LINE__)
+#else
+#define SLANG_OPTIX_HANDLE_ERROR(RESULT) SLANG_FAIL
+#endif
+
+#define SLANG_OPTIX_RETURN_ON_FAIL(EXPR) do { auto _res = EXPR; if(_isError(_res)) return SLANG_OPTIX_HANDLE_ERROR(_res); } while(0)
+
+void _optixLogCallback(unsigned int level, const char* tag, const char* message, void* userData)
+{
+    fprintf(stderr, "optix: %s (%s)\n",
+        message,
+        tag);
+}
+
+#endif
+
 class MemoryCUDAResource : public CUDAResource
 {
 public:
@@ -1202,17 +1240,238 @@ static SlangResult _compute(CUcontext context, CUmodule module, const ShaderComp
     ScopeCUDAContext cudaContext;
     SLANG_RETURN_ON_FAIL(cudaContext.init(0));
 
-    const Index index = outputAndLayout.output.findKernelDescIndex(StageType::Compute);
-    if (index < 0)
+
+    switch( outputAndLayout.output.desc.pipelineType )
     {
+    default:
         return SLANG_FAIL;
-    }
 
-    const auto& kernel = outputAndLayout.output.kernelDescs[index];
+    case PipelineType::Compute:
+        {
+            const Index index = outputAndLayout.output.findKernelDescIndex(StageType::Compute);
+            if (index < 0)
+            {
+                return SLANG_FAIL;
+            }
+
+            const auto& kernel = outputAndLayout.output.kernelDescs[index];
+
+            ScopeCUDAModule cudaModule;
+            SLANG_RETURN_ON_FAIL(cudaModule.load(kernel.codeBegin));
+            SLANG_RETURN_ON_FAIL(_compute(cudaContext, cudaModule, outputAndLayout, dispatchSize, outContext));
+        }
+        break;
+
+    case PipelineType::RayTracing:
+        {
+#ifdef RENDER_TEST_OPTIX
+            SLANG_OPTIX_RETURN_ON_FAIL(optixInit());
+
+            OptixDeviceContextOptions optixOptions = {};
+
+            // TODO: set log callback
+            optixOptions.logCallbackFunction = &_optixLogCallback;
+            optixOptions.logCallbackLevel = 4;
+
+            OptixDeviceContext optixContext = nullptr;
+            SLANG_OPTIX_RETURN_ON_FAIL(optixDeviceContextCreate(cudaContext, &optixOptions, &optixContext));
+
+            enum { kOptixLogSize = 2*1024 };
+            char log[kOptixLogSize];
+            size_t logSize = sizeof(log);
+
+            OptixPipelineCompileOptions optixPipelineCompileOptions = {};
+
+            // We need to load modules from the PTX code available to us,
+            // and then also create program groups from the kernels
+            // in those modules.
+            //
+            // For now we will only support program groups with a single
+            // kernel in them, and will create one per entry point.
+            //
+            Index entryPointCount = outputAndLayout.output.kernelDescs.getCount();
+            List<OptixProgramGroup> optixProgramGroups;
+            List<String> names;
+
+            OptixShaderBindingTable optixSBT = {};
+
+            for( Index ee = 0; ee < entryPointCount; ++ee )
+            {
+                auto& kernel = outputAndLayout.output.kernelDescs[ee];
 
-    ScopeCUDAModule cudaModule;
-    SLANG_RETURN_ON_FAIL(cudaModule.load(kernel.codeBegin));
-    SLANG_RETURN_ON_FAIL(_compute(cudaContext, cudaModule, outputAndLayout, dispatchSize, outContext));
+                OptixModuleCompileOptions optixModuleCompileOptions = {};
+
+                OptixModule optixModule;
+                SLANG_OPTIX_RETURN_ON_FAIL(optixModuleCreateFromPTX(
+                    optixContext,
+                    &optixModuleCompileOptions,
+                    &optixPipelineCompileOptions,
+                    (char const*) kernel.codeBegin,
+                    kernel.getCodeSize(),
+                    log,
+                    &logSize,
+                    &optixModule));
+
+
+                OptixProgramGroupOptions optixProgramGroupOptions = {};
+
+                OptixProgramGroupDesc optixProgramGroupDesc = {};
+                optixProgramGroupDesc.kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
+                optixProgramGroupDesc.raygen.module = optixModule;
+
+                String name = String("__raygen__") + kernel.entryPointName;
+                names.add(name);
+                optixProgramGroupDesc.raygen.entryFunctionName = name.begin();
+
+                OptixProgramGroup optixProgramGroup = nullptr;
+                SLANG_OPTIX_RETURN_ON_FAIL(optixProgramGroupCreate(
+                    optixContext,
+                    &optixProgramGroupDesc,
+                    1,
+                    &optixProgramGroupOptions,
+                    log,
+                    &logSize,
+                    &optixProgramGroup));
+
+                optixProgramGroups.add(optixProgramGroup);
+
+                {
+                    CUdeviceptr rayGenRecordPtr;
+                    size_t rayGenRecordSize = OPTIX_SBT_RECORD_HEADER_SIZE;
+
+                    SLANG_CUDA_RETURN_ON_FAIL(cudaMalloc((void**) &rayGenRecordPtr, rayGenRecordSize));
+
+                    struct { char data[OPTIX_SBT_RECORD_HEADER_SIZE]; } rayGenRecordData;
+                    SLANG_OPTIX_RETURN_ON_FAIL(optixSbtRecordPackHeader(optixProgramGroup, &rayGenRecordData));
+
+                    SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy(
+                        (void*) rayGenRecordPtr,
+                        &rayGenRecordData,
+                        rayGenRecordSize,
+                        cudaMemcpyHostToDevice));
+
+                    optixSBT.raygenRecord = rayGenRecordPtr;
+                }
+            }
+
+
+
+            OptixPipeline optixPipeline = nullptr;
+
+            OptixPipelineLinkOptions optixPipelineLinkOptions = {};
+            optixPipelineLinkOptions.maxTraceDepth = 5;
+            optixPipelineLinkOptions.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_FULL;
+            optixPipelineLinkOptions.overrideUsesMotionBlur = false;
+            SLANG_OPTIX_RETURN_ON_FAIL(optixPipelineCreate(
+                optixContext,
+                &optixPipelineCompileOptions,
+                &optixPipelineLinkOptions,
+                optixProgramGroups.getBuffer(),
+                (unsigned int)optixProgramGroups.getCount(),
+                log,
+                &logSize,
+                &optixPipeline));
+
+
+            {
+                // The OptiX API complains if we don't fill in a miss record
+                // in the SBT, so we will create a dummy one here to represent
+                // the lack of any miss shaders.
+                //
+                OptixProgramGroupOptions optixProgramGroupOptions = {};
+                OptixProgramGroupDesc missGroupDesc = {};
+                missGroupDesc.kind = OPTIX_PROGRAM_GROUP_KIND_MISS;
+                OptixProgramGroup missProgramGroup;
+                SLANG_OPTIX_RETURN_ON_FAIL(optixProgramGroupCreate(
+                    optixContext,
+                    &missGroupDesc,
+                    1,
+                    &optixProgramGroupOptions,
+                    log,
+                    &logSize,
+                    &missProgramGroup));
+
+
+                CUdeviceptr missRecordPtr;
+                size_t missRecordSize = OPTIX_SBT_RECORD_HEADER_SIZE;
+
+                SLANG_CUDA_RETURN_ON_FAIL(cudaMalloc((void**) &missRecordPtr, missRecordSize));
+
+                struct { char data[OPTIX_SBT_RECORD_HEADER_SIZE]; } missRecordData;
+                SLANG_OPTIX_RETURN_ON_FAIL(optixSbtRecordPackHeader(missProgramGroup, &missRecordData));
+
+                SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy(
+                    (void*) missRecordPtr,
+                    &missRecordData,
+                    missRecordSize,
+                    cudaMemcpyHostToDevice));
+
+                optixSBT.missRecordBase = missRecordPtr;
+                optixSBT.missRecordCount = 1;
+                optixSBT.missRecordStrideInBytes = missRecordSize;
+            }
+            {
+                // Okay, we also need a dummy hit group.
+
+                OptixProgramGroupOptions optixProgramGroupOptions = {};
+                OptixProgramGroupDesc hitGroupDesc = {};
+                hitGroupDesc.kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
+                OptixProgramGroup programGroup;
+                SLANG_OPTIX_RETURN_ON_FAIL(optixProgramGroupCreate(
+                    optixContext,
+                    &hitGroupDesc,
+                    1,
+                    &optixProgramGroupOptions,
+                    log,
+                    &logSize,
+                    &programGroup));
+
+
+                CUdeviceptr recordPtr;
+                size_t recordSize = OPTIX_SBT_RECORD_HEADER_SIZE;
+
+                SLANG_CUDA_RETURN_ON_FAIL(cudaMalloc((void**) &recordPtr, recordSize));
+
+                struct { char data[OPTIX_SBT_RECORD_HEADER_SIZE]; } recordData;
+                SLANG_OPTIX_RETURN_ON_FAIL(optixSbtRecordPackHeader(programGroup, &recordData));
+
+                SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy(
+                    (void*) recordPtr,
+                    &recordData,
+                    recordSize,
+                    cudaMemcpyHostToDevice));
+
+                optixSBT.hitgroupRecordBase = recordPtr;
+                optixSBT.hitgroupRecordCount = 1;
+                optixSBT.hitgroupRecordStrideInBytes = recordSize;
+            }
+
+            ScopeCUDAStream cudaStream;
+
+            CUdeviceptr globalParams = 0;
+            size_t globalParamsSize = 0;
+
+            unsigned int gridSizeX = 1;
+            unsigned int gridSizeY = 1;
+            unsigned int gridSizeZ = 1;
+
+
+            SLANG_OPTIX_RETURN_ON_FAIL(optixLaunch(
+                optixPipeline,
+                cudaStream,
+                globalParams,
+                globalParamsSize,
+                &optixSBT,
+                gridSizeX,
+                gridSizeY,
+                gridSizeZ));
+
+
+            SLANG_RETURN_ON_FAIL(cudaStream.sync());
+#endif
+        }
+        break;
+    }
 
     return SLANG_OK;
 }
diff --git a/tools/render-test/options.cpp b/tools/render-test/options.cpp
index 8331de07e..c2afe78ac 100644
--- a/tools/render-test/options.cpp
+++ b/tools/render-test/options.cpp
@@ -152,6 +152,10 @@ SlangResult parseOptions(int argc, const char*const* argv, Slang::WriterHelper s
         {
             gOptions.shaderType = ShaderProgramType::GraphicsCompute;
         }
+        else if (strcmp(arg, "-rt") == 0)
+        {
+            gOptions.shaderType = ShaderProgramType::RayTracing;
+        }
         else if( strcmp(arg, "-use-dxil") == 0 )
         {
             gOptions.useDXIL = true;
diff --git a/tools/render-test/options.h b/tools/render-test/options.h
index a8b7d5884..f2f0a8ab6 100644
--- a/tools/render-test/options.h
+++ b/tools/render-test/options.h
@@ -36,7 +36,8 @@ struct Options
     {
         Graphics,
         Compute,
-        GraphicsCompute
+        GraphicsCompute,
+        RayTracing,
     };
 
     char const* appName = "render-test";
diff --git a/tools/render-test/render-test-main.cpp b/tools/render-test/render-test-main.cpp
index 1d88ee500..f966d150e 100644
--- a/tools/render-test/render-test-main.cpp
+++ b/tools/render-test/render-test-main.cpp
@@ -212,7 +212,7 @@ SlangResult RenderTestApp::initialize(SlangSession* session, Renderer* renderer,
 
 Result RenderTestApp::_initializeShaders(SlangSession* session, Renderer* renderer, Options::ShaderProgramType shaderType, const ShaderCompilerUtil::Input& input)
 {
-    SLANG_RETURN_ON_FAIL(ShaderCompilerUtil::compileWithLayout(session, gOptions.sourcePath, gOptions.compileArgs, gOptions.shaderType, input,  m_compilationOutput));
+    SLANG_RETURN_ON_FAIL(ShaderCompilerUtil::compileWithLayout(session, gOptions, input,  m_compilationOutput));
     m_shaderInputLayout = m_compilationOutput.layout;
     m_shaderProgram = renderer->createProgram(m_compilationOutput.output.desc);
     return m_shaderProgram ? SLANG_OK : SLANG_FAIL;
@@ -500,6 +500,25 @@ static SlangResult _innerMain(Slang::StdWriters* stdWriters, SlangSession* sessi
             break;
     }
 
+    switch( gOptions.shaderType )
+    {
+    case Options::ShaderProgramType::Graphics:
+    case Options::ShaderProgramType::GraphicsCompute:
+        input.pipelineType = PipelineType::Graphics;
+        break;
+
+    case Options::ShaderProgramType::Compute:
+        input.pipelineType = PipelineType::Compute;
+        break;
+
+    case Options::ShaderProgramType::RayTracing:
+        input.pipelineType = PipelineType::RayTracing;
+        break;
+
+    default:
+        break;
+    }
+
     if (gOptions.sourceLanguage != SLANG_SOURCE_LANGUAGE_UNKNOWN)
     {
         input.sourceLanguage = gOptions.sourceLanguage;
@@ -554,7 +573,7 @@ static SlangResult _innerMain(Slang::StdWriters* stdWriters, SlangSession* sessi
         }
 
         ShaderCompilerUtil::OutputAndLayout compilationAndLayout;
-        SLANG_RETURN_ON_FAIL(ShaderCompilerUtil::compileWithLayout(session, gOptions.sourcePath, gOptions.compileArgs, gOptions.shaderType, input, compilationAndLayout));
+        SLANG_RETURN_ON_FAIL(ShaderCompilerUtil::compileWithLayout(session, gOptions, input, compilationAndLayout));
 
         {
             // Get the shared library -> it contains the executable code, we need to keep around if we recompile
@@ -575,7 +594,7 @@ static SlangResult _innerMain(Slang::StdWriters* stdWriters, SlangSession* sessi
                 // We just want CPP, so we get suitable reflection
                 slangInput.target = SLANG_CPP_SOURCE;
 
-                SLANG_RETURN_ON_FAIL(ShaderCompilerUtil::compileWithLayout(session, gOptions.sourcePath, gOptions.compileArgs, gOptions.shaderType, slangInput, compilationAndLayout));
+                SLANG_RETURN_ON_FAIL(ShaderCompilerUtil::compileWithLayout(session, gOptions, slangInput, compilationAndLayout));
             }
 
             // calculate binding
diff --git a/tools/render-test/slang-support.cpp b/tools/render-test/slang-support.cpp
index 5afcc6d24..3e5cc9a1c 100644
--- a/tools/render-test/slang-support.cpp
+++ b/tools/render-test/slang-support.cpp
@@ -18,6 +18,37 @@ using namespace Slang;
 static const char vertexEntryPointName[] = "vertexMain";
 static const char fragmentEntryPointName[] = "fragmentMain";
 static const char computeEntryPointName[] = "computeMain";
+static const char rtEntryPointName[] = "raygenMain";
+
+static gfx::StageType _translateStage(SlangStage slangStage)
+{
+    switch(slangStage)
+    {
+    default:
+        SLANG_ASSERT(!"unhandled case");
+        return gfx::StageType::Unknown;
+
+#define CASE(FROM, TO) \
+    case SLANG_STAGE_##FROM: return gfx::StageType::TO
+
+    CASE(VERTEX,    Vertex);
+    CASE(HULL,      Hull);
+    CASE(DOMAIN,    Domain);
+    CASE(GEOMETRY,  Geometry);
+    CASE(FRAGMENT,  Fragment);
+
+    CASE(COMPUTE,   Compute);
+
+    CASE(RAY_GENERATION,    RayGeneration);
+    CASE(INTERSECTION,      Intersection);
+    CASE(ANY_HIT,           AnyHit);
+    CASE(CLOSEST_HIT,       ClosestHit);
+    CASE(MISS,              Miss);
+    CASE(CALLABLE,          Callable);
+
+#undef CASE
+    }
+}
 
 /* static */ SlangResult ShaderCompilerUtil::compileProgram(SlangSession* session, const Input& input, const ShaderCompileRequest& request, Output& out)
 {
@@ -80,46 +111,12 @@ static const char computeEntryPointName[] = "computeMain";
     // the `-xslang <arg>` option to `render-test`.
     SLANG_RETURN_ON_FAIL(spProcessCommandLineArguments(slangRequest, input.args, input.argCount)); 
 
-    int computeTranslationUnit = 0;
-    int vertexTranslationUnit = 0;
-    int fragmentTranslationUnit = 0;
-    char const* vertexEntryPointName = request.vertexShader.name;
-    char const* fragmentEntryPointName = request.fragmentShader.name;
-    char const* computeEntryPointName = request.computeShader.name;
-
     const auto sourceLanguage = input.sourceLanguage;
 
-    if (sourceLanguage == SLANG_SOURCE_LANGUAGE_GLSL)
-    {
-        // GLSL presents unique challenges because, frankly, it got the whole
-        // compilation model wrong. One aspect of working around this is that
-        // we will compile the same source file multiple times: once per
-        // entry point, and we will have different preprocessor definitions
-        // active in each case.
-
-        vertexTranslationUnit = spAddTranslationUnit(slangRequest, sourceLanguage, nullptr);
-        spAddTranslationUnitSourceString(slangRequest, vertexTranslationUnit, request.source.path, request.source.dataBegin);
-        spTranslationUnit_addPreprocessorDefine(slangRequest, vertexTranslationUnit, "__GLSL_VERTEX__", "1");
-        vertexEntryPointName = "main";
-
-        fragmentTranslationUnit = spAddTranslationUnit(slangRequest, sourceLanguage, nullptr);
-        spAddTranslationUnitSourceString(slangRequest, fragmentTranslationUnit, request.source.path, request.source.dataBegin);
-        spTranslationUnit_addPreprocessorDefine(slangRequest, fragmentTranslationUnit, "__GLSL_FRAGMENT__", "1");
-        fragmentEntryPointName = "main";
-
-        computeTranslationUnit = spAddTranslationUnit(slangRequest, sourceLanguage, nullptr);
-        spAddTranslationUnitSourceString(slangRequest, computeTranslationUnit, request.source.path, request.source.dataBegin);
-        spTranslationUnit_addPreprocessorDefine(slangRequest, computeTranslationUnit, "__GLSL_COMPUTE__", "1");
-        computeEntryPointName = "main";
-    }
-    else
+    int translationUnitIndex = 0;
     {
-        int translationUnit = spAddTranslationUnit(slangRequest, sourceLanguage, nullptr);
-        spAddTranslationUnitSourceString(slangRequest, translationUnit, request.source.path, request.source.dataBegin);
-
-        vertexTranslationUnit = translationUnit;
-        fragmentTranslationUnit = translationUnit;
-        computeTranslationUnit = translationUnit;
+        translationUnitIndex = spAddTranslationUnit(slangRequest, sourceLanguage, nullptr);
+        spAddTranslationUnitSourceString(slangRequest, translationUnitIndex, request.source.path, request.source.dataBegin);
     }
 
     const int globalSpecializationArgCount = int(request.globalSpecializationArgs.getCount());
@@ -137,105 +134,100 @@ static const char computeEntryPointName[] = "computeMain";
         }
     };
 
-    if (request.computeShader.name)
-    {
-        int computeEntryPointIndex = 0;
-        if(!gOptions.dontAddDefaultEntryPoints)
+   Index explicitEntryPointCount = request.entryPoints.getCount();
+   for(Index ee = 0; ee < explicitEntryPointCount; ++ee)
+   {
+        if(gOptions.dontAddDefaultEntryPoints)
         {
-            computeEntryPointIndex = spAddEntryPoint(slangRequest, computeTranslationUnit,
-                computeEntryPointName,
-                SLANG_STAGE_COMPUTE);
-
-            setEntryPointSpecializationArgs(computeEntryPointIndex);
+            // If default entry points are not to be added, then
+            // the `request.entryPoints` array should have been
+            // left empty.
+            //
+            SLANG_ASSERT(false);
         }
 
-        spSetLineDirectiveMode(slangRequest, SLANG_LINE_DIRECTIVE_MODE_NONE);
+       auto& entryPointInfo = request.entryPoints[ee];
+       int entryPointIndex = spAddEntryPoint(
+           slangRequest,
+           translationUnitIndex,
+           entryPointInfo.name,
+           entryPointInfo.slangStage);
+       SLANG_ASSERT(entryPointIndex == ee);
 
-        const SlangResult res = spCompile(slangRequest);
+       setEntryPointSpecializationArgs(entryPointIndex);
+   }
 
-        if (auto diagnostics = spGetDiagnosticOutput(slangRequest))
-        {
-            fprintf(stderr, "%s", diagnostics);
-        }
+    spSetLineDirectiveMode(slangRequest, SLANG_LINE_DIRECTIVE_MODE_NONE);
 
-        SLANG_RETURN_ON_FAIL(res);
+    const SlangResult res = spCompile(slangRequest);
 
-        // We are going to get the entry point count... lets check what we have
-        if (input.passThrough == SLANG_PASS_THROUGH_NONE)
-        {
-            auto reflection = spGetReflection(slangRequest);
-            // Get the amount of entry points in reflection
-            const int entryPointCount = int(spReflection_getEntryPointCount(reflection));
+    if (auto diagnostics = spGetDiagnosticOutput(slangRequest))
+    {
+        fprintf(stderr, "%s", diagnostics);
+    }
 
-            // Above code assumes there is an entry point
-            SLANG_ASSERT(entryPointCount && computeEntryPointIndex < entryPointCount);
+    SLANG_RETURN_ON_FAIL(res);
 
-            auto entryPoint = spReflection_getEntryPointByIndex(reflection, computeEntryPointIndex);
+    
+    List<ShaderCompileRequest::EntryPoint> actualEntryPoints;
+    if(input.passThrough == SLANG_PASS_THROUGH_NONE)
+    {
+        // In the case where pass-through compilation is not being used,
+        // we can use the Slang reflection information to discover what
+        // the entry points were, and then use those to drive the
+        // loading of code.
+        //
+        auto reflection = slang::ProgramLayout::get(slangRequest);
 
-            // Get the entry point name
-            const char* entryPointName = spReflectionEntryPoint_getName(entryPoint);
+        // Get the amount of entry points in reflection
+        Index entryPointCount = Index(reflection->getEntryPointCount());
 
-            SLANG_ASSERT(entryPointName);
-        }
+        // We must have at least one entry point (whether explicit or implicit)
+        SLANG_ASSERT(entryPointCount);
 
+        for(Index ee = 0; ee < entryPointCount; ++ee)
         {
-            size_t codeSize = 0;
-            char const* code = (char const*) spGetEntryPointCode(slangRequest, computeEntryPointIndex, &codeSize);
+            auto entryPoint = reflection->getEntryPointByIndex(ee);
+            const char* entryPointName = entryPoint->getName();
+            SLANG_ASSERT(entryPointName);
+
+            auto slangStage = entryPoint->getStage();
 
-            ShaderProgram::KernelDesc kernelDesc;
-            kernelDesc.stage = StageType::Compute;
-            kernelDesc.codeBegin = code;
-            kernelDesc.codeEnd = code + codeSize;
+            ShaderCompileRequest::EntryPoint entryPointInfo;
+            entryPointInfo.name = entryPointName;
+            entryPointInfo.slangStage = slangStage;
 
-            out.set(PipelineType::Compute, &kernelDesc, 1);
+            actualEntryPoints.add(entryPointInfo);
         }
     }
     else
     {
-        int vertexEntryPoint = 0;
-        int fragmentEntryPoint = 1;
-        if( !gOptions.dontAddDefaultEntryPoints )
-        {
-            vertexEntryPoint = spAddEntryPoint(slangRequest, vertexTranslationUnit, vertexEntryPointName, SLANG_STAGE_VERTEX);
-            fragmentEntryPoint = spAddEntryPoint(slangRequest, fragmentTranslationUnit, fragmentEntryPointName, SLANG_STAGE_FRAGMENT);
-
-            setEntryPointSpecializationArgs(vertexEntryPoint);
-            setEntryPointSpecializationArgs(fragmentEntryPoint);
-        }
-
-        const SlangResult res = spCompile(slangRequest);
-        if (auto diagnostics = spGetDiagnosticOutput(slangRequest))
-        {
-            // TODO(tfoley): re-enable when I get a logging solution in place
-//            OutputDebugStringA(diagnostics);
-            fprintf(stderr, "%s", diagnostics);
-        }
-
-        SLANG_RETURN_ON_FAIL(res);
-
-        {
-            size_t vertexCodeSize = 0;
-            char const* vertexCode = (char const*) spGetEntryPointCode(slangRequest, vertexEntryPoint, &vertexCodeSize);
+        actualEntryPoints = request.entryPoints;
+    }
 
-            size_t fragmentCodeSize = 0;
-            char const* fragmentCode = (char const*) spGetEntryPointCode(slangRequest, fragmentEntryPoint, &fragmentCodeSize);
+    List<ShaderProgram::KernelDesc> kernelDescs;
 
-            static const int kDescCount = 2;
+    Index actualEntryPointCount = actualEntryPoints.getCount();
+    for(Index ee = 0; ee < actualEntryPointCount; ++ee)
+    {
+        auto& actualEntryPoint = actualEntryPoints[ee];
 
-            ShaderProgram::KernelDesc kernelDescs[kDescCount];
+        size_t codeSize = 0;
+        char const* code = (char const*) spGetEntryPointCode(slangRequest, int(ee), &codeSize);
 
-            kernelDescs[0].stage = StageType::Vertex;
-            kernelDescs[0].codeBegin = vertexCode;
-            kernelDescs[0].codeEnd = vertexCode + vertexCodeSize;
+        auto gfxStage = _translateStage(actualEntryPoint.slangStage);
 
-            kernelDescs[1].stage = StageType::Fragment;
-            kernelDescs[1].codeBegin = fragmentCode;
-            kernelDescs[1].codeEnd = fragmentCode + fragmentCodeSize;
+        ShaderProgram::KernelDesc kernelDesc;
+        kernelDesc.stage = gfxStage;
+        kernelDesc.codeBegin = code;
+        kernelDesc.codeEnd = code + codeSize;
+        kernelDesc.entryPointName = actualEntryPoint.name;
 
-            out.set(PipelineType::Graphics, kernelDescs, kDescCount);
-        }
+        kernelDescs.add(kernelDesc);
     }
 
+    out.set(input.pipelineType, kernelDescs.getBuffer(), kernelDescs.getCount());
+
     return SLANG_OK;
 }
 
@@ -260,8 +252,12 @@ static const char computeEntryPointName[] = "computeMain";
     return SLANG_OK;
 }
 
-/* static */SlangResult ShaderCompilerUtil::compileWithLayout(SlangSession* session, const String& sourcePath, const Slang::List<Slang::CommandLine::Arg>& compileArgs, Options::ShaderProgramType shaderType, const ShaderCompilerUtil::Input& input, OutputAndLayout& output)
+/* static */SlangResult ShaderCompilerUtil::compileWithLayout(SlangSession* session, const Options& options, const ShaderCompilerUtil::Input& input, OutputAndLayout& output)
 {
+    String sourcePath = options.sourcePath;
+    auto& compileArgs = options.compileArgs;
+    auto shaderType = options.shaderType;
+
     List<char> sourceText;
     SLANG_RETURN_ON_FAIL(readSource(sourcePath, sourceText));
 
@@ -294,6 +290,7 @@ static const char computeEntryPointName[] = "computeMain";
             break;
 
         case Options::ShaderProgramType::Compute:
+        case Options::ShaderProgramType::RayTracing:
             layout.numRenderTargets = 0;
             break;
     }
@@ -317,17 +314,47 @@ static const char computeEntryPointName[] = "computeMain";
     compileRequest.compileArgs = compileArgs;
 
     compileRequest.source = sourceInfo;
-    if (shaderType == Options::ShaderProgramType::Graphics || shaderType == Options::ShaderProgramType::GraphicsCompute)
-    {
-        compileRequest.vertexShader.source = sourceInfo;
-        compileRequest.vertexShader.name = vertexEntryPointName;
-        compileRequest.fragmentShader.source = sourceInfo;
-        compileRequest.fragmentShader.name = fragmentEntryPointName;
-    }
-    else
+
+    // Now we will add the "default" entry point names/stages that
+    // are appropriate to the pipeline type being targetted, *unless*
+    // the options specify that we should leave out the default
+    // entry points and instead rely on the Slang compiler's built-in
+    // mechanisms for discovering entry points (e.g., `[shader(...)]`
+    // attributes).
+    //
+    if( !options.dontAddDefaultEntryPoints )
     {
-        compileRequest.computeShader.source = sourceInfo;
-        compileRequest.computeShader.name = computeEntryPointName;
+        if (shaderType == Options::ShaderProgramType::Graphics || shaderType == Options::ShaderProgramType::GraphicsCompute)
+        {
+            ShaderCompileRequest::EntryPoint vertexEntryPoint;
+            vertexEntryPoint.name = vertexEntryPointName;
+            vertexEntryPoint.slangStage = SLANG_STAGE_VERTEX;
+            compileRequest.entryPoints.add(vertexEntryPoint);
+
+            ShaderCompileRequest::EntryPoint fragmentEntryPoint;
+            fragmentEntryPoint.name = fragmentEntryPointName;
+            fragmentEntryPoint.slangStage = SLANG_STAGE_FRAGMENT;
+            compileRequest.entryPoints.add(fragmentEntryPoint);
+        }
+        else if( shaderType == Options::ShaderProgramType::RayTracing )
+        {
+            // Note: Current GPU ray tracing pipelines allow for an
+            // almost arbitrary mix of entry points for different stages
+            // to be used together (e.g., a single "program" might
+            // have multiple any-hit shaders, multiple miss shaders, etc.)
+            //
+            // Rather than try to define a fixed set of entry point
+            // names and stages that the testing will support, we will
+            // instead rely on `[shader(...)]` annotations to tell us
+            // what entry points are present in the input code.
+        }
+        else
+        {
+            ShaderCompileRequest::EntryPoint computeEntryPoint;
+            computeEntryPoint.name = computeEntryPointName;
+            computeEntryPoint.slangStage = SLANG_STAGE_COMPUTE;
+            compileRequest.entryPoints.add(computeEntryPoint);
+        }
     }
     compileRequest.globalSpecializationArgs = layout.globalSpecializationArgs;
     compileRequest.entryPointSpecializationArgs = layout.entryPointSpecializationArgs;
diff --git a/tools/render-test/slang-support.h b/tools/render-test/slang-support.h
index 97b85ff8f..99509914e 100644
--- a/tools/render-test/slang-support.h
+++ b/tools/render-test/slang-support.h
@@ -17,6 +17,7 @@ struct ShaderCompilerUtil
         SlangCompileTarget      target;
         SlangSourceLanguage     sourceLanguage;
         SlangPassThrough        passThrough;
+        PipelineType            pipelineType = PipelineType::Unknown;
         char const*             profile;
         const char**            args;
         int                     argCount;
@@ -24,7 +25,7 @@ struct ShaderCompilerUtil
 
     struct Output
     {
-        void set(PipelineType pipelineType, const ShaderProgram::KernelDesc* inKernelDescs, int kernelDescCount)
+        void set(PipelineType pipelineType, const ShaderProgram::KernelDesc* inKernelDescs, Slang::Index kernelDescCount)
         {
             kernelDescs.clear();
             kernelDescs.addRange(inKernelDescs, kernelDescCount);
@@ -82,7 +83,7 @@ struct ShaderCompilerUtil
         Slang::String sourcePath;
     };
 
-    static SlangResult compileWithLayout(SlangSession* session, const Slang::String& sourcePath, const Slang::List<Slang::CommandLine::Arg>& compileArgs, Options::ShaderProgramType shaderType, const ShaderCompilerUtil::Input& input, OutputAndLayout& output);
+    static SlangResult compileWithLayout(SlangSession* session, const Options& options, const ShaderCompilerUtil::Input& input, OutputAndLayout& output);
 
     static SlangResult readSource(const Slang::String& inSourcePath, List<char>& outSourceText);