diff options
| author | Yong He <yonghe@outlook.com> | 2021-06-30 14:59:18 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2021-06-30 14:59:18 -0700 |
| commit | a03d21a5f54cba913c3f52e2822a433de8f39fdd (patch) | |
| tree | 84d24d4355cc4b3e941da9eab57147cd9d297ee4 | |
| parent | 5395ef82535c283109b1ea6b89b737c5a39bf147 (diff) | |
[gfx] Add inline ray tracing support. (#1899)
29 files changed, 2384 insertions, 335 deletions
diff --git a/build/visual-studio/ray-tracing/ray-tracing.vcxproj b/build/visual-studio/ray-tracing/ray-tracing.vcxproj new file mode 100644 index 000000000..cc64fcd8c --- /dev/null +++ b/build/visual-studio/ray-tracing/ray-tracing.vcxproj @@ -0,0 +1,193 @@ +<?xml version="1.0" encoding="utf-8"?> +<Project DefaultTargets="Build" ToolsVersion="15.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003"> + <ItemGroup Label="ProjectConfigurations"> + <ProjectConfiguration Include="Debug|Win32"> + <Configuration>Debug</Configuration> + <Platform>Win32</Platform> + </ProjectConfiguration> + <ProjectConfiguration Include="Debug|x64"> + <Configuration>Debug</Configuration> + <Platform>x64</Platform> + </ProjectConfiguration> + <ProjectConfiguration Include="Release|Win32"> + <Configuration>Release</Configuration> + <Platform>Win32</Platform> + </ProjectConfiguration> + <ProjectConfiguration Include="Release|x64"> + <Configuration>Release</Configuration> + <Platform>x64</Platform> + </ProjectConfiguration> + </ItemGroup> + <PropertyGroup Label="Globals"> + <ProjectGuid>{71AC0F50-5DFD-FA91-8661-E95372118EFB}</ProjectGuid> + <IgnoreWarnCompileDuplicatedFilename>true</IgnoreWarnCompileDuplicatedFilename> + <Keyword>Win32Proj</Keyword> + <RootNamespace>ray-tracing</RootNamespace> + </PropertyGroup> + <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" /> + <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'" Label="Configuration"> + <ConfigurationType>Application</ConfigurationType> + <UseDebugLibraries>true</UseDebugLibraries> + <CharacterSet>Unicode</CharacterSet> + <PlatformToolset>v141</PlatformToolset> + </PropertyGroup> + <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'" Label="Configuration"> + <ConfigurationType>Application</ConfigurationType> + <UseDebugLibraries>true</UseDebugLibraries> + <CharacterSet>Unicode</CharacterSet> + <PlatformToolset>v141</PlatformToolset> + </PropertyGroup> + <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'" Label="Configuration"> + <ConfigurationType>Application</ConfigurationType> + <UseDebugLibraries>false</UseDebugLibraries> + <CharacterSet>Unicode</CharacterSet> + <PlatformToolset>v141</PlatformToolset> + </PropertyGroup> + <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" Label="Configuration"> + <ConfigurationType>Application</ConfigurationType> + <UseDebugLibraries>false</UseDebugLibraries> + <CharacterSet>Unicode</CharacterSet> + <PlatformToolset>v141</PlatformToolset> + </PropertyGroup> + <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" /> + <ImportGroup Label="ExtensionSettings"> + </ImportGroup> + <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'"> + <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" /> + </ImportGroup> + <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Debug|x64'"> + <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" /> + </ImportGroup> + <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Release|Win32'"> + <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" /> + </ImportGroup> + <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Release|x64'"> + <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" /> + </ImportGroup> + <PropertyGroup Label="UserMacros" /> + <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'"> + <LinkIncremental>true</LinkIncremental> + <OutDir>..\..\..\bin\windows-x86\debug\</OutDir> + <IntDir>..\..\..\intermediate\windows-x86\debug\ray-tracing\</IntDir> + <TargetName>ray-tracing</TargetName> + <TargetExt>.exe</TargetExt> + </PropertyGroup> + <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'"> + <LinkIncremental>true</LinkIncremental> + <OutDir>..\..\..\bin\windows-x64\debug\</OutDir> + <IntDir>..\..\..\intermediate\windows-x64\debug\ray-tracing\</IntDir> + <TargetName>ray-tracing</TargetName> + <TargetExt>.exe</TargetExt> + </PropertyGroup> + <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'"> + <LinkIncremental>false</LinkIncremental> + <OutDir>..\..\..\bin\windows-x86\release\</OutDir> + <IntDir>..\..\..\intermediate\windows-x86\release\ray-tracing\</IntDir> + <TargetName>ray-tracing</TargetName> + <TargetExt>.exe</TargetExt> + </PropertyGroup> + <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'"> + <LinkIncremental>false</LinkIncremental> + <OutDir>..\..\..\bin\windows-x64\release\</OutDir> + <IntDir>..\..\..\intermediate\windows-x64\release\ray-tracing\</IntDir> + <TargetName>ray-tracing</TargetName> + <TargetExt>.exe</TargetExt> + </PropertyGroup> + <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'"> + <ClCompile> + <PrecompiledHeader>NotUsing</PrecompiledHeader> + <WarningLevel>Level3</WarningLevel> + <PreprocessorDefinitions>_DEBUG;%(PreprocessorDefinitions)</PreprocessorDefinitions> + <AdditionalIncludeDirectories>..\..\..;..\..\..\tools;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories> + <DebugInformationFormat>EditAndContinue</DebugInformationFormat> + <Optimization>Disabled</Optimization> + <RuntimeLibrary>MultiThreadedDebug</RuntimeLibrary> + </ClCompile> + <Link> + <SubSystem>Windows</SubSystem> + <GenerateDebugInformation>true</GenerateDebugInformation> + </Link> + </ItemDefinitionGroup> + <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'"> + <ClCompile> + <PrecompiledHeader>NotUsing</PrecompiledHeader> + <WarningLevel>Level3</WarningLevel> + <PreprocessorDefinitions>_DEBUG;%(PreprocessorDefinitions)</PreprocessorDefinitions> + <AdditionalIncludeDirectories>..\..\..;..\..\..\tools;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories> + <DebugInformationFormat>EditAndContinue</DebugInformationFormat> + <Optimization>Disabled</Optimization> + <RuntimeLibrary>MultiThreadedDebug</RuntimeLibrary> + </ClCompile> + <Link> + <SubSystem>Windows</SubSystem> + <GenerateDebugInformation>true</GenerateDebugInformation> + </Link> + </ItemDefinitionGroup> + <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'"> + <ClCompile> + <PrecompiledHeader>NotUsing</PrecompiledHeader> + <WarningLevel>Level3</WarningLevel> + <PreprocessorDefinitions>NDEBUG;%(PreprocessorDefinitions)</PreprocessorDefinitions> + <AdditionalIncludeDirectories>..\..\..;..\..\..\tools;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories> + <Optimization>Full</Optimization> + <FunctionLevelLinking>true</FunctionLevelLinking> + <IntrinsicFunctions>true</IntrinsicFunctions> + <MinimalRebuild>false</MinimalRebuild> + <StringPooling>true</StringPooling> + <RuntimeLibrary>MultiThreaded</RuntimeLibrary> + </ClCompile> + <Link> + <SubSystem>Windows</SubSystem> + <EnableCOMDATFolding>true</EnableCOMDATFolding> + <OptimizeReferences>true</OptimizeReferences> + </Link> + </ItemDefinitionGroup> + <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'"> + <ClCompile> + <PrecompiledHeader>NotUsing</PrecompiledHeader> + <WarningLevel>Level3</WarningLevel> + <PreprocessorDefinitions>NDEBUG;%(PreprocessorDefinitions)</PreprocessorDefinitions> + <AdditionalIncludeDirectories>..\..\..;..\..\..\tools;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories> + <Optimization>Full</Optimization> + <FunctionLevelLinking>true</FunctionLevelLinking> + <IntrinsicFunctions>true</IntrinsicFunctions> + <MinimalRebuild>false</MinimalRebuild> + <StringPooling>true</StringPooling> + <RuntimeLibrary>MultiThreaded</RuntimeLibrary> + </ClCompile> + <Link> + <SubSystem>Windows</SubSystem> + <EnableCOMDATFolding>true</EnableCOMDATFolding> + <OptimizeReferences>true</OptimizeReferences> + </Link> + </ItemDefinitionGroup> + <ItemGroup> + <ClCompile Include="..\..\..\examples\ray-tracing\main.cpp" /> + </ItemGroup> + <ItemGroup> + <None Include="..\..\..\examples\ray-tracing\shaders.slang" /> + </ItemGroup> + <ItemGroup> + <ProjectReference Include="..\example-base\example-base.vcxproj"> + <Project>{37BED5B5-23FA-D81F-8C0C-F1167867813A}</Project> + </ProjectReference> + <ProjectReference Include="..\slang\slang.vcxproj"> + <Project>{DB00DA62-0533-4AFD-B59F-A67D5B3A0808}</Project> + </ProjectReference> + <ProjectReference Include="..\gfx\gfx.vcxproj"> + <Project>{222F7498-B40C-4F3F-A704-DDEB91A4484A}</Project> + </ProjectReference> + <ProjectReference Include="..\gfx-util\gfx-util.vcxproj"> + <Project>{F5ADB74E-02A7-44FB-AA3B-FC02F8AC7A4B}</Project> + </ProjectReference> + <ProjectReference Include="..\platform\platform.vcxproj"> + <Project>{3565FE5E-4FA3-11EB-AE93-0242AC130002}</Project> + </ProjectReference> + <ProjectReference Include="..\core\core.vcxproj"> + <Project>{F9BE7957-8399-899E-0C49-E714FDDD4B65}</Project> + </ProjectReference> + </ItemGroup> + <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" /> + <ImportGroup Label="ExtensionTargets"> + </ImportGroup> +</Project>
\ No newline at end of file diff --git a/build/visual-studio/ray-tracing/ray-tracing.vcxproj.filters b/build/visual-studio/ray-tracing/ray-tracing.vcxproj.filters new file mode 100644 index 000000000..007b0b2e1 --- /dev/null +++ b/build/visual-studio/ray-tracing/ray-tracing.vcxproj.filters @@ -0,0 +1,18 @@ +<?xml version="1.0" encoding="utf-8"?> +<Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003"> + <ItemGroup> + <Filter Include="Source Files"> + <UniqueIdentifier>{E9C7FDCE-D52A-8D73-7EB0-C5296AF258F6}</UniqueIdentifier> + </Filter> + </ItemGroup> + <ItemGroup> + <ClCompile Include="..\..\..\examples\ray-tracing\main.cpp"> + <Filter>Source Files</Filter> + </ClCompile> + </ItemGroup> + <ItemGroup> + <None Include="..\..\..\examples\ray-tracing\shaders.slang"> + <Filter>Source Files</Filter> + </None> + </ItemGroup> +</Project>
\ No newline at end of file diff --git a/examples/example-base/example-base.cpp b/examples/example-base/example-base.cpp index 04f938697..c45e3cc27 100644 --- a/examples/example-base/example-base.cpp +++ b/examples/example-base/example-base.cpp @@ -9,7 +9,11 @@ using namespace Slang; using namespace gfx; -Slang::Result WindowedAppBase::initializeBase(const char* title, int width, int height) +Slang::Result WindowedAppBase::initializeBase( + const char* title, + int width, + int height, + DeviceType deviceType) { // Create a window for our application to render into. // @@ -30,7 +34,7 @@ Slang::Result WindowedAppBase::initializeBase(const char* title, int width, int gfxEnableDebugLayer(); #endif IDevice::Desc deviceDesc = {}; - // deviceDesc.slang.targetFlags = SLANG_TARGET_FLAG_DUMP_IR; + deviceDesc.deviceType = deviceType; gfx::Result res = gfxCreateDevice(&deviceDesc, gDevice.writeRef()); if (SLANG_FAILED(res)) return res; diff --git a/examples/example-base/example-base.h b/examples/example-base/example-base.h index a1c46aa53..b898f7417 100644 --- a/examples/example-base/example-base.h +++ b/examples/example-base/example-base.h @@ -26,7 +26,11 @@ protected: Slang::ComPtr<gfx::IRenderPassLayout> gRenderPass; Slang::ComPtr<gfx::ICommandQueue> gQueue; - Slang::Result initializeBase(const char* titile, int width, int height); + Slang::Result initializeBase( + const char* titile, + int width, + int height, + gfx::DeviceType deviceType = gfx::DeviceType::Default); void createSwapchainFramebuffers(); void mainLoop(); diff --git a/examples/gpu-printing/main.cpp b/examples/gpu-printing/main.cpp index 11a1553a3..03ab5b51f 100644 --- a/examples/gpu-printing/main.cpp +++ b/examples/gpu-printing/main.cpp @@ -109,7 +109,7 @@ Result execute() printBufferDesc.defaultState = ResourceState::UnorderedAccess; printBufferDesc.allowedStates = ResourceStateSet( ResourceState::CopySource, ResourceState::CopyDestination, ResourceState::UnorderedAccess); - printBufferDesc.cpuAccessFlags = IResource::AccessFlag::Read; // | Resource::AccessFlag::Write; + printBufferDesc.cpuAccessFlags = AccessFlag::Read; // | Resource::AccessFlag::Write; auto printBuffer = gDevice->createBufferResource(printBufferDesc); IResourceView::Desc printBufferViewDesc; diff --git a/examples/ray-tracing/README.md b/examples/ray-tracing/README.md new file mode 100644 index 000000000..fd7cc1cec --- /dev/null +++ b/examples/ray-tracing/README.md @@ -0,0 +1,9 @@ +Slang "Ray Tracing" Example +=========================== + +The goal of this example is to demonstrate how to use hardware ray-tracing in Slang. + +The `shaders.slang` file contains a compute shader that traces primary rays from camera and shade intersections with basic lighting + ray-traced shadows. The file also defines a vertex and a fragment shader entry point for displaying the ray-traced image produced by the compute shader. + +The `main.cpp` file contains the C++ application code, showing how to use the Slang API to load and compile the shader code, and how to use a graphics API abstraction layer implemented in `tools/gfx` to initiate hardware ray-tracing. +Note that this abstraction layer is *not* required in order to work with Slang, and it is just there to help us write example and test applications more conveniently.
\ No newline at end of file diff --git a/examples/ray-tracing/main.cpp b/examples/ray-tracing/main.cpp new file mode 100644 index 000000000..a5093ad18 --- /dev/null +++ b/examples/ray-tracing/main.cpp @@ -0,0 +1,645 @@ +// main.cpp + +// This file implements an example of hardware ray-tracing using +// Slang shaders and the `gfx` graphics API. + +#include <slang.h> +#include "slang-gfx.h" +#include "gfx-util/shader-cursor.h" +#include "tools/platform/window.h" +#include "tools/platform/vector-math.h" +#include "slang-com-ptr.h" +#include "source/core/slang-basic.h" +#include "examples/example-base/example-base.h" + +using namespace gfx; +using namespace Slang; + +struct Uniforms +{ + float screenWidth, screenHeight; + float focalLength = 24.0f, frameHeight = 24.0f; + float cameraDir[4]; + float cameraUp[4]; + float cameraRight[4]; + float cameraPosition[4]; + float lightDir[4]; +}; + +struct Vertex +{ + float position[3]; +}; + +// Define geometry data for our test scene. +// The scene contains a floor plane, and a cube placed on top of it at the center. +static const int kVertexCount = 24; +static const Vertex kVertexData[kVertexCount] = +{ + // Floor plane + {{-100.0f, 0, 100.0f}}, + {{100.0f, 0, 100.0f}}, + {{100.0f, 0, -100.0f}}, + {{-100.0f, 0, -100.0f}}, + // Cube face (+y). + {{-1.0f, 2.0, 1.0f}}, + {{1.0f, 2.0, 1.0f}}, + {{1.0f, 2.0, -1.0f}}, + {{-1.0f, 2.0, -1.0f}}, + // Cube face (+z). + {{-1.0f, 0.0, 1.0f}}, + {{1.0f, 0.0, 1.0f}}, + {{1.0f, 2.0, 1.0f}}, + {{-1.0f, 2.0, 1.0f}}, + // Cube face (-z). + {{-1.0f, 0.0, -1.0f}}, + {{-1.0f, 2.0, -1.0f}}, + {{1.0f, 2.0, -1.0f}}, + {{1.0f, 0.0, -1.0f}}, + // Cube face (-x). + {{-1.0f, 0.0, -1.0f}}, + {{-1.0f, 0.0, 1.0f}}, + {{-1.0f, 2.0, 1.0f}}, + {{-1.0f, 2.0, -1.0f}}, + // Cube face (+x). + {{1.0f, 2.0, -1.0f}}, + {{1.0f, 2.0, 1.0f}}, + {{1.0f, 0.0, 1.0f}}, + {{1.0f, 0.0, -1.0f}}, +}; +static const int kIndexCount = 36; +static const int kIndexData[kIndexCount] = +{ + 0, 1, 2, 0, 2, 3, + 4, 5, 6, 4, 6, 7, + 8, 9, 10, 8, 10, 11, + 12, 13, 14, 12, 14, 15, + 16, 17, 18, 16, 18, 19, + 20, 21, 22, 20, 22, 23 +}; + +struct Primitive +{ + float data[4]; + float color[4]; +}; +static const int kPrimitiveCount = 12; +static const Primitive kPrimitiveData[kPrimitiveCount] = +{ + {{0.0f, 1.0f, 0.0f, 0.0f}, {0.75f, 0.8f, 0.85f, 1.0f}}, + {{0.0f, 1.0f, 0.0f, 0.0f}, {0.75f, 0.8f, 0.85f, 1.0f}}, + {{0.0f, 1.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}}, + {{0.0f, 1.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}}, + {{0.0f, 0.0f, 1.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}}, + {{0.0f, 0.0f, 1.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}}, + {{0.0f, 0.0f, -1.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}}, + {{0.0f, 0.0f, -1.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}}, + {{-1.0f, 0.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}}, + {{-1.0f, 0.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}}, + {{1.0f, 0.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}}, + {{1.0f, 0.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}}, +}; + + +// We need to use a rasterization pipeline to copy the ray-traced image +// to the swapchain. To do so we need to render a full-screen triangle. +// We will define a small helper type that defines the data for such a triangle. +// +struct FullScreenTriangle +{ + struct Vertex + { + float position[2]; + }; + + enum + { + kVertexCount = 3 + }; + + static const Vertex kVertices[kVertexCount]; +}; +const FullScreenTriangle::Vertex FullScreenTriangle::kVertices[FullScreenTriangle::kVertexCount] = { + {{-1, -1}}, + {{-1, 3}}, + {{3, -1}}, +}; + +// The example application will be implemented as a `struct`, so that +// we can scope the resources it allocates without using global variables. +// +struct RayTracing : public WindowedAppBase +{ + + +Uniforms gUniforms = {}; + + +// Many Slang API functions return detailed diagnostic information +// (error messages, warnings, etc.) as a "blob" of data, or return +// a null blob pointer instead if there were no issues. +// +// For convenience, we define a subroutine that will dump the information +// in a diagnostic blob if one is produced, and skip it otherwise. +// +void diagnoseIfNeeded(slang::IBlob* diagnosticsBlob) +{ + if( diagnosticsBlob != nullptr ) + { + printf("%s", (const char*) diagnosticsBlob->getBufferPointer()); +#ifdef _WIN32 + _Win32OutputDebugString((const char*)diagnosticsBlob->getBufferPointer()); +#endif + } +} + +// Load and compile shader code from souce. +gfx::Result loadShaderProgram( + gfx::IDevice* device, + gfx::PipelineType pipelineType, + gfx::IShaderProgram** outProgram) +{ + ComPtr<slang::ISession> slangSession; + slangSession = device->getSlangSession(); + + ComPtr<slang::IBlob> diagnosticsBlob; + slang::IModule* module = slangSession->loadModule("shaders", diagnosticsBlob.writeRef()); + diagnoseIfNeeded(diagnosticsBlob); + if(!module) + return SLANG_FAIL; + + Slang::List<slang::IComponentType*> componentTypes; + componentTypes.add(module); + if (pipelineType == PipelineType::Compute) + { + ComPtr<slang::IEntryPoint> computeEntryPoint; + SLANG_RETURN_ON_FAIL(module->findEntryPointByName("computeMain", computeEntryPoint.writeRef())); + componentTypes.add(computeEntryPoint); + } + else + { + ComPtr<slang::IEntryPoint> entryPoint; + SLANG_RETURN_ON_FAIL(module->findEntryPointByName("vertexMain", entryPoint.writeRef())); + componentTypes.add(entryPoint); + SLANG_RETURN_ON_FAIL(module->findEntryPointByName("fragmentMain", entryPoint.writeRef())); + componentTypes.add(entryPoint); + } + + ComPtr<slang::IComponentType> linkedProgram; + SlangResult result = slangSession->createCompositeComponentType( + componentTypes.getBuffer(), + componentTypes.getCount(), + linkedProgram.writeRef(), + diagnosticsBlob.writeRef()); + diagnoseIfNeeded(diagnosticsBlob); + SLANG_RETURN_ON_FAIL(result); + + gfx::IShaderProgram::Desc programDesc = {}; + programDesc.pipelineType = pipelineType; + programDesc.slangProgram = linkedProgram; + SLANG_RETURN_ON_FAIL(device->createProgram(programDesc, outProgram)); + + return SLANG_OK; +} + +ComPtr<gfx::IPipelineState> gPresentPipelineState; +ComPtr<gfx::IPipelineState> gRenderPipelineState; +ComPtr<gfx::IBufferResource> gFullScreenVertexBuffer; +ComPtr<gfx::IBufferResource> gVertexBuffer; +ComPtr<gfx::IBufferResource> gIndexBuffer; +ComPtr<gfx::IBufferResource> gPrimitiveBuffer; +ComPtr<gfx::IBufferResource> gTransformBuffer; +ComPtr<gfx::IResourceView> gPrimitiveBufferSRV; +ComPtr<gfx::IBufferResource> gInstanceBuffer; +ComPtr<gfx::IBufferResource> gBLASBuffer; +ComPtr<gfx::IAccelerationStructure> gBLAS; +ComPtr<gfx::IBufferResource> gTLASBuffer; +ComPtr<gfx::IAccelerationStructure> gTLAS; +ComPtr<gfx::ITextureResource> gResultTexture; +ComPtr<gfx::IResourceView> gResultTextureUAV; + +uint64_t lastTime = 0; + +// glm::vec3 lightDir = normalize(glm::vec3(10, 10, 10)); +// glm::vec3 lightColor = glm::vec3(1, 1, 1); + +glm::vec3 cameraPosition = glm::vec3(-2.53f, 2.72f, 4.3f); +float cameraOrientationAngles[2] = {-0.475f, -0.35f}; // Spherical angles (theta, phi). + +float translationScale = 0.5f; +float rotationScale = 0.01f; + +// In order to control camera movement, we will +// use good old WASD +bool wPressed = false; +bool aPressed = false; +bool sPressed = false; +bool dPressed = false; + +bool isMouseDown = false; +float lastMouseX = 0.0f; +float lastMouseY = 0.0f; + +void setKeyState(platform::KeyCode key, bool state) +{ + switch (key) + { + default: + break; + case platform::KeyCode::W: + wPressed = state; + break; + case platform::KeyCode::A: + aPressed = state; + break; + case platform::KeyCode::S: + sPressed = state; + break; + case platform::KeyCode::D: + dPressed = state; + break; + } +} +void onKeyDown(platform::KeyEventArgs args) { setKeyState(args.key, true); } +void onKeyUp(platform::KeyEventArgs args) { setKeyState(args.key, false); } + +void onMouseDown(platform::MouseEventArgs args) +{ + isMouseDown = true; + lastMouseX = (float)args.x; + lastMouseY = (float)args.y; +} + +void onMouseMove(platform::MouseEventArgs args) +{ + if (isMouseDown) + { + float deltaX = args.x - lastMouseX; + float deltaY = args.y - lastMouseY; + + cameraOrientationAngles[0] += -deltaX * rotationScale; + cameraOrientationAngles[1] += -deltaY * rotationScale; + lastMouseX = (float)args.x; + lastMouseY = (float)args.y; + } +} +void onMouseUp(platform::MouseEventArgs args) { isMouseDown = false; } + +Slang::Result initialize() +{ + initializeBase("Ray Tracing", 1024, 768, DeviceType::Vulkan); + gWindow->events.mouseMove = [this](const platform::MouseEventArgs& e) { onMouseMove(e); }; + gWindow->events.mouseUp = [this](const platform::MouseEventArgs& e) { onMouseUp(e); }; + gWindow->events.mouseDown = [this](const platform::MouseEventArgs& e) { onMouseDown(e); }; + gWindow->events.keyDown = [this](const platform::KeyEventArgs& e) { onKeyDown(e); }; + gWindow->events.keyUp = [this](const platform::KeyEventArgs& e) { onKeyUp(e); }; + + IBufferResource::Desc vertexBufferDesc; + vertexBufferDesc.type = IResource::Type::Buffer; + vertexBufferDesc.sizeInBytes = kVertexCount * sizeof(Vertex); + vertexBufferDesc.defaultState = ResourceState::UnorderedAccess; + gVertexBuffer = gDevice->createBufferResource(vertexBufferDesc, &kVertexData[0]); + if(!gVertexBuffer) return SLANG_FAIL; + + IBufferResource::Desc indexBufferDesc; + indexBufferDesc.type = IResource::Type::Buffer; + indexBufferDesc.sizeInBytes = kIndexCount * sizeof(int32_t); + indexBufferDesc.defaultState = ResourceState::UnorderedAccess; + gIndexBuffer = gDevice->createBufferResource(indexBufferDesc, &kIndexData[0]); + if (!gIndexBuffer) + return SLANG_FAIL; + + IBufferResource::Desc primitiveBufferDesc; + primitiveBufferDesc.type = IResource::Type::Buffer; + primitiveBufferDesc.sizeInBytes = kPrimitiveCount * sizeof(Primitive); + primitiveBufferDesc.defaultState = ResourceState::ShaderResource; + gPrimitiveBuffer = gDevice->createBufferResource(primitiveBufferDesc, &kPrimitiveData[0]); + if (!gPrimitiveBuffer) + return SLANG_FAIL; + + IResourceView::Desc primitiveSRVDesc = {}; + primitiveSRVDesc.format = Format::Unknown; + primitiveSRVDesc.type = IResourceView::Type::ShaderResource; + gPrimitiveBufferSRV = gDevice->createBufferView(gPrimitiveBuffer, primitiveSRVDesc); + + IBufferResource::Desc transformBufferDesc; + transformBufferDesc.type = IResource::Type::Buffer; + transformBufferDesc.sizeInBytes = sizeof(float) * 16; + transformBufferDesc.defaultState = ResourceState::UnorderedAccess; + float transformData[12] = { + 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f}; + gTransformBuffer = gDevice->createBufferResource(transformBufferDesc, &transformData); + if (!gTransformBuffer) + return SLANG_FAIL; + // Build bottom level acceleration structure. + { + IAccelerationStructure::BuildInputs accelerationStructureBuildInputs; + IAccelerationStructure::PrebuildInfo accelerationStructurePrebuildInfo; + accelerationStructureBuildInputs.descCount = 1; + accelerationStructureBuildInputs.kind = IAccelerationStructure::Kind::BottomLevel; + accelerationStructureBuildInputs.flags = + IAccelerationStructure::BuildFlags::AllowCompaction; + IAccelerationStructure::GeometryDesc geomDesc; + geomDesc.flags = IAccelerationStructure::GeometryFlags::Opaque; + geomDesc.type = IAccelerationStructure::GeometryType::Triangles; + geomDesc.content.triangles.indexCount = kIndexCount; + geomDesc.content.triangles.indexData = gIndexBuffer->getDeviceAddress(); + geomDesc.content.triangles.indexFormat = Format::R_UInt32; + geomDesc.content.triangles.vertexCount = kVertexCount; + geomDesc.content.triangles.vertexData = gVertexBuffer->getDeviceAddress(); + geomDesc.content.triangles.vertexFormat = Format::RGB_Float32; + geomDesc.content.triangles.vertexStride = sizeof(Vertex); + geomDesc.content.triangles.transform3x4 = gTransformBuffer->getDeviceAddress(); + accelerationStructureBuildInputs.geometryDescs = &geomDesc; + + // Query buffer size for acceleration structure build. + SLANG_RETURN_ON_FAIL(gDevice->getAccelerationStructurePrebuildInfo( + accelerationStructureBuildInputs, &accelerationStructurePrebuildInfo)); + // Allocate buffers for acceleration structure. + IBufferResource::Desc asDraftBufferDesc; + asDraftBufferDesc.type = IResource::Type::Buffer; + asDraftBufferDesc.defaultState = ResourceState::AccelerationStructure; + asDraftBufferDesc.sizeInBytes = accelerationStructurePrebuildInfo.resultDataMaxSize; + ComPtr<IBufferResource> draftBuffer = gDevice->createBufferResource(asDraftBufferDesc); + IBufferResource::Desc scratchBufferDesc; + scratchBufferDesc.type = IResource::Type::Buffer; + scratchBufferDesc.defaultState = ResourceState::UnorderedAccess; + scratchBufferDesc.sizeInBytes = accelerationStructurePrebuildInfo.scratchDataSize; + ComPtr<IBufferResource> scratchBuffer = gDevice->createBufferResource(scratchBufferDesc); + + // Build acceleration structure. + ComPtr<IQueryPool> compactedSizeQuery; + IQueryPool::Desc queryPoolDesc; + queryPoolDesc.count = 1; + queryPoolDesc.type = QueryType::AccelerationStructureCompactedSize; + SLANG_RETURN_ON_FAIL( + gDevice->createQueryPool(queryPoolDesc, compactedSizeQuery.writeRef())); + + ComPtr<IAccelerationStructure> draftAS; + IAccelerationStructure::CreateDesc draftCreateDesc; + draftCreateDesc.buffer = draftBuffer; + draftCreateDesc.kind = IAccelerationStructure::Kind::BottomLevel; + draftCreateDesc.offset = 0; + draftCreateDesc.size = accelerationStructurePrebuildInfo.resultDataMaxSize; + SLANG_RETURN_ON_FAIL( + gDevice->createAccelerationStructure(draftCreateDesc, draftAS.writeRef())); + + auto commandBuffer = gTransientHeaps[0]->createCommandBuffer(); + auto encoder = commandBuffer->encodeRayTracingCommands(); + IAccelerationStructure::BuildDesc buildDesc = {}; + buildDesc.dest = draftAS; + buildDesc.inputs = accelerationStructureBuildInputs; + buildDesc.scratchData = scratchBuffer->getDeviceAddress(); + AccelerationStructureQueryDesc compactedSizeQueryDesc = {}; + compactedSizeQueryDesc.queryPool = compactedSizeQuery; + compactedSizeQueryDesc.queryType = QueryType::AccelerationStructureCompactedSize; + encoder->buildAccelerationStructure(buildDesc, 1, &compactedSizeQueryDesc); + encoder->endEncoding(); + commandBuffer->close(); + gQueue->executeCommandBuffer(commandBuffer); + gQueue->wait(); + + uint64_t compactedSize = 0; + compactedSizeQuery->getResult(0, 1, &compactedSize); + IBufferResource::Desc asBufferDesc; + asBufferDesc.type = IResource::Type::Buffer; + asBufferDesc.defaultState = ResourceState::AccelerationStructure; + asBufferDesc.sizeInBytes = compactedSize; + gBLASBuffer = gDevice->createBufferResource(asBufferDesc); + IAccelerationStructure::CreateDesc createDesc; + createDesc.buffer = gBLASBuffer; + createDesc.kind = IAccelerationStructure::Kind::BottomLevel; + createDesc.offset = 0; + createDesc.size = compactedSize; + gDevice->createAccelerationStructure(createDesc, gBLAS.writeRef()); + + commandBuffer = gTransientHeaps[0]->createCommandBuffer(); + encoder = commandBuffer->encodeRayTracingCommands(); + encoder->copyAccelerationStructure(gBLAS, draftAS, AccelerationStructureCopyMode::Compact); + encoder->endEncoding(); + commandBuffer->close(); + gQueue->executeCommandBuffer(commandBuffer); + gQueue->wait(); + } + + // Build top level acceleration structure. + { + List<IAccelerationStructure::InstanceDesc> instanceDescs; + instanceDescs.setCount(1); + instanceDescs[0].accelerationStructure = gBLAS->getDeviceAddress(); + instanceDescs[0].flags = + IAccelerationStructure::GeometryInstanceFlags::TriangleFacingCullDisable; + instanceDescs[0].instanceContributionToHitGroupIndex = 0; + instanceDescs[0].instanceID = 0; + instanceDescs[0].instanceMask = 0xFF; + float transformMatrix[] = {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f}; + memcpy(&instanceDescs[0].transform[0][0], transformMatrix, sizeof(float) * 12); + + IBufferResource::Desc instanceBufferDesc; + instanceBufferDesc.type = IResource::Type::Buffer; + instanceBufferDesc.sizeInBytes = + instanceDescs.getCount() * sizeof(IAccelerationStructure::InstanceDesc); + instanceBufferDesc.defaultState = ResourceState::UnorderedAccess; + gInstanceBuffer = gDevice->createBufferResource(instanceBufferDesc, instanceDescs.getBuffer()); + if (!gInstanceBuffer) + return SLANG_FAIL; + + IAccelerationStructure::BuildInputs accelerationStructureBuildInputs = {}; + IAccelerationStructure::PrebuildInfo accelerationStructurePrebuildInfo = {}; + accelerationStructureBuildInputs.descCount = 1; + accelerationStructureBuildInputs.kind = IAccelerationStructure::Kind::TopLevel; + accelerationStructureBuildInputs.instanceDescs = gInstanceBuffer->getDeviceAddress(); + + // Query buffer size for acceleration structure build. + SLANG_RETURN_ON_FAIL(gDevice->getAccelerationStructurePrebuildInfo( + accelerationStructureBuildInputs, &accelerationStructurePrebuildInfo)); + + IBufferResource::Desc asBufferDesc; + asBufferDesc.type = IResource::Type::Buffer; + asBufferDesc.defaultState = ResourceState::AccelerationStructure; + asBufferDesc.sizeInBytes = accelerationStructurePrebuildInfo.resultDataMaxSize; + gTLASBuffer = gDevice->createBufferResource(asBufferDesc); + + IBufferResource::Desc scratchBufferDesc; + scratchBufferDesc.type = IResource::Type::Buffer; + scratchBufferDesc.defaultState = ResourceState::UnorderedAccess; + scratchBufferDesc.sizeInBytes = accelerationStructurePrebuildInfo.scratchDataSize; + ComPtr<IBufferResource> scratchBuffer = gDevice->createBufferResource(scratchBufferDesc); + + IAccelerationStructure::CreateDesc createDesc; + createDesc.buffer = gTLASBuffer; + createDesc.kind = IAccelerationStructure::Kind::TopLevel; + createDesc.offset = 0; + createDesc.size = accelerationStructurePrebuildInfo.resultDataMaxSize; + SLANG_RETURN_ON_FAIL(gDevice->createAccelerationStructure(createDesc, gTLAS.writeRef())); + + auto commandBuffer = gTransientHeaps[0]->createCommandBuffer(); + auto encoder = commandBuffer->encodeRayTracingCommands(); + IAccelerationStructure::BuildDesc buildDesc = {}; + buildDesc.dest = gTLAS; + buildDesc.inputs = accelerationStructureBuildInputs; + buildDesc.scratchData = scratchBuffer->getDeviceAddress(); + encoder->buildAccelerationStructure(buildDesc, 0, nullptr); + encoder->endEncoding(); + commandBuffer->close(); + gQueue->executeCommandBuffer(commandBuffer); + gQueue->wait(); + } + + IBufferResource::Desc fullScreenVertexBufferDesc; + fullScreenVertexBufferDesc.type = IResource::Type::Buffer; + fullScreenVertexBufferDesc.sizeInBytes = + FullScreenTriangle::kVertexCount * sizeof(FullScreenTriangle::Vertex); + fullScreenVertexBufferDesc.defaultState = ResourceState::VertexBuffer; + gFullScreenVertexBuffer = gDevice->createBufferResource( + fullScreenVertexBufferDesc, &FullScreenTriangle::kVertices[0]); + if (!gFullScreenVertexBuffer) + return SLANG_FAIL; + + InputElementDesc inputElements[] = { + {"POSITION", 0, Format::RG_Float32, offsetof(FullScreenTriangle::Vertex, position)}, + }; + auto inputLayout = gDevice->createInputLayout(&inputElements[0], SLANG_COUNT_OF(inputElements)); + if (!inputLayout) + return SLANG_FAIL; + + ComPtr<IShaderProgram> shaderProgram; + SLANG_RETURN_ON_FAIL(loadShaderProgram(gDevice, PipelineType::Graphics, shaderProgram.writeRef())); + GraphicsPipelineStateDesc desc; + desc.inputLayout = inputLayout; + desc.program = shaderProgram; + desc.framebufferLayout = gFramebufferLayout; + gPresentPipelineState = gDevice->createGraphicsPipelineState(desc); + if (!gPresentPipelineState) + return SLANG_FAIL; + + ComPtr<IShaderProgram> computeProgram; + SLANG_RETURN_ON_FAIL( + loadShaderProgram(gDevice, PipelineType::Compute, computeProgram.writeRef())); + ComputePipelineStateDesc computeDesc; + computeDesc.program = computeProgram; + gRenderPipelineState = gDevice->createComputePipelineState(computeDesc); + if (!gRenderPipelineState) + return SLANG_FAIL; + + createResultTexture(); + return SLANG_OK; +} + +void createResultTexture() +{ + ITextureResource::Desc resultTextureDesc = {}; + resultTextureDesc.type = IResource::Type::Texture2D; + resultTextureDesc.numMipLevels = 1; + resultTextureDesc.size.width = windowWidth; + resultTextureDesc.size.height = windowHeight; + resultTextureDesc.size.depth = 1; + resultTextureDesc.defaultState = ResourceState::UnorderedAccess; + resultTextureDesc.format = Format::RGBA_Float16; + gResultTexture = gDevice->createTextureResource(resultTextureDesc); + IResourceView::Desc resultUAVDesc = {}; + resultUAVDesc.format = resultTextureDesc.format; + resultUAVDesc.type = IResourceView::Type::UnorderedAccess; + gResultTextureUAV = gDevice->createTextureView(gResultTexture, resultUAVDesc); +} + +virtual void windowSizeChanged() override +{ + WindowedAppBase::windowSizeChanged(); + createResultTexture(); +} + +glm::vec3 getVectorFromSphericalAngles(float theta, float phi) +{ + auto sinTheta = sin(theta); + auto cosTheta = cos(theta); + auto sinPhi = sin(phi); + auto cosPhi = cos(phi); + return glm::vec3(-sinTheta * cosPhi, sinPhi, -cosTheta * cosPhi); +} +void updateUniforms() +{ + gUniforms.screenWidth = (float)windowWidth; + gUniforms.screenHeight = (float)windowHeight; + if (!lastTime) + lastTime = getCurrentTime(); + uint64_t currentTime = getCurrentTime(); + float deltaTime = float(double(currentTime - lastTime) / double(getTimerFrequency())); + lastTime = currentTime; + + auto camDir = + getVectorFromSphericalAngles(cameraOrientationAngles[0], cameraOrientationAngles[1]); + auto camUp = getVectorFromSphericalAngles( + cameraOrientationAngles[0], cameraOrientationAngles[1] + glm::pi<float>() * 0.5f); + auto camRight = glm::cross(camDir, camUp); + + glm::vec3 movement = glm::vec3(0); + if (wPressed) + movement += camDir; + if (sPressed) + movement -= camDir; + if (aPressed) + movement -= camRight; + if (dPressed) + movement += camRight; + + cameraPosition += deltaTime * translationScale * movement; + + memcpy(gUniforms.cameraDir, &camDir, sizeof(float) * 3); + memcpy(gUniforms.cameraUp, &camUp, sizeof(float) * 3); + memcpy(gUniforms.cameraRight, &camRight, sizeof(float) * 3); + memcpy(gUniforms.cameraPosition, &cameraPosition, sizeof(float) * 3); + auto lightDir = glm::normalize(glm::vec3(1.0f, 3.0f, 2.0f)); + memcpy(gUniforms.lightDir, &lightDir, sizeof(float) * 3); +} + +virtual void renderFrame(int frameBufferIndex) override +{ + updateUniforms(); + { + ComPtr<ICommandBuffer> renderCommandBuffer = + gTransientHeaps[frameBufferIndex]->createCommandBuffer(); + auto renderEncoder = renderCommandBuffer->encodeComputeCommands(); + auto rootObject = renderEncoder->bindPipeline(gRenderPipelineState); + auto cursor = ShaderCursor(rootObject->getEntryPoint(0)); + cursor["resultTexture"].setResource(gResultTextureUAV); + cursor["uniforms"].setData(&gUniforms, sizeof(Uniforms)); + cursor["sceneBVH"].setResource(gTLAS); + cursor["primitiveBuffer"].setResource(gPrimitiveBufferSRV); + renderEncoder->dispatchCompute((windowWidth + 15) / 16, (windowHeight + 15) / 16, 1); + renderEncoder->endEncoding(); + renderCommandBuffer->close(); + gQueue->executeCommandBuffer(renderCommandBuffer); + } + + { + ComPtr<ICommandBuffer> presentCommandBuffer = + gTransientHeaps[frameBufferIndex]->createCommandBuffer(); + auto presentEncoder = presentCommandBuffer->encodeRenderCommands( + gRenderPass, gFramebuffers[frameBufferIndex]); + gfx::Viewport viewport = {}; + viewport.maxZ = 1.0f; + viewport.extentX = (float)windowWidth; + viewport.extentY = (float)windowHeight; + presentEncoder->setViewportAndScissor(viewport); + auto rootObject = presentEncoder->bindPipeline(gPresentPipelineState); + auto cursor = ShaderCursor(rootObject->getEntryPoint(1)); + cursor["t"].setResource(gResultTextureUAV); + presentEncoder->setVertexBuffer( + 0, gFullScreenVertexBuffer, sizeof(FullScreenTriangle::Vertex)); + presentEncoder->setPrimitiveTopology(PrimitiveTopology::TriangleList); + presentEncoder->draw(3); + presentEncoder->endEncoding(); + presentCommandBuffer->close(); + gQueue->executeCommandBuffer(presentCommandBuffer); + } + // With that, we are done drawing for one frame, and ready for the next. + // + gSwapchain->present(); +} + +}; + +// This macro instantiates an appropriate main function to +// run the application defined above. +PLATFORM_UI_MAIN(innerMain<RayTracing>) diff --git a/examples/ray-tracing/shaders.slang b/examples/ray-tracing/shaders.slang new file mode 100644 index 000000000..0aff55435 --- /dev/null +++ b/examples/ray-tracing/shaders.slang @@ -0,0 +1,144 @@ +// shaders.slang + +struct Uniforms +{ + float screenWidth, screenHeight; + float focalLength, frameHeight; + float4 cameraDir; + float4 cameraUp; + float4 cameraRight; + float4 cameraPosition; + float4 lightDir; +}; + +struct Primitive +{ + float4 data0; + float4 color; + float3 getNormal() { return data0.xyz; } + float3 getColor() { return color.xyz; } +}; + +bool traceRayFirstHit( + RaytracingAccelerationStructure sceneBVH, + float3 rayOrigin, + float3 rayDir, + out float t, + out int primitiveIndex) +{ + RayDesc ray; + ray.Origin = rayOrigin; + ray.TMin = 0.01f; + ray.Direction = rayDir; + ray.TMax = 1e4f; + RayQuery<RAY_FLAG_SKIP_PROCEDURAL_PRIMITIVES | + RAY_FLAG_ACCEPT_FIRST_HIT_AND_END_SEARCH> q; + let rayFlags = RAY_FLAG_SKIP_PROCEDURAL_PRIMITIVES | + RAY_FLAG_ACCEPT_FIRST_HIT_AND_END_SEARCH; + + q.TraceRayInline( + sceneBVH, + rayFlags, + 0xff, + ray); + q.Proceed(); + + if(q.CommittedStatus() == COMMITTED_TRIANGLE_HIT) + { + t = q.CommittedRayT(); + primitiveIndex = q.CommittedPrimitiveIndex(); + return true; + } + return false; +} + +bool traceRayNearestHit( + RaytracingAccelerationStructure sceneBVH, + float3 rayOrigin, + float3 rayDir, + out float t, + out int primitiveIndex) +{ + RayDesc ray; + ray.Origin = rayOrigin; + ray.TMin = 0.01f; + ray.Direction = rayDir; + ray.TMax = 1e4f; + RayQuery<RAY_FLAG_NONE> q; + let rayFlags = RAY_FLAG_NONE; + + q.TraceRayInline( + sceneBVH, + rayFlags, + 0xff, + ray); + + q.Proceed(); + if(q.CommittedStatus() == COMMITTED_TRIANGLE_HIT) + { + t = q.CommittedRayT(); + primitiveIndex = q.CommittedPrimitiveIndex(); + return true; + } + return false; +} + +[shader("compute")] +[numthreads(16,16,1)] +void computeMain( + uint3 threadIdx : SV_DispatchThreadID, + uniform RWTexture2D resultTexture, + uniform RaytracingAccelerationStructure sceneBVH, + uniform StructuredBuffer<Primitive> primitiveBuffer, + uniform Uniforms uniforms) +{ + if (threadIdx.x >= (int)uniforms.screenWidth) return; + if (threadIdx.y >= (int)uniforms.screenHeight) return; + + float frameWidth = uniforms.screenWidth / uniforms.screenHeight * uniforms.frameHeight; + float imageY = (threadIdx.y / uniforms.screenHeight - 0.5f) * uniforms.frameHeight; + float imageX = (threadIdx.x / uniforms.screenWidth - 0.5f) * frameWidth; + float imageZ = uniforms.focalLength; + float3 rayDir = normalize(uniforms.cameraDir.xyz*imageZ - uniforms.cameraUp.xyz * imageY + uniforms.cameraRight.xyz * imageX); + + float4 resultColor = 0; + + int primitiveIndex; + float intersectionT; + if (traceRayNearestHit(sceneBVH, uniforms.cameraPosition.xyz, rayDir, intersectionT, primitiveIndex)) + { + float3 hitLocation = uniforms.cameraPosition.xyz + rayDir * intersectionT; + float3 shadowRayDir = uniforms.lightDir.xyz; + float shadow = 1.0; + float shadowIntersectionT; + int shadowPrimitiveIndex; + if (traceRayFirstHit(sceneBVH, hitLocation, shadowRayDir, shadowIntersectionT, shadowPrimitiveIndex)) + { + shadow = 0.0f; + } + float3 normal = primitiveBuffer[primitiveIndex].getNormal(); + float3 color = primitiveBuffer[primitiveIndex].getColor(); + float ndotl = max(0.0, shadow * dot(normal, uniforms.lightDir.xyz)); + float intensity = ndotl * 0.7 + 0.3; + resultColor = float4(color * intensity, 1.0f); + } + resultTexture[threadIdx.xy] = resultColor; +} + +/// Vertex and fragment shader for displaying the final image. + +[shader("vertex")] +float4 vertexMain(float2 position : POSITION) + : SV_Position +{ + return float4(position, 0.5, 1.0); +} + +[shader("fragment")] +float4 fragmentMain( + float4 sv_position : SV_Position, + uniform RWTexture2D t) + : SV_Target +{ + return t.Load(sv_position.xy); +} diff --git a/examples/shader-object/main.cpp b/examples/shader-object/main.cpp index 71c5de983..8240aa63b 100644 --- a/examples/shader-object/main.cpp +++ b/examples/shader-object/main.cpp @@ -171,7 +171,7 @@ int main() ResourceState::CopyDestination, ResourceState::CopySource); bufferDesc.defaultState = ResourceState::UnorderedAccess; - bufferDesc.cpuAccessFlags = IResource::AccessFlag::Write | IResource::AccessFlag::Read; + bufferDesc.cpuAccessFlags = AccessFlag::Write | AccessFlag::Read; ComPtr<gfx::IBufferResource> numbersBuffer; SLANG_RETURN_ON_FAIL(device->createBufferResource( diff --git a/premake5.lua b/premake5.lua index 1717cdcf2..77c69b345 100644 --- a/premake5.lua +++ b/premake5.lua @@ -652,6 +652,8 @@ example "hello-world" -- Let's go ahead and set up the projects for our other example now. example "triangle" +example "ray-tracing" + example "gpu-printing" kind "ConsoleApp" diff --git a/slang-gfx.h b/slang-gfx.h index c80d7e0fb..effecb620 100644 --- a/slang-gfx.h +++ b/slang-gfx.h @@ -214,6 +214,7 @@ enum class ResourceState CopyDestination, ResolveSource, ResolveDestination, + AccelerationStructure, _Count }; @@ -242,6 +243,17 @@ private: void add() {} }; + +/// Combinations describe how a resource can be accessed (typically by the host/cpu) +struct AccessFlag +{ + enum Enum + { + Read = 0x1, + Write = 0x2 + }; +}; + class IResource: public ISlangUnknown { public: @@ -258,16 +270,6 @@ public: CountOf, }; - /// Combinations describe how a resource can be accessed (typically by the host/cpu) - struct AccessFlag - { - enum Enum - { - Read = 0x1, - Write = 0x2 - }; - }; - /// Base class for Descs struct DescBase { @@ -474,6 +476,7 @@ public: DepthStencil, ShaderResource, UnorderedAccess, + AccelerationStructure, }; struct RenderTargetDesc @@ -495,11 +498,163 @@ public: RenderTargetDesc renderTarget; }; }; -#define SLANG_UUID_IResourceView \ +#define SLANG_UUID_IResourceView \ { \ 0x7b6c4926, 0x884, 0x408c, { 0xad, 0x8a, 0x50, 0x3a, 0x8e, 0x23, 0x98, 0xa4 } \ } +class IAccelerationStructure : public IResourceView +{ +public: + enum class Kind + { + TopLevel, + BottomLevel + }; + + struct BuildFlags + { + enum Enum + { + None, + AllowUpdate = 1, + AllowCompaction = 2, + PreferFastTrace = 4, + PreferFastBuild = 8, + MinimizeMemory = 16, + PerformUpdate = 32 + }; + }; + + enum class GeometryType + { + Triangles, ProcedurePrimitives + }; + + struct GeometryFlags + { + enum Enum + { + None, + Opaque = 1, + NoDuplicateAnyHitInvocation = 2 + }; + }; + + struct TriangleDesc + { + DeviceAddress transform3x4; + Format indexFormat; + Format vertexFormat; + uint32_t indexCount; + uint32_t vertexCount; + DeviceAddress indexData; + DeviceAddress vertexData; + uint64_t vertexStride; + }; + + struct ProceduralAABB + { + float minX; + float minY; + float minZ; + float maxX; + float maxY; + float maxZ; + }; + + struct ProceduralAABBDesc + { + // Number of AABBs. + uint64_t count; + + /// Pointer to an array of `ProceduralAABB` values in device memory. + DeviceAddress data; + + /// Stride in bytes of the AABB values array. + uint64_t stride; + }; + + struct GeometryDesc + { + GeometryType type; + GeometryFlags::Enum flags; + union + { + TriangleDesc triangles; + ProceduralAABBDesc proceduralAABBs; + } content; + }; + + struct GeometryInstanceFlags + { + enum Enum : uint32_t + { + None = 0, + TriangleFacingCullDisable = 0x00000001, + TriangleFrontCounterClockwise = 0x00000002, + ForceOpaque = 0x00000004, + NoOpaque = 0x00000008 + }; + }; + + struct InstanceDesc + { + float transform[3][4]; + uint32_t instanceID : 24; + uint32_t instanceMask : 8; + uint32_t instanceContributionToHitGroupIndex : 24; + GeometryInstanceFlags::Enum flags : 8; + DeviceAddress accelerationStructure; + }; + + struct PrebuildInfo + { + uint64_t resultDataMaxSize; + uint64_t scratchDataSize; + uint64_t updateScratchDataSize; + }; + + struct BuildInputs + { + Kind kind; + + BuildFlags::Enum flags; + + int32_t descCount; + + /// Array of `InstanceDesc` values in device memory. + /// Used when `kind` is `TopLevel`. + DeviceAddress instanceDescs; + + /// Array of `GeometryDesc` values. + /// Used when `kind` is `BottomLevel`. + const GeometryDesc* geometryDescs; + }; + + struct CreateDesc + { + Kind kind; + IBufferResource* buffer; + uint64_t offset; + uint64_t size; + }; + + struct BuildDesc + { + BuildInputs inputs; + IAccelerationStructure* source; + IAccelerationStructure* dest; + DeviceAddress scratchData; + }; + + virtual SLANG_NO_THROW DeviceAddress SLANG_MCALL getDeviceAddress() = 0; +}; +#define SLANG_UUID_IAccelerationStructure \ + { \ + 0xa5cdda3c, 0x1d4e, 0x4df7, { 0x8e, 0xf2, 0xb7, 0x3f, 0xce, 0x4, 0xde, 0x3b } \ + } + struct ShaderOffset { SlangInt uniformOffset = 0; @@ -854,6 +1009,8 @@ public: enum class QueryType { Timestamp, + AccelerationStructureCompactedSize, + AccelerationStructureSerializedSize, }; class IQueryPool : public ISlangUnknown @@ -871,16 +1028,12 @@ public: { 0xc2cc3784, 0x12da, 0x480a, { 0xa8, 0x74, 0x8b, 0x31, 0x96, 0x1c, 0xa4, 0x36 } } -class ICommandEncoder : public ISlangUnknown +class ICommandEncoder { public: virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() = 0; virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* queryPool, SlangInt queryIndex) = 0; }; -#define SLANG_UUID_ICommandEncoder \ - { \ - 0xbd0717f8, 0xc4a7, 0x4603, { 0x94, 0xd4, 0x6f, 0x8f, 0x95, 0x16, 0x91, 0x47 } \ - } class IRenderCommandEncoder : public ICommandEncoder { @@ -933,10 +1086,6 @@ public: drawIndexed(UInt indexCount, UInt startIndex = 0, UInt baseVertex = 0) = 0; virtual SLANG_NO_THROW void SLANG_MCALL setStencilReference(uint32_t referenceValue) = 0; }; -#define SLANG_UUID_IRenderCommandEncoder \ - { \ - 0x39417cf7, 0x8d97, 0x43a9, { 0xbb, 0x9f, 0x2f, 0x35, 0xe9, 0x11, 0xd0, 0x42 } \ - } class IComputeCommandEncoder : public ICommandEncoder { @@ -956,10 +1105,6 @@ public: } virtual SLANG_NO_THROW void SLANG_MCALL dispatchCompute(int x, int y, int z) = 0; }; -#define SLANG_UUID_IComputeCommandEncoder \ - { \ - 0x65400452, 0xc877, 0x478f, { 0x91, 0x7d, 0x48, 0xd5, 0x41, 0x6f, 0x39, 0xab } \ - } class IResourceCommandEncoder : public ICommandEncoder { @@ -973,9 +1118,50 @@ public: virtual SLANG_NO_THROW void SLANG_MCALL uploadBufferData(IBufferResource* dst, size_t offset, size_t size, void* data) = 0; }; -#define SLANG_UUID_IResourceCommandEncoder \ + +enum class AccelerationStructureCopyMode +{ + Clone, Compact +}; + +struct AccelerationStructureQueryDesc +{ + QueryType queryType; + + IQueryPool* queryPool; + + int32_t firstQueryIndex; +}; + +class IRayTracingCommandEncoder : public ICommandEncoder +{ +public: + virtual SLANG_NO_THROW void SLANG_MCALL buildAccelerationStructure( + const IAccelerationStructure::BuildDesc& desc, + int propertyQueryCount, + AccelerationStructureQueryDesc* queryDescs) = 0; + virtual SLANG_NO_THROW void SLANG_MCALL copyAccelerationStructure( + IAccelerationStructure* dest, + IAccelerationStructure* src, + AccelerationStructureCopyMode mode) = 0; + virtual SLANG_NO_THROW void SLANG_MCALL queryAccelerationStructureProperties( + int accelerationStructureCount, + IAccelerationStructure* const* accelerationStructures, + int queryCount, + AccelerationStructureQueryDesc* queryDescs) = 0; + virtual SLANG_NO_THROW void SLANG_MCALL + serializeAccelerationStructure(DeviceAddress dest, IAccelerationStructure* source) = 0; + virtual SLANG_NO_THROW void SLANG_MCALL + deserializeAccelerationStructure(IAccelerationStructure* dest, DeviceAddress source) = 0; + virtual SLANG_NO_THROW void memoryBarrier( + int count, + IAccelerationStructure* const* structures, + AccessFlag::Enum sourceAccess, + AccessFlag::Enum destAccess) = 0; +}; +#define SLANG_UUID_IRayTracingCommandEncoder \ { \ - 0x5fe87643, 0x7ad7, 0x4177, { 0x8b, 0xd1, 0xd7, 0x84, 0xad, 0xcf, 0x3d, 0xce } \ + 0x9a672b87, 0x5035, 0x45e3, { 0x96, 0x7c, 0x1f, 0x85, 0xcd, 0xb3, 0x63, 0x4f } \ } class ICommandBuffer : public ISlangUnknown @@ -990,29 +1176,38 @@ public: IRenderPassLayout* renderPass, IFramebuffer* framebuffer, IRenderCommandEncoder** outEncoder) = 0; - ComPtr<IRenderCommandEncoder> + IRenderCommandEncoder* encodeRenderCommands(IRenderPassLayout* renderPass, IFramebuffer* framebuffer) { - ComPtr<IRenderCommandEncoder> result; - encodeRenderCommands(renderPass, framebuffer, result.writeRef()); + IRenderCommandEncoder* result; + encodeRenderCommands(renderPass, framebuffer, &result); return result; } virtual SLANG_NO_THROW void SLANG_MCALL encodeComputeCommands(IComputeCommandEncoder** outEncoder) = 0; - ComPtr<IComputeCommandEncoder> encodeComputeCommands() + IComputeCommandEncoder* encodeComputeCommands() { - ComPtr<IComputeCommandEncoder> result; - encodeComputeCommands(result.writeRef()); + IComputeCommandEncoder* result; + encodeComputeCommands(&result); return result; } virtual SLANG_NO_THROW void SLANG_MCALL encodeResourceCommands(IResourceCommandEncoder** outEncoder) = 0; - ComPtr<IResourceCommandEncoder> encodeResourceCommands() + IResourceCommandEncoder* encodeResourceCommands() + { + IResourceCommandEncoder* result; + encodeResourceCommands(&result); + return result; + } + + virtual SLANG_NO_THROW void SLANG_MCALL + encodeRayTracingCommands(IRayTracingCommandEncoder** outEncoder) = 0; + IRayTracingCommandEncoder* encodeRayTracingCommands() { - ComPtr<IResourceCommandEncoder> result; - encodeResourceCommands(result.writeRef()); + IRayTracingCommandEncoder* result; + encodeRayTracingCommands(&result); return result; } @@ -1407,6 +1602,15 @@ public: virtual SLANG_NO_THROW Result SLANG_MCALL createQueryPool( const IQueryPool::Desc& desc, IQueryPool** outPool) = 0; + + + virtual SLANG_NO_THROW Result SLANG_MCALL getAccelerationStructurePrebuildInfo( + const IAccelerationStructure::BuildInputs& buildInputs, + IAccelerationStructure::PrebuildInfo* outPrebuildInfo) = 0; + + virtual SLANG_NO_THROW Result SLANG_MCALL createAccelerationStructure( + const IAccelerationStructure::CreateDesc& desc, + IAccelerationStructure** outView) = 0; }; #define SLANG_UUID_IDevice \ @@ -33,6 +33,8 @@ Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "hello-world", "build\visual EndProject Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "model-viewer", "build\visual-studio\model-viewer\model-viewer.vcxproj", "{2F8724C6-1BC3-2730-84D5-3F277030D04A}" EndProject +Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "ray-tracing", "build\visual-studio\ray-tracing\ray-tracing.vcxproj", "{71AC0F50-5DFD-FA91-8661-E95372118EFB}" +EndProject Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "shader-object", "build\visual-studio\shader-object\shader-object.vcxproj", "{25512BFB-1138-EDF2-BA88-5310A64E6659}" EndProject Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "shader-toy", "build\visual-studio\shader-toy\shader-toy.vcxproj", "{0FC5DE93-FBEA-A8FA-E430-2EC6D0F5CDC6}" @@ -185,6 +187,14 @@ 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 + {71AC0F50-5DFD-FA91-8661-E95372118EFB}.Debug|Win32.ActiveCfg = Debug|Win32 + {71AC0F50-5DFD-FA91-8661-E95372118EFB}.Debug|Win32.Build.0 = Debug|Win32 + {71AC0F50-5DFD-FA91-8661-E95372118EFB}.Debug|x64.ActiveCfg = Debug|x64 + {71AC0F50-5DFD-FA91-8661-E95372118EFB}.Debug|x64.Build.0 = Debug|x64 + {71AC0F50-5DFD-FA91-8661-E95372118EFB}.Release|Win32.ActiveCfg = Release|Win32 + {71AC0F50-5DFD-FA91-8661-E95372118EFB}.Release|Win32.Build.0 = Release|Win32 + {71AC0F50-5DFD-FA91-8661-E95372118EFB}.Release|x64.ActiveCfg = Release|x64 + {71AC0F50-5DFD-FA91-8661-E95372118EFB}.Release|x64.Build.0 = Release|x64 {25512BFB-1138-EDF2-BA88-5310A64E6659}.Debug|Win32.ActiveCfg = Debug|Win32 {25512BFB-1138-EDF2-BA88-5310A64E6659}.Debug|Win32.Build.0 = Debug|Win32 {25512BFB-1138-EDF2-BA88-5310A64E6659}.Debug|x64.ActiveCfg = Debug|x64 @@ -282,6 +292,7 @@ Global {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} + {71AC0F50-5DFD-FA91-8661-E95372118EFB} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231} {25512BFB-1138-EDF2-BA88-5310A64E6659} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231} {0FC5DE93-FBEA-A8FA-E430-2EC6D0F5CDC6} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231} {3BB99068-27C9-3C39-9082-A1577CB12BD2} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231} diff --git a/tools/gfx/cuda/render-cuda.cpp b/tools/gfx/cuda/render-cuda.cpp index 5a7e7babf..ffd197e57 100644 --- a/tools/gfx/cuda/render-cuda.cpp +++ b/tools/gfx/cuda/render-cuda.cpp @@ -921,22 +921,6 @@ public: : public IComputeCommandEncoder { public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IComputeCommandEncoder) - { - *outObject = static_cast<IComputeCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - - public: CommandWriter* m_writer; CommandBufferImpl* m_commandBuffer; RefPtr<ShaderObjectBase> m_rootObject; @@ -982,22 +966,6 @@ public: : public IResourceCommandEncoder { public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IResourceCommandEncoder) - { - *outObject = static_cast<IResourceCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - - public: CommandWriter* m_writer; void init(CommandBufferImpl* cmdBuffer) @@ -1037,6 +1005,12 @@ public: *outEncoder = &m_resourceCommandEncoder; } + virtual SLANG_NO_THROW void SLANG_MCALL + encodeRayTracingCommands(IRayTracingCommandEncoder** outEncoder) override + { + *outEncoder = nullptr; + } + virtual SLANG_NO_THROW void SLANG_MCALL close() override {} }; diff --git a/tools/gfx/d3d11/render-d3d11.cpp b/tools/gfx/d3d11/render-d3d11.cpp index f2fe9b744..d74888812 100644 --- a/tools/gfx/d3d11/render-d3d11.cpp +++ b/tools/gfx/d3d11/render-d3d11.cpp @@ -1536,7 +1536,7 @@ protected: bufferDesc.defaultState = ResourceState::ConstantBuffer; bufferDesc.allowedStates = ResourceStateSet(ResourceState::ConstantBuffer, ResourceState::CopyDestination); - bufferDesc.cpuAccessFlags |= IResource::AccessFlag::Write; + bufferDesc.cpuAccessFlags |= AccessFlag::Write; SLANG_RETURN_ON_FAIL( device->createBufferResource(bufferDesc, nullptr, bufferResourcePtr.writeRef())); m_ordinaryDataBuffer = static_cast<BufferResourceImpl*>(bufferResourcePtr.get()); @@ -2506,10 +2506,10 @@ static int _calcResourceAccessFlags(int accessFlags) switch (accessFlags) { case 0: return 0; - case IResource::AccessFlag::Read: return D3D11_CPU_ACCESS_READ; - case IResource::AccessFlag::Write: return D3D11_CPU_ACCESS_WRITE; - case IResource::AccessFlag::Read | - IResource::AccessFlag::Write: return D3D11_CPU_ACCESS_READ | D3D11_CPU_ACCESS_WRITE; + case AccessFlag::Read: return D3D11_CPU_ACCESS_READ; + case AccessFlag::Write: return D3D11_CPU_ACCESS_WRITE; + case AccessFlag::Read | + AccessFlag::Write: return D3D11_CPU_ACCESS_READ | D3D11_CPU_ACCESS_WRITE; default: assert(!"Invalid flags"); return 0; } } @@ -2663,11 +2663,11 @@ Result D3D11Device::createBufferResource(const IBufferResource::Desc& descIn, co bufferDesc.BindFlags = d3dBindFlags; // For read we'll need to do some staging bufferDesc.CPUAccessFlags = - _calcResourceAccessFlags(descIn.cpuAccessFlags & IResource::AccessFlag::Write); + _calcResourceAccessFlags(descIn.cpuAccessFlags & AccessFlag::Write); bufferDesc.Usage = D3D11_USAGE_DEFAULT; // If written by CPU, make it dynamic - if ((descIn.cpuAccessFlags & IResource::AccessFlag::Write) && + if ((descIn.cpuAccessFlags & AccessFlag::Write) && !descIn.allowedStates.contains(ResourceState::UnorderedAccess)) { bufferDesc.Usage = D3D11_USAGE_DYNAMIC; @@ -2698,7 +2698,7 @@ Result D3D11Device::createBufferResource(const IBufferResource::Desc& descIn, co } } - if (srcDesc.cpuAccessFlags & IResource::AccessFlag::Write) + if (srcDesc.cpuAccessFlags & AccessFlag::Write) { bufferDesc.CPUAccessFlags |= D3D11_CPU_ACCESS_WRITE; } @@ -2711,8 +2711,8 @@ Result D3D11Device::createBufferResource(const IBufferResource::Desc& descIn, co SLANG_RETURN_ON_FAIL(m_device->CreateBuffer(&bufferDesc, initData ? &subResourceData : nullptr, buffer->m_buffer.writeRef())); buffer->m_d3dUsage = bufferDesc.Usage; - if ((srcDesc.cpuAccessFlags & IResource::AccessFlag::Read) || - ((srcDesc.cpuAccessFlags & IResource::AccessFlag::Write) && bufferDesc.Usage != D3D11_USAGE_DYNAMIC)) + if ((srcDesc.cpuAccessFlags & AccessFlag::Read) || + ((srcDesc.cpuAccessFlags & AccessFlag::Write) && bufferDesc.Usage != D3D11_USAGE_DYNAMIC)) { D3D11_BUFFER_DESC bufDesc = {}; bufDesc.BindFlags = 0; diff --git a/tools/gfx/d3d12/render-d3d12.cpp b/tools/gfx/d3d12/render-d3d12.cpp index 79cbe2f9a..b8c95ef3a 100644 --- a/tools/gfx/d3d12/render-d3d12.cpp +++ b/tools/gfx/d3d12/render-d3d12.cpp @@ -578,7 +578,7 @@ public: bufferDesc.allowedStates = ResourceStateSet(ResourceState::ConstantBuffer, ResourceState::CopyDestination); bufferDesc.sizeInBytes = desc.constantBufferSize; - bufferDesc.cpuAccessFlags |= IResource::AccessFlag::Write; + bufferDesc.cpuAccessFlags |= AccessFlag::Write; SLANG_RETURN_ON_FAIL(device->createBufferResource( bufferDesc, nullptr, @@ -3098,21 +3098,6 @@ public: , public PipelineCommandEncoder { public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IRenderCommandEncoder) - { - *outObject = static_cast<IRenderCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - public: RefPtr<RenderPassLayoutImpl> m_renderPass; RefPtr<FramebufferImpl> m_framebuffer; @@ -3481,22 +3466,6 @@ public: , public PipelineCommandEncoder { public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IComputeCommandEncoder) - { - *outObject = static_cast<IComputeCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() { return 1; } - - public: virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override { PipelineCommandEncoder::endEncodingImpl(); @@ -3548,22 +3517,6 @@ public: class ResourceCommandEncoderImpl : public IResourceCommandEncoder { public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IResourceCommandEncoder) - { - *outObject = static_cast<IResourceCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() { return 1; } - - public: CommandBufferImpl* m_commandBuffer; void init(D3D12Device* renderer, CommandBufferImpl* commandBuffer) { @@ -3615,6 +3568,12 @@ public: *outEncoder = &m_resourceCommandEncoder; } + virtual SLANG_NO_THROW void SLANG_MCALL + encodeRayTracingCommands(IRayTracingCommandEncoder** outEncoder) override + { + *outEncoder = nullptr; + } + virtual SLANG_NO_THROW void SLANG_MCALL close() override { m_cmdList->Close(); } }; diff --git a/tools/gfx/debug-layer.cpp b/tools/gfx/debug-layer.cpp index 26aa36535..7ad0dd31c 100644 --- a/tools/gfx/debug-layer.cpp +++ b/tools/gfx/debug-layer.cpp @@ -126,9 +126,6 @@ SLANG_GFX_DEBUG_GET_INTERFACE_IMPL_PARENT(BufferResource, Resource) SLANG_GFX_DEBUG_GET_INTERFACE_IMPL_PARENT(TextureResource, Resource) SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(CommandBuffer) SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(CommandQueue) -SLANG_GFX_DEBUG_GET_INTERFACE_IMPL_PARENT(ComputeCommandEncoder, CommandEncoder) -SLANG_GFX_DEBUG_GET_INTERFACE_IMPL_PARENT(RenderCommandEncoder, CommandEncoder) -SLANG_GFX_DEBUG_GET_INTERFACE_IMPL_PARENT(ResourceCommandEncoder, CommandEncoder) SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(Framebuffer) SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(FramebufferLayout) SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(InputLayout) @@ -141,11 +138,48 @@ SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(ShaderProgram) SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(Swapchain) SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(TransientResourceHeap) SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(QueryPool) +SLANG_GFX_DEBUG_GET_INTERFACE_IMPL_PARENT(AccelerationStructure, ResourceView) #undef SLANG_GFX_DEBUG_GET_INTERFACE_IMPL #undef SLANG_GFX_DEBUG_GET_INTERFACE_IMPL_PARENT +// Utility conversion functions to get Debug* object or the inner object from a user provided +// pointer. +#define SLANG_GFX_DEBUG_GET_OBJ_IMPL(type) \ + static Debug##type* getDebugObj(I##type* ptr) { return static_cast<Debug##type*>(ptr); } \ + static I##type* getInnerObj(I##type* ptr) \ + { \ + if (!ptr) return nullptr; \ + auto debugObj = getDebugObj(ptr); \ + return debugObj->baseObject; \ + } + +SLANG_GFX_DEBUG_GET_OBJ_IMPL(Device) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(BufferResource) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(TextureResource) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(CommandBuffer) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(CommandQueue) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(ComputeCommandEncoder) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(RenderCommandEncoder) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(ResourceCommandEncoder) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(RayTracingCommandEncoder) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(Framebuffer) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(FramebufferLayout) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(InputLayout) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(RenderPassLayout) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(PipelineState) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(ResourceView) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(SamplerState) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(ShaderObject) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(ShaderProgram) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(Swapchain) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(TransientResourceHeap) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(QueryPool) +SLANG_GFX_DEBUG_GET_OBJ_IMPL(AccelerationStructure) + +#undef SLANG_GFX_DEBUG_GET_OBJ_IMPL + Result DebugDevice::getFeatures(const char** outFeatures, UInt bufferSize, UInt* outFeatureCount) { SLANG_GFX_API_FUNC; @@ -267,6 +301,30 @@ Result DebugDevice::createBufferView( return result; } +Result DebugDevice::getAccelerationStructurePrebuildInfo( + const IAccelerationStructure::BuildInputs& buildInputs, + IAccelerationStructure::PrebuildInfo* outPrebuildInfo) +{ + SLANG_GFX_API_FUNC; + + return baseObject->getAccelerationStructurePrebuildInfo(buildInputs, outPrebuildInfo); +} + +Result DebugDevice::createAccelerationStructure( + const IAccelerationStructure::CreateDesc& desc, + IAccelerationStructure** outAS) +{ + SLANG_GFX_API_FUNC; + auto innerDesc = desc; + innerDesc.buffer = getInnerObj(innerDesc.buffer); + RefPtr<DebugAccelerationStructure> outObject = new DebugAccelerationStructure(); + auto result = baseObject->createAccelerationStructure(innerDesc, outObject->baseObject.writeRef()); + if (SLANG_FAILED(result)) + return result; + returnComPtr(outAS, outObject); + return SLANG_OK; +} + Result DebugDevice::createFramebufferLayout( IFramebufferLayout::Desc const& desc, IFramebufferLayout** outFrameBuffer) @@ -540,7 +598,7 @@ void DebugCommandBuffer::encodeRenderCommands( framebuffer ? static_cast<DebugFramebuffer*>(framebuffer)->baseObject : nullptr; m_renderCommandEncoder.isOpen = true; baseObject->encodeRenderCommands( - innerRenderPass, innerFramebuffer, m_renderCommandEncoder.baseObject.writeRef()); + innerRenderPass, innerFramebuffer, &m_renderCommandEncoder.baseObject); if (m_renderCommandEncoder.baseObject) *outEncoder = &m_renderCommandEncoder; else @@ -553,8 +611,15 @@ void DebugCommandBuffer::encodeComputeCommands(IComputeCommandEncoder** outEncod checkCommandBufferOpenWhenCreatingEncoder(); checkEncodersClosedBeforeNewEncoder(); m_computeCommandEncoder.isOpen = true; - baseObject->encodeComputeCommands(m_computeCommandEncoder.baseObject.writeRef()); - *outEncoder = &m_computeCommandEncoder; + baseObject->encodeComputeCommands(&m_computeCommandEncoder.baseObject); + if (m_computeCommandEncoder.baseObject) + { + *outEncoder = &m_computeCommandEncoder; + } + else + { + *outEncoder = nullptr; + } } void DebugCommandBuffer::encodeResourceCommands(IResourceCommandEncoder** outEncoder) @@ -563,8 +628,32 @@ void DebugCommandBuffer::encodeResourceCommands(IResourceCommandEncoder** outEnc checkCommandBufferOpenWhenCreatingEncoder(); checkEncodersClosedBeforeNewEncoder(); m_resourceCommandEncoder.isOpen = true; - baseObject->encodeResourceCommands(m_resourceCommandEncoder.baseObject.writeRef()); - *outEncoder = &m_resourceCommandEncoder; + baseObject->encodeResourceCommands(&m_resourceCommandEncoder.baseObject); + if (m_resourceCommandEncoder.baseObject) + { + *outEncoder = &m_resourceCommandEncoder; + } + else + { + *outEncoder = nullptr; + } +} + +void DebugCommandBuffer::encodeRayTracingCommands(IRayTracingCommandEncoder** outEncoder) +{ + SLANG_GFX_API_FUNC; + checkCommandBufferOpenWhenCreatingEncoder(); + checkEncodersClosedBeforeNewEncoder(); + m_rayTracingCommandEncoder.isOpen = true; + baseObject->encodeRayTracingCommands(&m_rayTracingCommandEncoder.baseObject); + if (m_rayTracingCommandEncoder.baseObject) + { + *outEncoder = &m_rayTracingCommandEncoder; + } + else + { + *outEncoder = nullptr; + } } void DebugCommandBuffer::close() @@ -619,6 +708,7 @@ void DebugCommandBuffer::checkCommandBufferOpenWhenCreatingEncoder() void DebugComputeCommandEncoder::endEncoding() { SLANG_GFX_API_FUNC; + isOpen = false; baseObject->endEncoding(); } @@ -650,6 +740,7 @@ void DebugComputeCommandEncoder::writeTimestamp(IQueryPool* pool, SlangInt index void DebugRenderCommandEncoder::endEncoding() { SLANG_GFX_API_FUNC; + isOpen = false; baseObject->endEncoding(); } @@ -738,6 +829,7 @@ void DebugRenderCommandEncoder::writeTimestamp(IQueryPool* pool, SlangInt index) void DebugResourceCommandEncoder::endEncoding() { SLANG_GFX_API_FUNC; + isOpen = false; baseObject->endEncoding(); } @@ -771,6 +863,103 @@ void DebugResourceCommandEncoder::uploadBufferData( baseObject->uploadBufferData(dstImpl->baseObject, offset, size, data); } +void DebugRayTracingCommandEncoder::endEncoding() +{ + SLANG_GFX_API_FUNC; + isOpen = false; + baseObject->endEncoding(); +} + +SLANG_NO_THROW void SLANG_MCALL + DebugRayTracingCommandEncoder::writeTimestamp(IQueryPool* pool, SlangInt index) +{ + SLANG_GFX_API_FUNC; + baseObject->writeTimestamp(static_cast<DebugQueryPool*>(pool)->baseObject, index); +} + +void DebugRayTracingCommandEncoder::buildAccelerationStructure( + const IAccelerationStructure::BuildDesc& desc, + int propertyQueryCount, + AccelerationStructureQueryDesc* queryDescs) +{ + SLANG_GFX_API_FUNC; + IAccelerationStructure::BuildDesc innerDesc = desc; + innerDesc.dest = getInnerObj(innerDesc.dest); + innerDesc.source = getInnerObj(innerDesc.source); + List<AccelerationStructureQueryDesc> innerQueryDescs; + innerQueryDescs.addRange(queryDescs, propertyQueryCount); + for (auto& innerQueryDesc : innerQueryDescs) + { + innerQueryDesc.queryPool = getInnerObj(innerQueryDesc.queryPool); + } + baseObject->buildAccelerationStructure( + innerDesc, propertyQueryCount, innerQueryDescs.getBuffer()); +} + +void DebugRayTracingCommandEncoder::copyAccelerationStructure( + IAccelerationStructure* dest, + IAccelerationStructure* src, + AccelerationStructureCopyMode mode) +{ + SLANG_GFX_API_FUNC; + auto innerDest = getInnerObj(dest); + auto innerSrc = getInnerObj(src); + baseObject->copyAccelerationStructure(innerDest, innerSrc, mode); +} + +void DebugRayTracingCommandEncoder::queryAccelerationStructureProperties( + int accelerationStructureCount, + IAccelerationStructure* const* accelerationStructures, + int queryCount, + AccelerationStructureQueryDesc* queryDescs) +{ + SLANG_GFX_API_FUNC; + List<IAccelerationStructure*> innerAS; + for (int i = 0; i < accelerationStructureCount; i++) + { + innerAS.add(getInnerObj(accelerationStructures[i])); + } + List<AccelerationStructureQueryDesc> innerQueryDescs; + innerQueryDescs.addRange(queryDescs, queryCount); + for (auto& innerQueryDesc : innerQueryDescs) + { + innerQueryDesc.queryPool = getInnerObj(innerQueryDesc.queryPool); + } + baseObject->queryAccelerationStructureProperties( + accelerationStructureCount, innerAS.getBuffer(), queryCount, innerQueryDescs.getBuffer()); +} + +void DebugRayTracingCommandEncoder::serializeAccelerationStructure( + DeviceAddress dest, + IAccelerationStructure* source) +{ + SLANG_GFX_API_FUNC; + baseObject->serializeAccelerationStructure(dest, getInnerObj(source)); +} + +void DebugRayTracingCommandEncoder::deserializeAccelerationStructure( + IAccelerationStructure* dest, + DeviceAddress source) +{ + SLANG_GFX_API_FUNC; + baseObject->deserializeAccelerationStructure(getInnerObj(dest), source); +} + +void DebugRayTracingCommandEncoder::memoryBarrier( + int count, + IAccelerationStructure* const* structures, + AccessFlag::Enum sourceAccess, + AccessFlag::Enum destAccess) +{ + SLANG_GFX_API_FUNC; + List<IAccelerationStructure*> innerAS; + for (int i = 0; i < count; i++) + { + innerAS.add(getInnerObj(structures[i])); + } + baseObject->memoryBarrier(count, innerAS.getBuffer(), sourceAccess, destAccess); +} + const ICommandQueue::Desc& DebugCommandQueue::getDesc() { SLANG_GFX_API_FUNC; @@ -794,6 +983,14 @@ void DebugCommandQueue::executeCommandBuffers(uint32_t count, ICommandBuffer* co "before submitting to a command queue.", cmdBufferImpl->uid); } + if (i > 0) + { + if (cmdBufferImpl->m_transientHeap != getDebugObj(commandBuffers[0])->m_transientHeap) + { + GFX_DIAGNOSE_ERROR("Command buffers passed to a single executeCommandBuffers " + "call must be allocated from the same transient heap."); + } + } } baseObject->executeCommandBuffers(count, innerCommandBuffers.getBuffer()); } @@ -810,6 +1007,7 @@ Result DebugTransientResourceHeap::createCommandBuffer(ICommandBuffer** outComma { SLANG_GFX_API_FUNC; RefPtr<DebugCommandBuffer> outObject = new DebugCommandBuffer(); + outObject->m_transientHeap = this; auto result = baseObject->createCommandBuffer(outObject->baseObject.writeRef()); if (SLANG_FAILED(result)) return result; @@ -1009,14 +1207,25 @@ Result DebugRootShaderObject::setSpecializationArgs( const slang::SpecializationArg* args, uint32_t count) { + SLANG_GFX_API_FUNC; + return baseObject->setSpecializationArgs(offset, args, count); } Result DebugQueryPool::getResult(SlangInt index, SlangInt count, uint64_t* data) { - if (index < 0 || index + count >= desc.count) + SLANG_GFX_API_FUNC; + + if (index < 0 || index + count > desc.count) GFX_DIAGNOSE_ERROR("index is out of bounds."); return baseObject->getResult(index, count, data); } +DeviceAddress DebugAccelerationStructure::getDeviceAddress() +{ + SLANG_GFX_API_FUNC; + + return baseObject->getDeviceAddress(); +} + } // namespace gfx diff --git a/tools/gfx/debug-layer.h b/tools/gfx/debug-layer.h index 6225260dd..8ab1146d0 100644 --- a/tools/gfx/debug-layer.h +++ b/tools/gfx/debug-layer.h @@ -23,6 +23,15 @@ public: Slang::ComPtr<TInterface> baseObject; }; +template <typename TInterface> +class UnownedDebugObject + : public TInterface + , public DebugObjectBase +{ +public: + TInterface* baseObject = nullptr; +}; + class DebugDevice : public DebugObject<IDevice> { public: @@ -57,6 +66,12 @@ public: IBufferResource* buffer, IResourceView::Desc const& desc, IResourceView** outView) override; + virtual SLANG_NO_THROW Result SLANG_MCALL getAccelerationStructurePrebuildInfo( + const IAccelerationStructure::BuildInputs& buildInputs, + IAccelerationStructure::PrebuildInfo* outPrebuildInfo) override; + virtual SLANG_NO_THROW Result SLANG_MCALL createAccelerationStructure( + const IAccelerationStructure::CreateDesc& desc, + IAccelerationStructure** outView) override; virtual SLANG_NO_THROW Result SLANG_MCALL createFramebufferLayout( IFramebufferLayout::Desc const& desc, IFramebufferLayout** outFrameBuffer) override; @@ -147,6 +162,16 @@ public: IResourceView* getInterface(const Slang::Guid& guid); }; +class DebugAccelerationStructure : public DebugObject<IAccelerationStructure> +{ +public: + SLANG_COM_OBJECT_IUNKNOWN_ALL; + +public: + IAccelerationStructure* getInterface(const Slang::Guid& guid); + virtual SLANG_NO_THROW DeviceAddress SLANG_MCALL getDeviceAddress() override; +}; + class DebugSamplerState : public DebugObject<ISamplerState> { public: @@ -228,16 +253,9 @@ public: class DebugCommandBuffer; -class DebugComputeCommandEncoder : public DebugObject<IComputeCommandEncoder> +class DebugComputeCommandEncoder : public UnownedDebugObject<IComputeCommandEncoder> { public: - SLANG_COM_OBJECT_IUNKNOWN_QUERY_INTERFACE; - -public: - IComputeCommandEncoder* getInterface(const Slang::Guid& guid); - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override; virtual SLANG_NO_THROW Result SLANG_MCALL bindPipeline(IPipelineState* state, IShaderObject** outRootShaderObject) override; @@ -249,16 +267,9 @@ public: bool isOpen = false; }; -class DebugRenderCommandEncoder : public DebugObject<IRenderCommandEncoder> +class DebugRenderCommandEncoder : public UnownedDebugObject<IRenderCommandEncoder> { public: - SLANG_COM_OBJECT_IUNKNOWN_QUERY_INTERFACE; - -public: - IRenderCommandEncoder* getInterface(const Slang::Guid& guid); - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override; virtual SLANG_NO_THROW Result SLANG_MCALL bindPipeline(IPipelineState* state, IShaderObject** outRootShaderObject) override; @@ -287,16 +298,9 @@ public: bool isOpen = false; }; -class DebugResourceCommandEncoder : public DebugObject<IResourceCommandEncoder> +class DebugResourceCommandEncoder : public UnownedDebugObject<IResourceCommandEncoder> { public: - SLANG_COM_OBJECT_IUNKNOWN_QUERY_INTERFACE; - -public: - IResourceCommandEncoder* getInterface(const Slang::Guid& guid); - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override; virtual SLANG_NO_THROW void SLANG_MCALL copyBuffer( IBufferResource* dst, @@ -313,15 +317,56 @@ public: bool isOpen = false; }; +class DebugRayTracingCommandEncoder : public UnownedDebugObject<IRayTracingCommandEncoder> +{ +public: + virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override; + virtual SLANG_NO_THROW void SLANG_MCALL + writeTimestamp(IQueryPool* pool, SlangInt index) override; + virtual SLANG_NO_THROW void SLANG_MCALL buildAccelerationStructure( + const IAccelerationStructure::BuildDesc& desc, + int propertyQueryCount, + AccelerationStructureQueryDesc* queryDescs) override; + virtual SLANG_NO_THROW void SLANG_MCALL copyAccelerationStructure( + IAccelerationStructure* dest, + IAccelerationStructure* src, + AccelerationStructureCopyMode mode) override; + virtual SLANG_NO_THROW void SLANG_MCALL queryAccelerationStructureProperties( + int accelerationStructureCount, + IAccelerationStructure* const* accelerationStructures, + int queryCount, + AccelerationStructureQueryDesc* queryDescs) override; + virtual SLANG_NO_THROW void SLANG_MCALL + serializeAccelerationStructure(DeviceAddress dest, IAccelerationStructure* source) override; + virtual SLANG_NO_THROW void SLANG_MCALL deserializeAccelerationStructure( + IAccelerationStructure* dest, + DeviceAddress source) override; + virtual SLANG_NO_THROW void memoryBarrier( + int count, + IAccelerationStructure* const* structures, + AccessFlag::Enum sourceAccess, + AccessFlag::Enum destAccess) override; + +public: + DebugCommandBuffer* commandBuffer; + bool isOpen = false; +}; + +class DebugTransientResourceHeap; + class DebugCommandBuffer : public DebugObject<ICommandBuffer> { public: SLANG_COM_OBJECT_IUNKNOWN_ALL; +public: + DebugTransientResourceHeap* m_transientHeap; + private: DebugRenderCommandEncoder m_renderCommandEncoder; DebugComputeCommandEncoder m_computeCommandEncoder; DebugResourceCommandEncoder m_resourceCommandEncoder; + DebugRayTracingCommandEncoder m_rayTracingCommandEncoder; public: DebugCommandBuffer(); @@ -334,6 +379,8 @@ public: encodeComputeCommands(IComputeCommandEncoder** outEncoder) override; virtual SLANG_NO_THROW void SLANG_MCALL encodeResourceCommands(IResourceCommandEncoder** outEncoder) override; + virtual SLANG_NO_THROW void SLANG_MCALL + encodeRayTracingCommands(IRayTracingCommandEncoder** outEncoder) override; virtual SLANG_NO_THROW void SLANG_MCALL close() override; private: diff --git a/tools/gfx/immediate-renderer-base.cpp b/tools/gfx/immediate-renderer-base.cpp index 8fffbfdfa..eae6c82b0 100644 --- a/tools/gfx/immediate-renderer-base.cpp +++ b/tools/gfx/immediate-renderer-base.cpp @@ -50,22 +50,6 @@ public: : public IRenderCommandEncoder { public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IRenderCommandEncoder) - { - *outObject = static_cast<IRenderCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - - public: CommandWriter* m_writer; CommandBufferImpl* m_commandBuffer; virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override {} @@ -188,22 +172,6 @@ public: : public IComputeCommandEncoder { public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IComputeCommandEncoder) - { - *outObject = static_cast<IComputeCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - - public: CommandWriter* m_writer; CommandBufferImpl* m_commandBuffer; @@ -252,22 +220,6 @@ public: : public IResourceCommandEncoder { public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IResourceCommandEncoder) - { - *outObject = static_cast<IResourceCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - - public: CommandWriter* m_writer; void init(CommandBufferImpl* cmdBuffer) @@ -307,6 +259,12 @@ public: *outEncoder = &m_resourceCommandEncoder; } + virtual SLANG_NO_THROW void SLANG_MCALL + encodeRayTracingCommands(IRayTracingCommandEncoder** outEncoder) override + { + *outEncoder = nullptr; + } + virtual SLANG_NO_THROW void SLANG_MCALL close() override { } void execute() diff --git a/tools/gfx/open-gl/render-gl.cpp b/tools/gfx/open-gl/render-gl.cpp index 81a1a6ccd..92f8680d7 100644 --- a/tools/gfx/open-gl/render-gl.cpp +++ b/tools/gfx/open-gl/render-gl.cpp @@ -1257,7 +1257,7 @@ public: bufferDesc.defaultState = ResourceState::ConstantBuffer; bufferDesc.allowedStates = ResourceStateSet(ResourceState::ConstantBuffer, ResourceState::CopyDestination); - bufferDesc.cpuAccessFlags |= IResource::AccessFlag::Write; + bufferDesc.cpuAccessFlags |= AccessFlag::Write; SLANG_RETURN_ON_FAIL( device->createBufferResource(bufferDesc, nullptr, bufferResourcePtr.writeRef())); m_ordinaryDataBuffer = static_cast<BufferResourceImpl*>(bufferResourcePtr.get()); diff --git a/tools/gfx/renderer-shared.cpp b/tools/gfx/renderer-shared.cpp index 9d6c85309..c88081547 100644 --- a/tools/gfx/renderer-shared.cpp +++ b/tools/gfx/renderer-shared.cpp @@ -25,13 +25,11 @@ const Slang::Guid GfxGUID::IID_IDevice = SLANG_UUID_IDevice; const Slang::Guid GfxGUID::IID_IShaderObject = SLANG_UUID_IShaderObject; const Slang::Guid GfxGUID::IID_IRenderPassLayout = SLANG_UUID_IRenderPassLayout; -const Slang::Guid GfxGUID::IID_ICommandEncoder = SLANG_UUID_ICommandEncoder; -const Slang::Guid GfxGUID::IID_IRenderCommandEncoder = SLANG_UUID_IRenderCommandEncoder; -const Slang::Guid GfxGUID::IID_IComputeCommandEncoder = SLANG_UUID_IComputeCommandEncoder; -const Slang::Guid GfxGUID::IID_IResourceCommandEncoder = SLANG_UUID_IResourceCommandEncoder; +const Slang::Guid GfxGUID::IID_IRayTracingCommandEncoder = SLANG_UUID_IRayTracingCommandEncoder; const Slang::Guid GfxGUID::IID_ICommandBuffer = SLANG_UUID_ICommandBuffer; const Slang::Guid GfxGUID::IID_ICommandQueue = SLANG_UUID_ICommandQueue; const Slang::Guid GfxGUID::IID_IQueryPool = SLANG_UUID_IQueryPool; +const Slang::Guid GfxGUID::IID_IAccelerationStructure = SLANG_UUID_IAccelerationStructure; StageType translateStage(SlangStage slangStage) @@ -119,6 +117,14 @@ IResourceView* ResourceViewBase::getInterface(const Guid& guid) return nullptr; } +IAccelerationStructure* AccelerationStructureBase::getInterface(const Slang::Guid& guid) +{ + if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IResourceView || + guid == GfxGUID::IID_IAccelerationStructure) + return static_cast<IAccelerationStructure*>(this); + return nullptr; +} + IShaderObject* ShaderObjectBase::getInterface(const Guid& guid) { if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IShaderObject) @@ -279,6 +285,24 @@ SLANG_NO_THROW Result SLANG_MCALL RendererBase::createShaderObject( return createShaderObject(shaderObjectLayout, outObject); } +Result RendererBase::getAccelerationStructurePrebuildInfo( + const IAccelerationStructure::BuildInputs& buildInputs, + IAccelerationStructure::PrebuildInfo* outPrebuildInfo) +{ + SLANG_UNUSED(buildInputs); + SLANG_UNUSED(outPrebuildInfo); + return SLANG_E_NOT_AVAILABLE; +} + +Result RendererBase::createAccelerationStructure( + const IAccelerationStructure::CreateDesc& desc, + IAccelerationStructure** outView) +{ + SLANG_UNUSED(desc); + SLANG_UNUSED(outView); + return SLANG_E_NOT_AVAILABLE; +} + Result RendererBase::getShaderObjectLayout( slang::TypeReflection* type, ShaderObjectContainerType container, diff --git a/tools/gfx/renderer-shared.h b/tools/gfx/renderer-shared.h index 127987726..e3580b9b6 100644 --- a/tools/gfx/renderer-shared.h +++ b/tools/gfx/renderer-shared.h @@ -33,9 +33,11 @@ struct GfxGUID static const Slang::Guid IID_IRenderCommandEncoder; static const Slang::Guid IID_IComputeCommandEncoder; static const Slang::Guid IID_IResourceCommandEncoder; + static const Slang::Guid IID_IRayTracingCommandEncoder; static const Slang::Guid IID_ICommandBuffer; static const Slang::Guid IID_ICommandQueue; static const Slang::Guid IID_IQueryPool; + static const Slang::Guid IID_IAccelerationStructure; }; // We use a `BreakableReference` to avoid the cyclic reference situation in gfx implementation. @@ -252,6 +254,15 @@ public: IResourceView* getInterface(const Slang::Guid& guid); }; +class AccelerationStructureBase + : public IAccelerationStructure + , public Slang::ComObject +{ +public: + SLANG_COM_OBJECT_IUNKNOWN_ALL + IAccelerationStructure* getInterface(const Slang::Guid& guid); +}; + class RendererBase; typedef uint32_t ShaderComponentID; @@ -1061,6 +1072,18 @@ public: ShaderObjectContainerType containerType, IShaderObject** outObject) SLANG_OVERRIDE; + // Provides a default implementation that returns SLANG_E_NOT_AVAILABLE for platforms + // without ray tracing support. + virtual SLANG_NO_THROW Result SLANG_MCALL getAccelerationStructurePrebuildInfo( + const IAccelerationStructure::BuildInputs& buildInputs, + IAccelerationStructure::PrebuildInfo* outPrebuildInfo) override; + + // Provides a default implementation that returns SLANG_E_NOT_AVAILABLE for platforms + // without ray tracing support. + virtual SLANG_NO_THROW Result SLANG_MCALL createAccelerationStructure( + const IAccelerationStructure::CreateDesc& desc, + IAccelerationStructure** outView) override; + Result getShaderObjectLayout( slang::TypeReflection* type, ShaderObjectContainerType container, diff --git a/tools/gfx/simple-transient-resource-heap.h b/tools/gfx/simple-transient-resource-heap.h index d8ab3517d..4340d49df 100644 --- a/tools/gfx/simple-transient-resource-heap.h +++ b/tools/gfx/simple-transient-resource-heap.h @@ -35,7 +35,7 @@ public: bufferDesc.allowedStates = ResourceStateSet(ResourceState::ConstantBuffer, ResourceState::CopyDestination); bufferDesc.defaultState = ResourceState::ConstantBuffer; bufferDesc.sizeInBytes = desc.constantBufferSize; - bufferDesc.cpuAccessFlags = IResource::AccessFlag::Write; + bufferDesc.cpuAccessFlags = AccessFlag::Write; SLANG_RETURN_ON_FAIL( device->createBufferResource(bufferDesc, nullptr, m_constantBuffer.writeRef())); return SLANG_OK; diff --git a/tools/gfx/transient-resource-heap-base.h b/tools/gfx/transient-resource-heap-base.h index ef8a61616..f11a91c14 100644 --- a/tools/gfx/transient-resource-heap-base.h +++ b/tools/gfx/transient-resource-heap-base.h @@ -45,7 +45,7 @@ public: bufferDesc.allowedStates = ResourceStateSet(ResourceState::ConstantBuffer, ResourceState::CopyDestination); bufferDesc.sizeInBytes = desc.constantBufferSize; - bufferDesc.cpuAccessFlags = IResource::AccessFlag::Write; + bufferDesc.cpuAccessFlags = AccessFlag::Write; SLANG_RETURN_ON_FAIL( m_device->createBufferResource(bufferDesc, nullptr, bufferPtr.writeRef())); m_constantBuffers.add(static_cast<TBufferResource*>(bufferPtr.get())); @@ -90,7 +90,7 @@ public: bufferDesc.defaultState = ResourceState::ConstantBuffer; bufferDesc.allowedStates = ResourceStateSet(ResourceState::ConstantBuffer, ResourceState::CopyDestination); - bufferDesc.cpuAccessFlags |= IResource::AccessFlag::Write; + bufferDesc.cpuAccessFlags |= AccessFlag::Write; size_t lastConstantBufferSize = 0; if (m_constantBuffers.getCount()) { diff --git a/tools/gfx/vulkan/render-vk.cpp b/tools/gfx/vulkan/render-vk.cpp index 97a527ca6..10cc7aae5 100644 --- a/tools/gfx/vulkan/render-vk.cpp +++ b/tools/gfx/vulkan/render-vk.cpp @@ -120,6 +120,15 @@ public: size_t offset, size_t size, ISlangBlob** outBlob) override; + + virtual SLANG_NO_THROW Result SLANG_MCALL getAccelerationStructurePrebuildInfo( + const IAccelerationStructure::BuildInputs& buildInputs, + IAccelerationStructure::PrebuildInfo* outPrebuildInfo) override; + + virtual SLANG_NO_THROW Result SLANG_MCALL createAccelerationStructure( + const IAccelerationStructure::CreateDesc& desc, + IAccelerationStructure** outView) override; + void waitForGpu(); virtual SLANG_NO_THROW const DeviceInfo& SLANG_MCALL getDeviceInfo() const override { @@ -128,6 +137,40 @@ public: /// Dtor ~VKDevice(); +public: + // Float16 features + VkPhysicalDeviceFloat16Int8FeaturesKHR float16Features = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FLOAT16_INT8_FEATURES_KHR}; + // 16 bit storage features + VkPhysicalDevice16BitStorageFeatures storage16BitFeatures = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_16BIT_STORAGE_FEATURES_KHR}; + // AtomicInt64 features + VkPhysicalDeviceShaderAtomicInt64FeaturesKHR atomicInt64Features = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_ATOMIC_INT64_FEATURES_KHR}; + // Atomic Float features + VkPhysicalDeviceShaderAtomicFloatFeaturesEXT atomicFloatFeatures = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_ATOMIC_FLOAT_FEATURES_EXT}; + // Timeline Semaphore features + VkPhysicalDeviceTimelineSemaphoreFeatures timelineFeatures = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_TIMELINE_SEMAPHORE_FEATURES}; + // Extended dynamic state features + VkPhysicalDeviceExtendedDynamicStateFeaturesEXT extendedDynamicStateFeatures = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTENDED_DYNAMIC_STATE_FEATURES_EXT}; + // Subgroup extended type features + VkPhysicalDeviceShaderSubgroupExtendedTypesFeatures shaderSubgroupExtendedTypeFeatures = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_SUBGROUP_EXTENDED_TYPES_FEATURES}; + // Acceleration structure features + VkPhysicalDeviceAccelerationStructureFeaturesKHR accelerationStructureFeatures = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ACCELERATION_STRUCTURE_FEATURES_KHR}; + // Ray query (inline ray-tracing) features + VkPhysicalDeviceRayQueryFeaturesKHR rayQueryFeatures = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_RAY_QUERY_FEATURES_KHR}; + // Buffer device address features + VkPhysicalDeviceBufferDeviceAddressFeatures bufferDeviceAddressFeatures = { + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_BUFFER_DEVICE_ADDRESS_FEATURES}; + +public: + class Buffer { public: @@ -303,6 +346,28 @@ public: VkDeviceSize size; }; + class AccelerationStructureImpl : public AccelerationStructureBase + { + public: + VkAccelerationStructureKHR m_vkHandle = VK_NULL_HANDLE; + RefPtr<BufferResourceImpl> m_buffer; + VkDeviceSize m_offset; + VkDeviceSize m_size; + RefPtr<VKDevice> m_device; + public: + virtual SLANG_NO_THROW DeviceAddress SLANG_MCALL getDeviceAddress() override + { + return m_buffer->getDeviceAddress() + m_offset; + } + ~AccelerationStructureImpl() + { + if (m_device) + { + m_device->m_api.vkDestroyAccelerationStructureKHR(m_device->m_api.m_device, m_vkHandle, nullptr); + } + } + }; + class FramebufferLayoutImpl : public FramebufferLayoutBase { public: @@ -2101,7 +2166,6 @@ public: class PipelineCommandEncoder : public RefObject { public: - bool m_isOpen = false; CommandBufferImpl* m_commandBuffer; VkCommandBuffer m_vkCommandBuffer; VkCommandBuffer m_vkPreCommandBuffer = VK_NULL_HANDLE; @@ -2130,7 +2194,6 @@ public: void endEncodingImpl() { - m_isOpen = false; for (auto& pipeline : m_boundPipelines) pipeline = VK_NULL_HANDLE; } @@ -2668,6 +2731,36 @@ public: } } + static void writeAccelerationStructureDescriptor( + RootBindingContext& context, + BindingOffset const& offset, + VkDescriptorType descriptorType, + ArrayView<RefPtr<ResourceViewImpl>> resourceViews) + { + auto descriptorSet = context.descriptorSets[offset.bindingSet]; + + Index count = resourceViews.getCount(); + for (Index i = 0; i < count; ++i) + { + auto accelerationStructure = static_cast<AccelerationStructureImpl*>( + static_cast<IResourceView*>(resourceViews[i].Ptr())); + + VkWriteDescriptorSetAccelerationStructureKHR writeAS = {}; + writeAS.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET_ACCELERATION_STRUCTURE_KHR; + writeAS.accelerationStructureCount = 1; + writeAS.pAccelerationStructures = &accelerationStructure->m_vkHandle; + VkWriteDescriptorSet write = {}; + write.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET; + write.descriptorCount = 1; + write.descriptorType = descriptorType; + write.dstArrayElement = uint32_t(i); + write.dstBinding = offset.binding; + write.dstSet = descriptorSet; + write.pNext = &writeAS; + writeDescriptor(context, write); + } + } + static void writeTextureDescriptor( RootBindingContext& context, BindingOffset const& offset, @@ -2881,7 +2974,15 @@ public: VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER, m_resourceViews.getArrayView(baseIndex, count)); break; - + case slang::BindingType::RayTracingAccelerationStructure: + rangeOffset.bindingSet += bindingRangeInfo.setOffset; + rangeOffset.binding += bindingRangeInfo.bindingOffset; + writeAccelerationStructureDescriptor( + context, + rangeOffset, + VK_DESCRIPTOR_TYPE_ACCELERATION_STRUCTURE_KHR, + m_resourceViews.getArrayView(baseIndex, count)); + break; case slang::BindingType::VaryingInput: case slang::BindingType::VaryingOutput: break; @@ -3473,7 +3574,6 @@ public: VkCommandBuffer m_commandBuffer; VkCommandBuffer m_preCommandBuffer = VK_NULL_HANDLE; VkCommandPool m_pool; - VkFence m_fence; VKDevice* m_renderer; BreakableReference<TransientResourceHeapImpl> m_transientHeap; bool m_isPreCommandBufferEmpty = true; @@ -3485,13 +3585,11 @@ public: Result init( VKDevice* renderer, VkCommandPool pool, - VkFence fence, TransientResourceHeapImpl* transientHeap) { m_renderer = renderer; m_transientHeap = transientHeap; m_pool = pool; - m_fence = fence; auto& api = renderer->m_api; VkCommandBufferAllocateInfo allocInfo = {}; @@ -3576,21 +3674,6 @@ public: VkIndexType m_boundIndexFormat; public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IRenderCommandEncoder) - { - *outObject = static_cast<IRenderCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - void beginPass(IRenderPassLayout* renderPass, IFramebuffer* framebuffer) { FramebufferImpl* framebufferImpl = static_cast<FramebufferImpl*>(framebuffer); @@ -3610,7 +3693,6 @@ public: beginInfo.pClearValues = framebufferImpl->m_clearValues; auto& api = *m_api; api.vkCmdBeginRenderPass(m_vkCommandBuffer, &beginInfo, VK_SUBPASS_CONTENTS_INLINE); - m_isOpen = true; } virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override @@ -3825,7 +3907,6 @@ public: m_renderCommandEncoder = new RenderCommandEncoder(); m_renderCommandEncoder->init(this); } - assert(!m_renderCommandEncoder->m_isOpen); m_renderCommandEncoder->beginPass(renderPass, framebuffer); *outEncoder = m_renderCommandEncoder.Ptr(); } @@ -3835,21 +3916,6 @@ public: , public PipelineCommandEncoder { public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IComputeCommandEncoder) - { - *outObject = static_cast<IComputeCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - public: virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override { endEncodingImpl(); @@ -3893,7 +3959,6 @@ public: m_computeCommandEncoder = new ComputeCommandEncoder(); m_computeCommandEncoder->init(this); } - assert(!m_computeCommandEncoder->m_isOpen); *outEncoder = m_computeCommandEncoder.Ptr(); } @@ -3904,21 +3969,6 @@ public: public: CommandBufferImpl* m_commandBuffer; public: - virtual SLANG_NO_THROW SlangResult SLANG_MCALL - queryInterface(SlangUUID const& uuid, void** outObject) override - { - if (uuid == GfxGUID::IID_ISlangUnknown || uuid == GfxGUID::IID_ICommandEncoder || - uuid == GfxGUID::IID_IResourceCommandEncoder) - { - *outObject = static_cast<IResourceCommandEncoder*>(this); - return SLANG_OK; - } - *outObject = nullptr; - return SLANG_E_NO_INTERFACE; - } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL addRef() override { return 1; } - virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() override { return 1; } - public: virtual SLANG_NO_THROW void SLANG_MCALL copyBuffer( IBufferResource* dst, size_t dstOffset, @@ -4006,6 +4056,264 @@ public: *outEncoder = m_resourceCommandEncoder.Ptr(); } + class RayTracingCommandEncoder + : public IRayTracingCommandEncoder + , public RefObject + { + public: + CommandBufferImpl* m_commandBuffer; + + public: + void init(CommandBufferImpl* commandBuffer) { m_commandBuffer = commandBuffer; } + + inline VkAccessFlags translateAccelerationStructureAccessFlag(AccessFlag::Enum access) + { + VkAccessFlags result = 0; + if (access & AccessFlag::Read) + result |= VK_ACCESS_ACCELERATION_STRUCTURE_READ_BIT_KHR | + VK_ACCESS_SHADER_READ_BIT | VK_ACCESS_TRANSFER_READ_BIT; + if (access & AccessFlag::Write) + result |= VK_ACCESS_ACCELERATION_STRUCTURE_WRITE_BIT_KHR; + return result; + } + + inline void _memoryBarrier( + int count, + IAccelerationStructure* const* structures, + AccessFlag::Enum srcAccess, + AccessFlag::Enum destAccess) + { + ShortList<VkBufferMemoryBarrier> memBarriers; + memBarriers.setCount(count); + for (int i = 0; i < count; i++) + { + memBarriers[i].sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER; + memBarriers[i].pNext = nullptr; + memBarriers[i].dstAccessMask = + translateAccelerationStructureAccessFlag(destAccess); + memBarriers[i].srcAccessMask = + translateAccelerationStructureAccessFlag(srcAccess); + memBarriers[i].srcQueueFamilyIndex = + m_commandBuffer->m_renderer->m_queueFamilyIndex; + memBarriers[i].dstQueueFamilyIndex = + m_commandBuffer->m_renderer->m_queueFamilyIndex; + + auto asImpl = static_cast<AccelerationStructureImpl*>(structures[i]); + memBarriers[i].buffer = asImpl->m_buffer->m_buffer.m_buffer; + memBarriers[i].offset = asImpl->m_offset; + memBarriers[i].size = asImpl->m_size; + } + m_commandBuffer->m_renderer->m_api.vkCmdPipelineBarrier( + m_commandBuffer->m_commandBuffer, + VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR | + VK_PIPELINE_STAGE_BOTTOM_OF_PIPE_BIT, + VK_PIPELINE_STAGE_ACCELERATION_STRUCTURE_BUILD_BIT_KHR | + VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT | VK_PIPELINE_STAGE_TRANSFER_BIT, + 0, + 0, + nullptr, + (uint32_t)memBarriers.getCount(), + memBarriers.getArrayView().getBuffer(), + 0, + nullptr); + } + + inline void _queryAccelerationStructureProperties( + int accelerationStructureCount, + IAccelerationStructure* const* accelerationStructures, + int queryCount, + AccelerationStructureQueryDesc* queryDescs) + { + ShortList<VkAccelerationStructureKHR> vkHandles; + vkHandles.setCount(accelerationStructureCount); + for (int i = 0; i < accelerationStructureCount; i++) + { + vkHandles[i] = + static_cast<AccelerationStructureImpl*>(accelerationStructures[i]) + ->m_vkHandle; + } + auto vkHandlesView = vkHandles.getArrayView(); + for (int i = 0; i < queryCount; i++) + { + VkQueryType queryType; + switch (queryDescs[i].queryType) + { + case QueryType::AccelerationStructureCompactedSize: + queryType = VK_QUERY_TYPE_ACCELERATION_STRUCTURE_COMPACTED_SIZE_KHR; + break; + case QueryType::AccelerationStructureSerializedSize: + queryType = VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_SIZE_KHR; + break; + default: + getDebugCallback()->handleMessage(DebugMessageType::Error, DebugMessageSource::Layer, + "Invalid query type for use in queryAccelerationStructureProperties."); + return; + } + auto queryPool = static_cast<QueryPoolImpl*>(queryDescs[i].queryPool)->m_pool; + m_commandBuffer->m_renderer->m_api.vkCmdResetQueryPool( + m_commandBuffer->m_commandBuffer, + queryPool, + (uint32_t)queryDescs[i].firstQueryIndex, + 1); + m_commandBuffer->m_renderer->m_api + .vkCmdWriteAccelerationStructuresPropertiesKHR( + m_commandBuffer->m_commandBuffer, + accelerationStructureCount, + vkHandlesView.getBuffer(), + queryType, + queryPool, + queryDescs[i].firstQueryIndex); + } + } + + virtual SLANG_NO_THROW void SLANG_MCALL buildAccelerationStructure( + const IAccelerationStructure::BuildDesc& desc, + int propertyQueryCount, + AccelerationStructureQueryDesc* queryDescs) override + { + AccelerationStructureBuildGeometryInfoBuilder geomInfoBuilder; + if (geomInfoBuilder.build(desc.inputs, getDebugCallback()) != SLANG_OK) + return; + + if (desc.dest) + { + geomInfoBuilder.buildInfo.dstAccelerationStructure = + static_cast<AccelerationStructureImpl*>(desc.dest)->m_vkHandle; + } + if (desc.source) + { + geomInfoBuilder.buildInfo.srcAccelerationStructure = + static_cast<AccelerationStructureImpl*>(desc.source)->m_vkHandle; + } + geomInfoBuilder.buildInfo.scratchData.deviceAddress = desc.scratchData; + + List<VkAccelerationStructureBuildRangeInfoKHR> rangeInfos; + rangeInfos.setCount(geomInfoBuilder.primitiveCounts.getCount()); + for (Index i = 0; i < geomInfoBuilder.primitiveCounts.getCount(); i++) + { + auto& rangeInfo = rangeInfos[i]; + rangeInfo.primitiveCount = geomInfoBuilder.primitiveCounts[i]; + rangeInfo.firstVertex = 0; + rangeInfo.primitiveOffset = 0; + rangeInfo.transformOffset = 0; + } + + auto rangeInfoPtr = rangeInfos.getBuffer(); + m_commandBuffer->m_renderer->m_api.vkCmdBuildAccelerationStructuresKHR( + m_commandBuffer->m_commandBuffer, 1, &geomInfoBuilder.buildInfo, &rangeInfoPtr); + + if (propertyQueryCount) + { + _memoryBarrier(1, &desc.dest, AccessFlag::Write, AccessFlag::Read); + _queryAccelerationStructureProperties( + 1, &desc.dest, propertyQueryCount, queryDescs); + } + } + + virtual SLANG_NO_THROW void SLANG_MCALL copyAccelerationStructure( + IAccelerationStructure* dest, + IAccelerationStructure* src, + AccelerationStructureCopyMode mode) override + { + VkCopyAccelerationStructureInfoKHR copyInfo = { + VK_STRUCTURE_TYPE_COPY_ACCELERATION_STRUCTURE_INFO_KHR}; + copyInfo.src = static_cast<AccelerationStructureImpl*>(src)->m_vkHandle; + copyInfo.dst = static_cast<AccelerationStructureImpl*>(dest)->m_vkHandle; + switch (mode) + { + case AccelerationStructureCopyMode::Clone: + copyInfo.mode = VK_COPY_ACCELERATION_STRUCTURE_MODE_CLONE_KHR; + break; + case AccelerationStructureCopyMode::Compact: + copyInfo.mode = VK_COPY_ACCELERATION_STRUCTURE_MODE_COMPACT_KHR; + break; + default: + getDebugCallback()->handleMessage( + DebugMessageType::Error, + DebugMessageSource::Layer, + "Unsupported AccelerationStructureCopyMode."); + return; + } + m_commandBuffer->m_renderer->m_api.vkCmdCopyAccelerationStructureKHR( + m_commandBuffer->m_commandBuffer, ©Info); + } + + virtual SLANG_NO_THROW void SLANG_MCALL queryAccelerationStructureProperties( + int accelerationStructureCount, + IAccelerationStructure* const* accelerationStructures, + int queryCount, + AccelerationStructureQueryDesc* queryDescs) override + { + _queryAccelerationStructureProperties( + accelerationStructureCount, accelerationStructures, queryCount, queryDescs); + } + + virtual SLANG_NO_THROW void SLANG_MCALL serializeAccelerationStructure( + DeviceAddress dest, + IAccelerationStructure* source) override + { + VkCopyAccelerationStructureToMemoryInfoKHR copyInfo = { + VK_STRUCTURE_TYPE_COPY_ACCELERATION_STRUCTURE_TO_MEMORY_INFO_KHR}; + copyInfo.src = static_cast<AccelerationStructureImpl*>(source)->m_vkHandle; + copyInfo.dst.deviceAddress = dest; + copyInfo.mode = VK_COPY_ACCELERATION_STRUCTURE_MODE_SERIALIZE_KHR; + m_commandBuffer->m_renderer->m_api.vkCmdCopyAccelerationStructureToMemoryKHR( + m_commandBuffer->m_commandBuffer, ©Info); + } + + virtual SLANG_NO_THROW void SLANG_MCALL deserializeAccelerationStructure( + IAccelerationStructure* dest, + DeviceAddress source) override + { + VkCopyMemoryToAccelerationStructureInfoKHR copyInfo = { + VK_STRUCTURE_TYPE_COPY_MEMORY_TO_ACCELERATION_STRUCTURE_INFO_KHR}; + copyInfo.src.deviceAddress = source; + copyInfo.dst = static_cast<AccelerationStructureImpl*>(dest)->m_vkHandle; + copyInfo.mode = VK_COPY_ACCELERATION_STRUCTURE_MODE_DESERIALIZE_KHR; + m_commandBuffer->m_renderer->m_api.vkCmdCopyMemoryToAccelerationStructureKHR( + m_commandBuffer->m_commandBuffer, ©Info); + } + + virtual SLANG_NO_THROW void memoryBarrier( + int count, + IAccelerationStructure* const* structures, + AccessFlag::Enum srcAccess, + AccessFlag::Enum destAccess) override + { + _memoryBarrier(count, structures, srcAccess, destAccess); + } + + virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() override + { + } + + virtual SLANG_NO_THROW void SLANG_MCALL + writeTimestamp(IQueryPool* queryPool, SlangInt index) override + { + _writeTimestamp( + &m_commandBuffer->m_renderer->m_api, + m_commandBuffer->m_commandBuffer, + queryPool, + index); + } + }; + + RefPtr<RayTracingCommandEncoder> m_rayTracingCommandEncoder; + + virtual SLANG_NO_THROW void SLANG_MCALL + encodeRayTracingCommands(IRayTracingCommandEncoder** outEncoder) override + { + if (!m_rayTracingCommandEncoder) + { + if (m_renderer->m_api.vkCmdBuildAccelerationStructuresKHR) + { + m_rayTracingCommandEncoder = new RayTracingCommandEncoder(); + m_rayTracingCommandEncoder->init(this); + } + } + *outEncoder = m_rayTracingCommandEncoder.Ptr(); + } + virtual SLANG_NO_THROW void SLANG_MCALL close() override { auto& vkAPI = m_renderer->m_api; @@ -4052,9 +4360,9 @@ public: RefPtr<VKDevice> m_renderer; VkQueue m_queue; uint32_t m_queueFamilyIndex; - VkSemaphore m_pendingWaitSemaphore = VK_NULL_HANDLE; + VkSemaphore m_pendingWaitSemaphores[2] = {VK_NULL_HANDLE, VK_NULL_HANDLE}; List<VkCommandBuffer> m_submitCommandBuffers; - static const int kSemaphoreCount = 2; + static const int kSemaphoreCount = 32; uint32_t m_currentSemaphoreIndex; VkSemaphore m_semaphores[kSemaphoreCount]; ~CommandQueueImpl() @@ -4114,26 +4422,37 @@ public: auto vkCmdBuf = cmdBufImpl->m_commandBuffer; m_submitCommandBuffers.add(vkCmdBuf); } - VkSemaphore waitSemaphore = m_pendingWaitSemaphore; VkSemaphore signalSemaphore = m_semaphores[m_currentSemaphoreIndex]; VkSubmitInfo submitInfo = {}; submitInfo.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; - VkPipelineStageFlags stageFlag = VK_PIPELINE_STAGE_BOTTOM_OF_PIPE_BIT; - submitInfo.pWaitDstStageMask = &stageFlag; + VkPipelineStageFlags stageFlag[] = { + VK_PIPELINE_STAGE_BOTTOM_OF_PIPE_BIT, VK_PIPELINE_STAGE_BOTTOM_OF_PIPE_BIT}; + submitInfo.pWaitDstStageMask = stageFlag; submitInfo.commandBufferCount = (uint32_t)m_submitCommandBuffers.getCount(); submitInfo.pCommandBuffers = m_submitCommandBuffers.getBuffer(); - if (m_pendingWaitSemaphore != VK_NULL_HANDLE) + Array<VkSemaphore, 2> waitSemaphores; + for (auto s : m_pendingWaitSemaphores) + { + if (s != VK_NULL_HANDLE) + { + waitSemaphores.add(s); + } + } + submitInfo.waitSemaphoreCount = (uint32_t)waitSemaphores.getCount(); + if (submitInfo.waitSemaphoreCount) { - submitInfo.waitSemaphoreCount = 1; - submitInfo.pWaitSemaphores = &waitSemaphore; + submitInfo.pWaitSemaphores = waitSemaphores.getBuffer(); } submitInfo.signalSemaphoreCount = 1; submitInfo.pSignalSemaphores = &signalSemaphore; - auto fence = static_cast<CommandBufferImpl*>(commandBuffers[0])->m_fence; + auto commandBufferImpl = static_cast<CommandBufferImpl*>(commandBuffers[0]); + auto fence = commandBufferImpl->m_transientHeap->getCurrentFence(); vkAPI.vkResetFences(vkAPI.m_device, 1, &fence); vkAPI.vkQueueSubmit(m_queue, 1, &submitInfo, fence); - m_pendingWaitSemaphore = signalSemaphore; + m_pendingWaitSemaphores[0] = signalSemaphore; + m_pendingWaitSemaphores[1] = VK_NULL_HANDLE; + commandBufferImpl->m_transientHeap->advanceFence(); m_currentSemaphoreIndex++; m_currentSemaphoreIndex = m_currentSemaphoreIndex % kSemaphoreCount; @@ -4149,16 +4468,37 @@ public: public: VkCommandPool m_commandPool; DescriptorSetAllocator m_descSetAllocator; - VkFence m_fence; + List<VkFence> m_fences; + Index m_fenceIndex = -1; List<RefPtr<CommandBufferImpl>> m_commandBufferPool; uint32_t m_commandBufferAllocId = 0; + VkFence getCurrentFence() + { + return m_fences[m_fenceIndex]; + } + void advanceFence() + { + m_fenceIndex++; + if (m_fenceIndex >= m_fences.getCount()) + { + m_fences.setCount(m_fenceIndex + 1); + VkFenceCreateInfo fenceCreateInfo = {}; + fenceCreateInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceCreateInfo.flags = VK_FENCE_CREATE_SIGNALED_BIT; + m_device->m_api.vkCreateFence( + m_device->m_api.m_device, &fenceCreateInfo, nullptr, &m_fences[m_fenceIndex]); + } + } Result init(const ITransientResourceHeap::Desc& desc, VKDevice* device); ~TransientResourceHeapImpl() { m_commandBufferPool = decltype(m_commandBufferPool)(); m_device->m_api.vkDestroyCommandPool(m_device->m_api.m_device, m_commandPool, nullptr); - m_device->m_api.vkDestroyFence(m_device->m_api.m_device, m_fence, nullptr); + for (auto fence : m_fences) + { + m_device->m_api.vkDestroyFence(m_device->m_api.m_device, fence, nullptr); + } m_descSetAllocator.close(); } public: @@ -4191,6 +4531,12 @@ public: case QueryType::Timestamp: createInfo.queryType = VK_QUERY_TYPE_TIMESTAMP; break; + case QueryType::AccelerationStructureCompactedSize: + createInfo.queryType = VK_QUERY_TYPE_ACCELERATION_STRUCTURE_COMPACTED_SIZE_KHR; + break; + case QueryType::AccelerationStructureSerializedSize: + createInfo.queryType = VK_QUERY_TYPE_ACCELERATION_STRUCTURE_SERIALIZATION_SIZE_KHR; + break; default: return SLANG_E_INVALID_ARG; } @@ -4554,13 +4900,22 @@ public: presentInfo.swapchainCount = 1; presentInfo.pSwapchains = &m_swapChain; presentInfo.pImageIndices = swapChainIndices; - if (m_queue->m_pendingWaitSemaphore != VK_NULL_HANDLE) + Array<VkSemaphore, 2> waitSemaphores; + for (auto s : m_queue->m_pendingWaitSemaphores) { - presentInfo.waitSemaphoreCount = 1; - presentInfo.pWaitSemaphores = &m_queue->m_pendingWaitSemaphore; + if (s != VK_NULL_HANDLE) + { + waitSemaphores.add(s); + } + } + presentInfo.waitSemaphoreCount = (uint32_t)waitSemaphores.getCount(); + if (presentInfo.waitSemaphoreCount) + { + presentInfo.pWaitSemaphores = waitSemaphores.getBuffer(); } m_api->vkQueuePresentKHR(m_queue->m_queue, &presentInfo); - m_queue->m_pendingWaitSemaphore = VK_NULL_HANDLE; + m_queue->m_pendingWaitSemaphores[0] = VK_NULL_HANDLE; + m_queue->m_pendingWaitSemaphores[1] = VK_NULL_HANDLE; return SLANG_OK; } virtual SLANG_NO_THROW int SLANG_MCALL acquireNextImage() override @@ -4584,7 +4939,7 @@ public: return m_currentImageIndex; } // Make the queue's next submit wait on `m_nextImageSemaphore`. - m_queue->m_pendingWaitSemaphore = m_nextImageSemaphore; + m_queue->m_pendingWaitSemaphores[1] = m_nextImageSemaphore; return m_currentImageIndex; } }; @@ -4730,7 +5085,6 @@ Result VKDevice::Buffer::init(const VulkanApi& api, size_t bufferSize, VkBufferU VkBufferCreateInfo bufferCreateInfo = { VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO }; bufferCreateInfo.size = bufferSize; bufferCreateInfo.usage = usage; - SLANG_VK_CHECK(api.vkCreateBuffer(api.m_device, &bufferCreateInfo, nullptr, &m_buffer)); VkMemoryRequirements memoryReqs = {}; @@ -4744,7 +5098,14 @@ Result VKDevice::Buffer::init(const VulkanApi& api, size_t bufferSize, VkBufferU VkMemoryAllocateInfo allocateInfo = { VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO }; allocateInfo.allocationSize = memoryReqs.size; allocateInfo.memoryTypeIndex = memoryTypeIndex; - + VkMemoryAllocateFlagsInfo flagInfo = {VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_FLAGS_INFO}; + if (usage & VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT) + { + flagInfo.deviceMask = 1; + flagInfo.flags = VK_MEMORY_ALLOCATE_DEVICE_ADDRESS_BIT; + allocateInfo.pNext = &flagInfo; + } + SLANG_VK_CHECK(api.vkAllocateMemory(api.m_device, &allocateInfo, nullptr, &m_memory)); SLANG_VK_CHECK(api.vkBindBufferMemory(api.m_device, m_buffer, m_memory, 0)); @@ -5042,24 +5403,6 @@ Result VKDevice::initVulkanInstanceAndDevice(bool useValidationLayer) const uint32_t majorVersion = VK_VERSION_MAJOR(basicProps.apiVersion); const uint32_t minorVersion = VK_VERSION_MINOR(basicProps.apiVersion); - // Need in this scope because it will be linked into the device creation (if it is available) - - // Float16 features - VkPhysicalDeviceFloat16Int8FeaturesKHR float16Features = { VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FLOAT16_INT8_FEATURES_KHR }; - // 16 bit storage features - VkPhysicalDevice16BitStorageFeatures storage16BitFeatures = { VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_16BIT_STORAGE_FEATURES_KHR }; - // AtomicInt64 features - VkPhysicalDeviceShaderAtomicInt64FeaturesKHR atomicInt64Features = { VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_ATOMIC_INT64_FEATURES_KHR }; - // Atomic Float features - VkPhysicalDeviceShaderAtomicFloatFeaturesEXT atomicFloatFeatures = { VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_ATOMIC_FLOAT_FEATURES_EXT }; - // Timeline Semaphore features - VkPhysicalDeviceTimelineSemaphoreFeatures timelineFeatures = { VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_TIMELINE_SEMAPHORE_FEATURES }; - // Extended dynamic state features - VkPhysicalDeviceExtendedDynamicStateFeaturesEXT extendedDynamicStateFeatures = { VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTENDED_DYNAMIC_STATE_FEATURES_EXT }; - // Subgroup extended type features - VkPhysicalDeviceShaderSubgroupExtendedTypesFeatures shaderSubgroupExtendedTypeFeatures = { - VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_SUBGROUP_EXTENDED_TYPES_FEATURES}; - // API version check, can't use vkGetPhysicalDeviceProperties2 yet since this device might not support it if (VK_MAKE_VERSION(majorVersion, minorVersion, 0) >= VK_API_VERSION_1_1 && m_api.vkGetPhysicalDeviceProperties2 && @@ -5069,6 +5412,18 @@ Result VKDevice::initVulkanInstanceAndDevice(bool useValidationLayer) VkPhysicalDeviceFeatures2 deviceFeatures2 = {}; deviceFeatures2.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_FEATURES_2; + // Buffer device address features + bufferDeviceAddressFeatures.pNext = deviceFeatures2.pNext; + deviceFeatures2.pNext = &bufferDeviceAddressFeatures; + + // Ray query features + rayQueryFeatures.pNext = deviceFeatures2.pNext; + deviceFeatures2.pNext = &rayQueryFeatures; + + // Acceleration structure features + accelerationStructureFeatures.pNext = deviceFeatures2.pNext; + deviceFeatures2.pNext = &accelerationStructureFeatures; + // Subgroup features shaderSubgroupExtendedTypeFeatures.pNext = deviceFeatures2.pNext; deviceFeatures2.pNext = &shaderSubgroupExtendedTypeFeatures; @@ -5174,6 +5529,31 @@ Result VKDevice::initVulkanInstanceAndDevice(bool useValidationLayer) deviceExtensions.add(VK_KHR_SHADER_SUBGROUP_EXTENDED_TYPES_EXTENSION_NAME); m_features.add("shader-subgroup-extended-types"); } + + if (accelerationStructureFeatures.accelerationStructure) + { + accelerationStructureFeatures.pNext = (void*)deviceCreateInfo.pNext; + deviceCreateInfo.pNext = &accelerationStructureFeatures; + deviceExtensions.add(VK_KHR_ACCELERATION_STRUCTURE_EXTENSION_NAME); + deviceExtensions.add(VK_KHR_DEFERRED_HOST_OPERATIONS_EXTENSION_NAME); + m_features.add("acceleration-structure"); + } + + if (rayQueryFeatures.rayQuery) + { + rayQueryFeatures.pNext = (void*)deviceCreateInfo.pNext; + deviceCreateInfo.pNext = &rayQueryFeatures; + deviceExtensions.add(VK_KHR_RAY_QUERY_EXTENSION_NAME); + m_features.add("ray-query"); + } + + if (bufferDeviceAddressFeatures.bufferDeviceAddress) + { + bufferDeviceAddressFeatures.pNext = (void*)deviceCreateInfo.pNext; + deviceCreateInfo.pNext = &bufferDeviceAddressFeatures; + deviceExtensions.add(VK_KHR_BUFFER_DEVICE_ADDRESS_EXTENSION_NAME); + m_features.add("buffer-device-address"); + } } m_queueFamilyIndex = m_api.findQueue(VK_QUEUE_GRAPHICS_BIT | VK_QUEUE_COMPUTE_BIT); @@ -5255,11 +5635,7 @@ Result VKDevice::TransientResourceHeapImpl::init( device->m_api.vkCreateCommandPool( device->m_api.m_device, &poolCreateInfo, nullptr, &m_commandPool); - VkFenceCreateInfo fenceCreateInfo = {}; - fenceCreateInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; - fenceCreateInfo.flags = VK_FENCE_CREATE_SIGNALED_BIT; - device->m_api.vkCreateFence(device->m_api.m_device, &fenceCreateInfo, nullptr, &m_fence); - + advanceFence(); return SLANG_OK; } @@ -5276,7 +5652,7 @@ Result VKDevice::TransientResourceHeapImpl::createCommandBuffer(ICommandBuffer** RefPtr<CommandBufferImpl> commandBuffer = new CommandBufferImpl(); SLANG_RETURN_ON_FAIL(commandBuffer->init( - m_device, m_commandPool, m_fence, this)); + m_device, m_commandPool, this)); m_commandBufferPool.add(commandBuffer); m_commandBufferAllocId++; returnComPtr(outCmdBuffer, commandBuffer); @@ -5287,12 +5663,15 @@ Result VKDevice::TransientResourceHeapImpl::synchronizeAndReset() { m_commandBufferAllocId = 0; auto& api = m_device->m_api; - if (api.vkWaitForFences(api.m_device, 1, &m_fence, 1, UINT64_MAX) != VK_SUCCESS) + if (api.vkWaitForFences( + api.m_device, (uint32_t)m_fences.getCount(), m_fences.getBuffer(), 1, UINT64_MAX) != + VK_SUCCESS) { return SLANG_FAIL; } api.vkResetCommandPool(api.m_device, m_commandPool, 0); m_descSetAllocator.reset(); + m_fenceIndex = 0; Super::reset(); return SLANG_OK; } @@ -5421,6 +5800,69 @@ SlangResult VKDevice::readBufferResource( return SLANG_OK; } +Result VKDevice::getAccelerationStructurePrebuildInfo( + const IAccelerationStructure::BuildInputs& buildInputs, + IAccelerationStructure::PrebuildInfo* outPrebuildInfo) +{ + if (!m_api.vkGetAccelerationStructureBuildSizesKHR) + { + return SLANG_E_NOT_AVAILABLE; + } + VkAccelerationStructureBuildSizesInfoKHR sizeInfo = { + VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_BUILD_SIZES_INFO_KHR}; + AccelerationStructureBuildGeometryInfoBuilder geomInfoBuilder; + SLANG_RETURN_ON_FAIL(geomInfoBuilder.build(buildInputs, getDebugCallback())); + m_api.vkGetAccelerationStructureBuildSizesKHR( + m_api.m_device, + VK_ACCELERATION_STRUCTURE_BUILD_TYPE_DEVICE_KHR, + &geomInfoBuilder.buildInfo, + geomInfoBuilder.primitiveCounts.getBuffer(), + &sizeInfo); + outPrebuildInfo->resultDataMaxSize = sizeInfo.accelerationStructureSize; + outPrebuildInfo->scratchDataSize = sizeInfo.buildScratchSize; + outPrebuildInfo->updateScratchDataSize = sizeInfo.updateScratchSize; + return SLANG_OK; +} + +Result VKDevice::createAccelerationStructure( + const IAccelerationStructure::CreateDesc& desc, + IAccelerationStructure** outAS) +{ + if (!m_api.vkCreateAccelerationStructureKHR) + { + return SLANG_E_NOT_AVAILABLE; + } + RefPtr<AccelerationStructureImpl> resultAS = new AccelerationStructureImpl(); + resultAS->m_offset = desc.offset; + resultAS->m_size = desc.size; + resultAS->m_buffer = static_cast<BufferResourceImpl*>(desc.buffer); + resultAS->m_device = this; + VkAccelerationStructureCreateInfoKHR createInfo = {VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_CREATE_INFO_KHR}; + createInfo.buffer = resultAS->m_buffer->m_buffer.m_buffer; + createInfo.offset = desc.offset; + createInfo.size = desc.size; + switch (desc.kind) + { + case IAccelerationStructure::Kind::BottomLevel: + createInfo.type = VK_ACCELERATION_STRUCTURE_TYPE_BOTTOM_LEVEL_KHR; + break; + case IAccelerationStructure::Kind::TopLevel: + createInfo.type = VK_ACCELERATION_STRUCTURE_TYPE_TOP_LEVEL_KHR; + break; + default: + getDebugCallback()->handleMessage( + DebugMessageType::Error, + DebugMessageSource::Layer, + "invalid value of IAccelerationStructure::Kind encountered in desc.kind"); + return SLANG_E_INVALID_ARG; + } + + SLANG_VK_RETURN_ON_FAIL(m_api.vkCreateAccelerationStructureKHR( + m_api.m_device, &createInfo, nullptr, &resultAS->m_vkHandle)); + returnComPtr(outAS, resultAS); + return SLANG_OK; +} + static VkBufferUsageFlagBits _calcBufferUsageFlags(ResourceState state) { switch (state) @@ -5448,6 +5890,8 @@ static VkBufferUsageFlagBits _calcBufferUsageFlags(ResourceState state) return VK_BUFFER_USAGE_TRANSFER_SRC_BIT; case ResourceState::CopyDestination: return VK_BUFFER_USAGE_TRANSFER_DST_BIT; + case ResourceState::AccelerationStructure: + return VK_BUFFER_USAGE_ACCELERATION_STRUCTURE_STORAGE_BIT_KHR; default: return VkBufferUsageFlagBits(0); } @@ -5510,7 +5954,7 @@ static VkImageUsageFlags _calcImageUsageFlags( { VkImageUsageFlags usage = _calcImageUsageFlags(states); - if ((cpuAccessFlags & IResource::AccessFlag::Write) || initData) + if ((cpuAccessFlags & AccessFlag::Write) || initData) { usage |= VK_IMAGE_USAGE_TRANSFER_DST_BIT; } @@ -5626,6 +6070,15 @@ void VKDevice::_transitionImageLayout(VkImage image, VkFormat format, const Text sourceStage = VK_PIPELINE_STAGE_BOTTOM_OF_PIPE_BIT; destinationStage = VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT; } + else if (oldLayout == VK_IMAGE_LAYOUT_UNDEFINED && + newLayout == VK_IMAGE_LAYOUT_GENERAL) + { + barrier.srcAccessMask = 0; + barrier.dstAccessMask = VK_ACCESS_MEMORY_READ_BIT | VK_ACCESS_MEMORY_WRITE_BIT; + + sourceStage = VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT; + destinationStage = VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT; + } else { assert(!"unsupported layout transition!"); @@ -5902,6 +6355,10 @@ Result VKDevice::createBufferResource(const IBufferResource::Desc& descIn, const VkMemoryPropertyFlags reqMemoryProperties = 0; VkBufferUsageFlags usage = _calcBufferUsageFlags(desc.allowedStates); + if (bufferDeviceAddressFeatures.bufferDeviceAddress) + { + usage |= VK_BUFFER_USAGE_SHADER_DEVICE_ADDRESS_BIT; + } if (initData) { @@ -5916,7 +6373,7 @@ Result VKDevice::createBufferResource(const IBufferResource::Desc& descIn, const RefPtr<BufferResourceImpl> buffer(new BufferResourceImpl(desc, this)); SLANG_RETURN_ON_FAIL(buffer->m_buffer.init(m_api, desc.sizeInBytes, usage, reqMemoryProperties)); - if ((desc.cpuAccessFlags & IResource::AccessFlag::Write) || initData) + if ((desc.cpuAccessFlags & AccessFlag::Write) || initData) { SLANG_RETURN_ON_FAIL(buffer->m_uploadBuffer.init(m_api, bufferSize, VK_BUFFER_USAGE_TRANSFER_SRC_BIT, VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT)); } diff --git a/tools/gfx/vulkan/vk-api.h b/tools/gfx/vulkan/vk-api.h index 4c024525b..5d044944f 100644 --- a/tools/gfx/vulkan/vk-api.h +++ b/tools/gfx/vulkan/vk-api.h @@ -159,7 +159,14 @@ namespace gfx { x(vkGetBufferDeviceAddress) \ x(vkGetBufferDeviceAddressKHR) \ x(vkGetBufferDeviceAddressEXT) \ - + x(vkCmdBuildAccelerationStructuresKHR) \ + x(vkCmdCopyAccelerationStructureKHR) \ + x(vkCmdCopyAccelerationStructureToMemoryKHR) \ + x(vkCmdCopyMemoryToAccelerationStructureKHR) \ + x(vkCmdWriteAccelerationStructuresPropertiesKHR) \ + x(vkCreateAccelerationStructureKHR) \ + x(vkDestroyAccelerationStructureKHR) \ + x(vkGetAccelerationStructureBuildSizesKHR) \ /* */ #define VK_API_ALL_GLOBAL_PROCS(x) \ diff --git a/tools/gfx/vulkan/vk-util.cpp b/tools/gfx/vulkan/vk-util.cpp index 5f7077753..4971e3f3d 100644 --- a/tools/gfx/vulkan/vk-util.cpp +++ b/tools/gfx/vulkan/vk-util.cpp @@ -1,5 +1,6 @@ // vk-util.cpp #include "vk-util.h" +#include "core/slang-math.h" #include <stdlib.h> #include <stdio.h> @@ -170,4 +171,145 @@ VkImageLayout VulkanUtil::mapResourceStateToLayout(ResourceState state) } } -} // renderer_test +Result AccelerationStructureBuildGeometryInfoBuilder::build( + const IAccelerationStructure::BuildInputs& buildInputs, + IDebugCallback* debugCallback) +{ + buildInfo.dstAccelerationStructure = VK_NULL_HANDLE; + switch (buildInputs.kind) + { + case IAccelerationStructure::Kind::BottomLevel: + buildInfo.type = VK_ACCELERATION_STRUCTURE_TYPE_BOTTOM_LEVEL_KHR; + break; + case IAccelerationStructure::Kind::TopLevel: + buildInfo.type = VK_ACCELERATION_STRUCTURE_TYPE_TOP_LEVEL_KHR; + break; + default: + debugCallback->handleMessage( + DebugMessageType::Error, + DebugMessageSource::Layer, + "invalid value of IAccelerationStructure::Kind encountered in buildInputs.kind"); + return SLANG_E_INVALID_ARG; + } + if (buildInputs.flags & IAccelerationStructure::BuildFlags::Enum::PerformUpdate) + { + buildInfo.mode = VK_BUILD_ACCELERATION_STRUCTURE_MODE_UPDATE_KHR; + } + else + { + buildInfo.mode = VK_BUILD_ACCELERATION_STRUCTURE_MODE_BUILD_KHR; + } + if (buildInputs.flags & IAccelerationStructure::BuildFlags::Enum::AllowCompaction) + { + buildInfo.flags |= VK_BUILD_ACCELERATION_STRUCTURE_ALLOW_COMPACTION_BIT_KHR; + } + if (buildInputs.flags & IAccelerationStructure::BuildFlags::Enum::AllowUpdate) + { + buildInfo.flags |= VK_BUILD_ACCELERATION_STRUCTURE_ALLOW_UPDATE_BIT_KHR; + } + if (buildInputs.flags & IAccelerationStructure::BuildFlags::Enum::MinimizeMemory) + { + buildInfo.flags |= VK_BUILD_ACCELERATION_STRUCTURE_LOW_MEMORY_BIT_KHR; + } + if (buildInputs.flags & IAccelerationStructure::BuildFlags::Enum::PreferFastBuild) + { + buildInfo.flags |= VK_BUILD_ACCELERATION_STRUCTURE_PREFER_FAST_BUILD_BIT_KHR; + } + if (buildInputs.flags & IAccelerationStructure::BuildFlags::Enum::PreferFastTrace) + { + buildInfo.flags |= VK_BUILD_ACCELERATION_STRUCTURE_PREFER_FAST_TRACE_BIT_KHR; + } + if (buildInputs.kind == IAccelerationStructure::Kind::BottomLevel) + { + m_geometryInfos.setCount(buildInputs.descCount); + primitiveCounts.setCount(buildInputs.descCount); + memset( + m_geometryInfos.getBuffer(), + 0, + sizeof(VkAccelerationStructureGeometryKHR) * buildInputs.descCount); + for (int i = 0; i < buildInputs.descCount; i++) + { + auto& geomDesc = buildInputs.geometryDescs[i]; + m_geometryInfos[i].sType = VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_KHR; + if (geomDesc.flags & IAccelerationStructure::GeometryFlags::NoDuplicateAnyHitInvocation) + { + m_geometryInfos[i].flags |= VK_GEOMETRY_NO_DUPLICATE_ANY_HIT_INVOCATION_BIT_KHR; + } + else if (geomDesc.flags & IAccelerationStructure::GeometryFlags::Opaque) + { + m_geometryInfos[i].flags |= VK_GEOMETRY_OPAQUE_BIT_KHR; + } + auto& vkGeomData = m_geometryInfos[i].geometry; + switch (geomDesc.type) + { + case IAccelerationStructure::GeometryType::Triangles: + m_geometryInfos[i].geometryType = VK_GEOMETRY_TYPE_TRIANGLES_KHR; + vkGeomData.triangles.sType = + VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_TRIANGLES_DATA_KHR; + vkGeomData.triangles.vertexFormat = + VulkanUtil::getVkFormat(geomDesc.content.triangles.vertexFormat); + vkGeomData.triangles.vertexData.deviceAddress = + geomDesc.content.triangles.vertexData; + vkGeomData.triangles.vertexStride = geomDesc.content.triangles.vertexStride; + vkGeomData.triangles.maxVertex = geomDesc.content.triangles.vertexCount - 1; + switch (geomDesc.content.triangles.indexFormat) + { + case Format::R_UInt32: + vkGeomData.triangles.indexType = VK_INDEX_TYPE_UINT32; + break; + case Format::R_UInt16: + vkGeomData.triangles.indexType = VK_INDEX_TYPE_UINT16; + break; + default: + debugCallback->handleMessage( + DebugMessageType::Error, + DebugMessageSource::Layer, + "unsupported value of Format encountered in " + "GeometryDesc::content.triangles.indexFormat"); + return SLANG_E_INVALID_ARG; + } + vkGeomData.triangles.indexData.deviceAddress = geomDesc.content.triangles.indexData; + vkGeomData.triangles.transformData.deviceAddress = + geomDesc.content.triangles.transform3x4; + primitiveCounts[i] = Slang::Math::Max( + geomDesc.content.triangles.vertexCount, + geomDesc.content.triangles.indexCount) / + 3; + break; + case IAccelerationStructure::GeometryType::ProcedurePrimitives: + m_geometryInfos[i].geometryType = VK_GEOMETRY_TYPE_AABBS_KHR; + vkGeomData.aabbs.sType = + VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_AABBS_DATA_KHR; + vkGeomData.aabbs.data.deviceAddress = geomDesc.content.proceduralAABBs.data; + vkGeomData.aabbs.stride = geomDesc.content.proceduralAABBs.stride; + primitiveCounts[i] = + (uint32_t)buildInputs.geometryDescs[i].content.proceduralAABBs.count; + break; + default: + debugCallback->handleMessage( + DebugMessageType::Error, + DebugMessageSource::Layer, + "invalid value of IAccelerationStructure::GeometryType encountered in " + "buildInputs.geometryDescs"); + return SLANG_E_INVALID_ARG; + } + } + buildInfo.geometryCount = buildInputs.descCount; + buildInfo.pGeometries = m_geometryInfos.getBuffer(); + } + else + { + m_vkInstanceInfo.geometryType = VK_GEOMETRY_TYPE_INSTANCES_KHR; + m_vkInstanceInfo.geometry.instances.sType = + VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_INSTANCES_DATA_KHR; + m_vkInstanceInfo.geometry.instances.arrayOfPointers = 0; + m_vkInstanceInfo.geometry.instances.data.deviceAddress = buildInputs.instanceDescs; + buildInfo.pGeometries = &m_vkInstanceInfo; + buildInfo.geometryCount = 1; + primitiveCounts.setCount(1); + primitiveCounts[0] = buildInputs.descCount; + } + return SLANG_OK; +} + +} // namespace gfx diff --git a/tools/gfx/vulkan/vk-util.h b/tools/gfx/vulkan/vk-util.h index a39fe5115..450e78ebd 100644 --- a/tools/gfx/vulkan/vk-util.h +++ b/tools/gfx/vulkan/vk-util.h @@ -1,6 +1,7 @@ // vk-util.h #pragma once +#include "core/slang-basic.h" #include "vk-api.h" #include "slang-gfx.h" @@ -46,4 +47,22 @@ struct VulkanUtil static VkImageLayout getImageLayoutFromState(ResourceState state); }; +struct AccelerationStructureBuildGeometryInfoBuilder +{ +public: + VkAccelerationStructureBuildGeometryInfoKHR buildInfo = { + VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_BUILD_GEOMETRY_INFO_KHR}; + Slang::List<uint32_t> primitiveCounts; + + Slang::Result build( + const IAccelerationStructure::BuildInputs& buildInputs, + IDebugCallback* debugCallback); + +private: + Slang::List<VkAccelerationStructureGeometryKHR> m_geometryInfos; + VkAccelerationStructureGeometryKHR m_vkInstanceInfo = { + VK_STRUCTURE_TYPE_ACCELERATION_STRUCTURE_GEOMETRY_KHR}; +}; + + } // renderer_test diff --git a/tools/platform/gui.cpp b/tools/platform/gui.cpp index ba6fc35f8..e4f269b95 100644 --- a/tools/platform/gui.cpp +++ b/tools/platform/gui.cpp @@ -216,7 +216,7 @@ void GUI::endFrame(ITransientResourceHeap* transientHeap, IFramebuffer* framebuf vertexBufferDesc.allowedStates = ResourceStateSet(ResourceState::VertexBuffer, ResourceState::CopyDestination); vertexBufferDesc.sizeInBytes = vertexCount * sizeof(ImDrawVert); - vertexBufferDesc.cpuAccessFlags = IResource::AccessFlag::Write; + vertexBufferDesc.cpuAccessFlags = AccessFlag::Write; auto vertexBuffer = device->createBufferResource(vertexBufferDesc); gfx::IBufferResource::Desc indexBufferDesc; @@ -225,7 +225,7 @@ void GUI::endFrame(ITransientResourceHeap* transientHeap, IFramebuffer* framebuf indexBufferDesc.allowedStates = ResourceStateSet(ResourceState::IndexBuffer, ResourceState::CopyDestination); indexBufferDesc.defaultState = ResourceState::IndexBuffer; - indexBufferDesc.cpuAccessFlags = IResource::AccessFlag::Write; + indexBufferDesc.cpuAccessFlags = AccessFlag::Write; auto indexBuffer = device->createBufferResource(indexBufferDesc); auto cmdBuf = transientHeap->createCommandBuffer(); auto encoder = cmdBuf->encodeResourceCommands(); @@ -253,7 +253,7 @@ void GUI::endFrame(ITransientResourceHeap* transientHeap, IFramebuffer* framebuf ResourceStateSet(ResourceState::ConstantBuffer, ResourceState::CopyDestination); constantBufferDesc.defaultState = ResourceState::ConstantBuffer; constantBufferDesc.sizeInBytes = sizeof(glm::mat4x4); - constantBufferDesc.cpuAccessFlags = IResource::AccessFlag::Write; + constantBufferDesc.cpuAccessFlags = AccessFlag::Write; auto constantBuffer = device->createBufferResource(constantBufferDesc); { diff --git a/tools/render-test/render-test-main.cpp b/tools/render-test/render-test-main.cpp index a846f4a33..03662a480 100644 --- a/tools/render-test/render-test-main.cpp +++ b/tools/render-test/render-test-main.cpp @@ -435,9 +435,7 @@ Result RenderTestApp::applyBinding(PipelineType pipelineType, ICommandEncoder* e { case PipelineType::Compute: { - ComPtr<IComputeCommandEncoder> computeEncoder; - encoder->queryInterface( - SLANG_UUID_IComputeCommandEncoder, (void**)computeEncoder.writeRef()); + IComputeCommandEncoder* computeEncoder = static_cast<IComputeCommandEncoder*>(encoder); auto rootObject = computeEncoder->bindPipeline(m_pipelineState); SLANG_RETURN_ON_FAIL(_assignVarsFromLayout( m_device, rootObject, m_compilationOutput.layout, m_outputPlan, slangReflection)); @@ -445,9 +443,7 @@ Result RenderTestApp::applyBinding(PipelineType pipelineType, ICommandEncoder* e break; case PipelineType::Graphics: { - ComPtr<IRenderCommandEncoder> renderEncoder; - encoder->queryInterface( - SLANG_UUID_IRenderCommandEncoder, (void**)renderEncoder.writeRef()); + IRenderCommandEncoder* renderEncoder = static_cast<IRenderCommandEncoder*>(encoder); auto rootObject = renderEncoder->bindPipeline(m_pipelineState); SLANG_RETURN_ON_FAIL(_assignVarsFromLayout( m_device, rootObject, m_compilationOutput.layout, m_outputPlan, slangReflection)); @@ -523,7 +519,7 @@ SlangResult RenderTestApp::initialize( IBufferResource::Desc vertexBufferDesc; vertexBufferDesc.type = IResource::Type::Buffer; vertexBufferDesc.sizeInBytes = kVertexCount * sizeof(Vertex); - vertexBufferDesc.cpuAccessFlags = IResource::AccessFlag::Write; + vertexBufferDesc.cpuAccessFlags = AccessFlag::Write; vertexBufferDesc.defaultState = ResourceState::VertexBuffer; vertexBufferDesc.allowedStates = ResourceStateSet(ResourceState::VertexBuffer); @@ -697,7 +693,7 @@ Result RenderTestApp::writeBindingOutput(const String& fileName) const size_t bufferSize = bufferDesc.sizeInBytes; ComPtr<ISlangBlob> blob; - if(bufferDesc.cpuAccessFlags & IResource::AccessFlag::Read) + if(bufferDesc.cpuAccessFlags & AccessFlag::Read) { // The buffer is already allocated for CPU access, so we can read it back directly. // @@ -708,7 +704,7 @@ Result RenderTestApp::writeBindingOutput(const String& fileName) // The buffer is not CPU-readable, so we will copy it using a staging buffer. auto stagingBufferDesc = bufferDesc; - stagingBufferDesc.cpuAccessFlags = IResource::AccessFlag::Read; + stagingBufferDesc.cpuAccessFlags = AccessFlag::Read; stagingBufferDesc.allowedStates = ResourceStateSet(ResourceState::CopyDestination, ResourceState::CopySource); stagingBufferDesc.defaultState = ResourceState::CopyDestination; @@ -720,8 +716,8 @@ Result RenderTestApp::writeBindingOutput(const String& fileName) SLANG_RETURN_ON_FAIL( m_transientHeap->createCommandBuffer(commandBuffer.writeRef())); - ComPtr<IResourceCommandEncoder> encoder; - commandBuffer->encodeResourceCommands(encoder.writeRef()); + IResourceCommandEncoder* encoder = nullptr; + commandBuffer->encodeResourceCommands(&encoder); encoder->copyBuffer(stagingBuffer, 0, bufferResource, 0, bufferSize); encoder->endEncoding(); |