summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--examples/gpu-printing/gpu-printing.vcxproj8
-rw-r--r--examples/gpu-printing/gpu-printing.vcxproj.filters25
-rw-r--r--prelude/slang-cuda-prelude.h4
-rw-r--r--premake5.lua33
-rw-r--r--slang.sln172
-rw-r--r--source/core/slang-downstream-compiler.h9
-rw-r--r--source/core/slang-nvrtc-compiler.cpp57
-rw-r--r--source/slang/hlsl.meta.slang1
-rw-r--r--source/slang/slang-compiler.cpp50
-rw-r--r--source/slang/slang-emit-cuda.cpp60
-rw-r--r--source/slang/slang.vcxproj13
-rw-r--r--source/slang/slang.vcxproj.filters45
-rw-r--r--tools/gfx/render.h15
-rw-r--r--tools/render-test/cuda/cuda-compute-util.cpp273
-rw-r--r--tools/render-test/options.cpp4
-rw-r--r--tools/render-test/options.h3
-rw-r--r--tools/render-test/render-test-main.cpp25
-rw-r--r--tools/render-test/slang-support.cpp267
-rw-r--r--tools/render-test/slang-support.h5
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);
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-01 10:05:45 +0000