summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--build/visual-studio/autodiff-texture/autodiff-texture.vcxproj307
-rw-r--r--build/visual-studio/autodiff-texture/autodiff-texture.vcxproj.filters33
-rw-r--r--examples/autodiff-texture/README.md4
-rw-r--r--examples/autodiff-texture/buildmip.slang25
-rw-r--r--examples/autodiff-texture/checkerboard.jpgbin0 -> 51262 bytes
-rw-r--r--examples/autodiff-texture/convert.slang25
-rw-r--r--examples/autodiff-texture/draw-quad.slang52
-rw-r--r--examples/autodiff-texture/learnmip.slang22
-rw-r--r--examples/autodiff-texture/main.cpp662
-rw-r--r--examples/autodiff-texture/reconstruct.slang48
-rw-r--r--examples/autodiff-texture/train.slang216
-rw-r--r--examples/example-base/example-base.cpp74
-rw-r--r--examples/example-base/example-base.h2
-rw-r--r--examples/heterogeneous-hello-world/README.md4
-rw-r--r--examples/heterogeneous-hello-world/main.slang25
-rw-r--r--premake5.lua12
-rw-r--r--slang.sln15
-rw-r--r--source/slang/slang-check-decl.cpp33
-rw-r--r--source/slang/slang-ir-autodiff-fwd.cpp13
-rw-r--r--tools/platform/windows/win-window.cpp1
20 files changed, 1515 insertions, 58 deletions
diff --git a/build/visual-studio/autodiff-texture/autodiff-texture.vcxproj b/build/visual-studio/autodiff-texture/autodiff-texture.vcxproj
new file mode 100644
index 000000000..c8bac6e9d
--- /dev/null
+++ b/build/visual-studio/autodiff-texture/autodiff-texture.vcxproj
@@ -0,0 +1,307 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project DefaultTargets="Build" 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="Debug|ARM64">
+ <Configuration>Debug</Configuration>
+ <Platform>ARM64</Platform>
+ </ProjectConfiguration>
+ <ProjectConfiguration Include="Debug aarch64|Win32">
+ <Configuration>Debug aarch64</Configuration>
+ <Platform>Win32</Platform>
+ </ProjectConfiguration>
+ <ProjectConfiguration Include="Debug aarch64|x64">
+ <Configuration>Debug aarch64</Configuration>
+ <Platform>x64</Platform>
+ </ProjectConfiguration>
+ <ProjectConfiguration Include="Debug aarch64|ARM64">
+ <Configuration>Debug aarch64</Configuration>
+ <Platform>ARM64</Platform>
+ </ProjectConfiguration>
+ <ProjectConfiguration Include="Release|Win32">
+ <Configuration>Release</Configuration>
+ <Platform>Win32</Platform>
+ </ProjectConfiguration>
+ <ProjectConfiguration Include="Release|x64">
+ <Configuration>Release</Configuration>
+ <Platform>x64</Platform>
+ </ProjectConfiguration>
+ <ProjectConfiguration Include="Release|ARM64">
+ <Configuration>Release</Configuration>
+ <Platform>ARM64</Platform>
+ </ProjectConfiguration>
+ <ProjectConfiguration Include="Release aarch64|Win32">
+ <Configuration>Release aarch64</Configuration>
+ <Platform>Win32</Platform>
+ </ProjectConfiguration>
+ <ProjectConfiguration Include="Release aarch64|x64">
+ <Configuration>Release aarch64</Configuration>
+ <Platform>x64</Platform>
+ </ProjectConfiguration>
+ <ProjectConfiguration Include="Release aarch64|ARM64">
+ <Configuration>Release aarch64</Configuration>
+ <Platform>ARM64</Platform>
+ </ProjectConfiguration>
+ </ItemGroup>
+ <PropertyGroup Label="Globals">
+ <ProjectGuid>{2F0A3884-1B72-FE51-84DD-12687064153E}</ProjectGuid>
+ <IgnoreWarnCompileDuplicatedFilename>true</IgnoreWarnCompileDuplicatedFilename>
+ <Keyword>Win32Proj</Keyword>
+ <RootNamespace>autodiff-texture</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>v142</PlatformToolset>
+ </PropertyGroup>
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'" Label="Configuration">
+ <ConfigurationType>Application</ConfigurationType>
+ <UseDebugLibraries>true</UseDebugLibraries>
+ <CharacterSet>Unicode</CharacterSet>
+ <PlatformToolset>v142</PlatformToolset>
+ </PropertyGroup>
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug aarch64|ARM64'" Label="Configuration">
+ <ConfigurationType>Application</ConfigurationType>
+ <UseDebugLibraries>true</UseDebugLibraries>
+ <CharacterSet>Unicode</CharacterSet>
+ <PlatformToolset>v142</PlatformToolset>
+ <WindowsSDKDesktopARM64Support>true</WindowsSDKDesktopARM64Support>
+ </PropertyGroup>
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'" Label="Configuration">
+ <ConfigurationType>Application</ConfigurationType>
+ <UseDebugLibraries>false</UseDebugLibraries>
+ <CharacterSet>Unicode</CharacterSet>
+ <PlatformToolset>v142</PlatformToolset>
+ </PropertyGroup>
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" Label="Configuration">
+ <ConfigurationType>Application</ConfigurationType>
+ <UseDebugLibraries>false</UseDebugLibraries>
+ <CharacterSet>Unicode</CharacterSet>
+ <PlatformToolset>v142</PlatformToolset>
+ </PropertyGroup>
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release aarch64|ARM64'" Label="Configuration">
+ <ConfigurationType>Application</ConfigurationType>
+ <UseDebugLibraries>false</UseDebugLibraries>
+ <CharacterSet>Unicode</CharacterSet>
+ <PlatformToolset>v142</PlatformToolset>
+ <WindowsSDKDesktopARM64Support>true</WindowsSDKDesktopARM64Support>
+ </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)'=='Debug aarch64|ARM64'">
+ <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>
+ <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Release aarch64|ARM64'">
+ <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\autodiff-texture\</IntDir>
+ <TargetName>autodiff-texture</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\autodiff-texture\</IntDir>
+ <TargetName>autodiff-texture</TargetName>
+ <TargetExt>.exe</TargetExt>
+ </PropertyGroup>
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug aarch64|ARM64'">
+ <LinkIncremental>true</LinkIncremental>
+ <OutDir>..\..\..\bin\windows-aarch64\debug\</OutDir>
+ <IntDir>..\..\..\intermediate\windows-aarch64\debug\autodiff-texture\</IntDir>
+ <TargetName>autodiff-texture</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\autodiff-texture\</IntDir>
+ <TargetName>autodiff-texture</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\autodiff-texture\</IntDir>
+ <TargetName>autodiff-texture</TargetName>
+ <TargetExt>.exe</TargetExt>
+ </PropertyGroup>
+ <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release aarch64|ARM64'">
+ <LinkIncremental>false</LinkIncremental>
+ <OutDir>..\..\..\bin\windows-aarch64\release\</OutDir>
+ <IntDir>..\..\..\intermediate\windows-aarch64\release\autodiff-texture\</IntDir>
+ <TargetName>autodiff-texture</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>ProgramDatabase</DebugInformationFormat>
+ <Optimization>Disabled</Optimization>
+ <MinimalRebuild>false</MinimalRebuild>
+ <RuntimeLibrary>MultiThreadedDebug</RuntimeLibrary>
+ <MultiProcessorCompilation>true</MultiProcessorCompilation>
+ </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>ProgramDatabase</DebugInformationFormat>
+ <Optimization>Disabled</Optimization>
+ <MinimalRebuild>false</MinimalRebuild>
+ <RuntimeLibrary>MultiThreadedDebug</RuntimeLibrary>
+ <MultiProcessorCompilation>true</MultiProcessorCompilation>
+ </ClCompile>
+ <Link>
+ <SubSystem>Windows</SubSystem>
+ <GenerateDebugInformation>true</GenerateDebugInformation>
+ </Link>
+ </ItemDefinitionGroup>
+ <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug aarch64|ARM64'">
+ <ClCompile>
+ <PrecompiledHeader>NotUsing</PrecompiledHeader>
+ <WarningLevel>Level3</WarningLevel>
+ <PreprocessorDefinitions>_DEBUG;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+ <AdditionalIncludeDirectories>..\..\..;..\..\..\tools;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+ <DebugInformationFormat>ProgramDatabase</DebugInformationFormat>
+ <Optimization>Disabled</Optimization>
+ <MinimalRebuild>false</MinimalRebuild>
+ <RuntimeLibrary>MultiThreadedDebug</RuntimeLibrary>
+ <MultiProcessorCompilation>true</MultiProcessorCompilation>
+ </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>
+ <MultiProcessorCompilation>true</MultiProcessorCompilation>
+ </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>
+ <MultiProcessorCompilation>true</MultiProcessorCompilation>
+ </ClCompile>
+ <Link>
+ <SubSystem>Windows</SubSystem>
+ <EnableCOMDATFolding>true</EnableCOMDATFolding>
+ <OptimizeReferences>true</OptimizeReferences>
+ </Link>
+ </ItemDefinitionGroup>
+ <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release aarch64|ARM64'">
+ <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>
+ <MultiProcessorCompilation>true</MultiProcessorCompilation>
+ </ClCompile>
+ <Link>
+ <SubSystem>Windows</SubSystem>
+ <EnableCOMDATFolding>true</EnableCOMDATFolding>
+ <OptimizeReferences>true</OptimizeReferences>
+ </Link>
+ </ItemDefinitionGroup>
+ <ItemGroup>
+ <ClCompile Include="..\..\..\examples\autodiff-texture\main.cpp" />
+ </ItemGroup>
+ <ItemGroup>
+ <None Include="..\..\..\examples\autodiff-texture\buildmip.slang" />
+ <None Include="..\..\..\examples\autodiff-texture\convert.slang" />
+ <None Include="..\..\..\examples\autodiff-texture\draw-quad.slang" />
+ <None Include="..\..\..\examples\autodiff-texture\learnmip.slang" />
+ <None Include="..\..\..\examples\autodiff-texture\reconstruct.slang" />
+ <None Include="..\..\..\examples\autodiff-texture\train.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/autodiff-texture/autodiff-texture.vcxproj.filters b/build/visual-studio/autodiff-texture/autodiff-texture.vcxproj.filters
new file mode 100644
index 000000000..17045ac96
--- /dev/null
+++ b/build/visual-studio/autodiff-texture/autodiff-texture.vcxproj.filters
@@ -0,0 +1,33 @@
+<?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\autodiff-texture\main.cpp">
+ <Filter>Source Files</Filter>
+ </ClCompile>
+ </ItemGroup>
+ <ItemGroup>
+ <None Include="..\..\..\examples\autodiff-texture\buildmip.slang">
+ <Filter>Source Files</Filter>
+ </None>
+ <None Include="..\..\..\examples\autodiff-texture\convert.slang">
+ <Filter>Source Files</Filter>
+ </None>
+ <None Include="..\..\..\examples\autodiff-texture\draw-quad.slang">
+ <Filter>Source Files</Filter>
+ </None>
+ <None Include="..\..\..\examples\autodiff-texture\learnmip.slang">
+ <Filter>Source Files</Filter>
+ </None>
+ <None Include="..\..\..\examples\autodiff-texture\reconstruct.slang">
+ <Filter>Source Files</Filter>
+ </None>
+ <None Include="..\..\..\examples\autodiff-texture\train.slang">
+ <Filter>Source Files</Filter>
+ </None>
+ </ItemGroup>
+</Project> \ No newline at end of file
diff --git a/examples/autodiff-texture/README.md b/examples/autodiff-texture/README.md
new file mode 100644
index 000000000..3c7328fa7
--- /dev/null
+++ b/examples/autodiff-texture/README.md
@@ -0,0 +1,4 @@
+Autodiff Texture Example
+===========================
+
+The goal of this example is to demonstrate how to use custom backward derivative functions to propagate derivatives backwards to a texture. \ No newline at end of file
diff --git a/examples/autodiff-texture/buildmip.slang b/examples/autodiff-texture/buildmip.slang
new file mode 100644
index 000000000..ebec7d754
--- /dev/null
+++ b/examples/autodiff-texture/buildmip.slang
@@ -0,0 +1,25 @@
+// A compute shader to build a mip-map layer using box filtering.
+
+cbuffer Uniforms
+{
+ uint dstWidth;
+ uint dstHeight;
+ RWTexture2D dstTexture;
+ RWTexture2D srcTexture;
+}
+
+[shader("compute")]
+[numthreads(16, 16, 1)]
+void computeMain(uint3 threadIdx : SV_DispatchThreadID)
+{
+ uint x = threadIdx.x;
+ uint y = threadIdx.y;
+ if (x >= dstWidth) return;
+ if (y >= dstHeight) return;
+ var val0 = srcTexture[uint2(x * 2, y * 2)];
+ var val1 = srcTexture[uint2(x * 2 + 1, y * 2)];
+ var val2 = srcTexture[uint2(x * 2, y * 2 + 1)];
+ var val3 = srcTexture[uint2(x * 2 + 1, y * 2 + 1)];
+
+ dstTexture[uint2(x, y)] = (val0 + val1 + val2 + val3) / 4;
+}
diff --git a/examples/autodiff-texture/checkerboard.jpg b/examples/autodiff-texture/checkerboard.jpg
new file mode 100644
index 000000000..5181d78cd
--- /dev/null
+++ b/examples/autodiff-texture/checkerboard.jpg
Binary files differ
diff --git a/examples/autodiff-texture/convert.slang b/examples/autodiff-texture/convert.slang
new file mode 100644
index 000000000..364cc286e
--- /dev/null
+++ b/examples/autodiff-texture/convert.slang
@@ -0,0 +1,25 @@
+// A compute shader to convert from a buffer into a mip-map texture.
+cbuffer Uniforms
+{
+ uint4 mipOffset[16];
+ uint dstLayer;
+ uint width;
+ uint height;
+ RWStructuredBuffer<float> srcBuffer;
+ RWTexture2D dstTexture;
+}
+
+[shader("compute")]
+[numthreads(16, 16, 1)]
+void computeMain(uint3 threadIdx : SV_DispatchThreadID)
+{
+ uint x = threadIdx.x;
+ uint y = threadIdx.y;
+ uint dstW = width >> dstLayer;
+ uint dstH = height >> dstLayer;
+ if (x >= dstW) return;
+ if (y >= dstH) return;
+ uint dstOffset = mipOffset[dstLayer / 4][dstLayer % 4] + (y * dstW + x) * 4;
+ var dstVal = float4(srcBuffer[dstOffset], srcBuffer[dstOffset + 1], srcBuffer[dstOffset + 2], srcBuffer[dstOffset + 3]);
+ dstTexture[uint2(x, y)] = dstVal;
+}
diff --git a/examples/autodiff-texture/draw-quad.slang b/examples/autodiff-texture/draw-quad.slang
new file mode 100644
index 000000000..55a33b46b
--- /dev/null
+++ b/examples/autodiff-texture/draw-quad.slang
@@ -0,0 +1,52 @@
+// Vertex and fragment shader to draw a textured quad on screen.
+
+cbuffer Uniforms
+{
+ int x;
+ int y;
+ int width;
+ int height;
+ int viewWidth;
+ int viewHeight;
+ Texture2D texture;
+ SamplerState sampler;
+}
+
+struct AssembledVertex
+{
+ float3 position : POSITION;
+};
+
+struct Fragment
+{
+ float4 color;
+};
+
+struct VertexStageOutput
+{
+ float2 uv : UV;
+ float4 sv_position : SV_Position;
+};
+
+[shader("vertex")]
+VertexStageOutput vertexMain(
+ AssembledVertex assembledVertex)
+{
+ VertexStageOutput output;
+
+ float3 position = assembledVertex.position;
+
+ output.uv = position.xy;
+ output.sv_position.x = (x + position.x * width) / (float)viewWidth * 2.0f - 1.0f;
+ output.sv_position.y = -((y + position.y * height) / (float)viewHeight * 2.0f - 1.0f);
+ output.sv_position.z = 0.5;
+ output.sv_position.w = 1.0;
+ return output;
+}
+
+[shader("fragment")]
+float4 fragmentMain(
+ float2 uv : UV) : SV_Target
+{
+ return float4(texture.Sample(sampler, uv).xyz, 1.0);
+}
diff --git a/examples/autodiff-texture/learnmip.slang b/examples/autodiff-texture/learnmip.slang
new file mode 100644
index 000000000..1434f1a66
--- /dev/null
+++ b/examples/autodiff-texture/learnmip.slang
@@ -0,0 +1,22 @@
+// A compute shader to add gradients to a mip-map texture.
+
+cbuffer Uniforms
+{
+ uint dstWidth;
+ uint dstHeight;
+ float learningRate;
+ RWTexture2D dstTexture;
+ RWTexture2D srcTexture;
+}
+
+[shader("compute")]
+[numthreads(16, 16, 1)]
+void computeMain(uint3 threadIdx : SV_DispatchThreadID)
+{
+ uint x = threadIdx.x;
+ uint y = threadIdx.y;
+ if (x >= dstWidth) return;
+ if (y >= dstHeight) return;
+ var val = srcTexture[uint2(x, y)];
+ dstTexture[uint2(x, y)] = float4((dstTexture[uint2(x, y)] - val * learningRate).xyz, 1.0);
+}
diff --git a/examples/autodiff-texture/main.cpp b/examples/autodiff-texture/main.cpp
new file mode 100644
index 000000000..7d1f809ee
--- /dev/null
+++ b/examples/autodiff-texture/main.cpp
@@ -0,0 +1,662 @@
+#include "examples/example-base/example-base.h"
+#include "gfx-util/shader-cursor.h"
+#include "slang-com-ptr.h"
+#include "slang-gfx.h"
+#include "source/core/slang-basic.h"
+#include "tools/platform/vector-math.h"
+#include "tools/platform/window.h"
+#include <slang.h>
+
+using namespace gfx;
+using namespace Slang;
+
+struct Vertex
+{
+ float position[3];
+};
+
+static const int kVertexCount = 4;
+static const Vertex kVertexData[kVertexCount] = {
+ {{0, 0, 0}},
+ {{0, 1, 0}},
+ {{1, 0, 0}},
+ {{1, 1, 0}},
+};
+
+struct AutoDiffTexture : public WindowedAppBase
+{
+
+ List<uint32_t> mipMapOffset;
+ int textureWidth;
+ int textureHeight;
+
+ void diagnoseIfNeeded(slang::IBlob* diagnosticsBlob)
+ {
+ if (diagnosticsBlob != nullptr)
+ {
+ printf("%s", (const char*)diagnosticsBlob->getBufferPointer());
+ }
+ }
+
+ gfx::Result loadRenderProgram(
+ gfx::IDevice* device, const char* fileName, const char* fragmentShader, gfx::IShaderProgram** outProgram)
+ {
+ ComPtr<slang::ISession> slangSession;
+ slangSession = device->getSlangSession();
+
+ ComPtr<slang::IBlob> diagnosticsBlob;
+ slang::IModule* module = slangSession->loadModule(fileName, diagnosticsBlob.writeRef());
+ diagnoseIfNeeded(diagnosticsBlob);
+ if (!module)
+ return SLANG_FAIL;
+
+ ComPtr<slang::IEntryPoint> vertexEntryPoint;
+ SLANG_RETURN_ON_FAIL(
+ module->findEntryPointByName("vertexMain", vertexEntryPoint.writeRef()));
+ ComPtr<slang::IEntryPoint> fragmentEntryPoint;
+ SLANG_RETURN_ON_FAIL(
+ module->findEntryPointByName(fragmentShader, fragmentEntryPoint.writeRef()));
+
+ Slang::List<slang::IComponentType*> componentTypes;
+ componentTypes.add(module);
+ int entryPointCount = 0;
+ int vertexEntryPointIndex = entryPointCount++;
+ componentTypes.add(vertexEntryPoint);
+
+ int fragmentEntryPointIndex = entryPointCount++;
+ componentTypes.add(fragmentEntryPoint);
+
+ 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.slangGlobalScope = linkedProgram;
+ SLANG_RETURN_ON_FAIL(device->createProgram(programDesc, outProgram));
+
+ return SLANG_OK;
+ }
+
+ gfx::Result loadComputeProgram(
+ gfx::IDevice* device, const char* fileName, gfx::IShaderProgram** outProgram)
+ {
+ ComPtr<slang::ISession> slangSession;
+ slangSession = device->getSlangSession();
+
+ ComPtr<slang::IBlob> diagnosticsBlob;
+ slang::IModule* module = slangSession->loadModule(fileName, diagnosticsBlob.writeRef());
+ diagnoseIfNeeded(diagnosticsBlob);
+ if (!module)
+ return SLANG_FAIL;
+
+ Slang::List<slang::IComponentType*> componentTypes;
+ componentTypes.add(module);
+ ComPtr<slang::IEntryPoint> computeEntryPoint;
+ SLANG_RETURN_ON_FAIL(
+ module->findEntryPointByName("computeMain", computeEntryPoint.writeRef()));
+ componentTypes.add(computeEntryPoint);
+
+ 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.slangGlobalScope = linkedProgram;
+ SLANG_RETURN_ON_FAIL(device->createProgram(programDesc, outProgram));
+
+ return SLANG_OK;
+ }
+
+ ComPtr<gfx::IPipelineState> gRefPipelineState;
+ ComPtr<gfx::IPipelineState> gIterPipelineState;
+ ComPtr<gfx::IPipelineState> gReconstructPipelineState;
+ ComPtr<gfx::IPipelineState> gConvertPipelineState;
+ ComPtr<gfx::IPipelineState> gBuildMipPipelineState;
+ ComPtr<gfx::IPipelineState> gLearnMipPipelineState;
+ ComPtr<gfx::IPipelineState> gDrawQuadPipelineState;
+
+ ComPtr<gfx::ITextureResource> gLearningTexture;
+ ComPtr<gfx::IResourceView> gLearningTextureSRV;
+ List<ComPtr<gfx::IResourceView>> gLearningTextureUAVs;
+
+ ComPtr<gfx::ITextureResource> gDiffTexture;
+ ComPtr<gfx::IResourceView> gDiffTextureSRV;
+ List<ComPtr<gfx::IResourceView>> gDiffTextureUAVs;
+
+ ComPtr<gfx::IBufferResource> gVertexBuffer;
+ ComPtr<gfx::IResourceView> gTexView;
+ ComPtr<gfx::ISamplerState> gSampler;
+ ComPtr<gfx::IFramebuffer> gRefFrameBuffer;
+ ComPtr<gfx::IFramebuffer> gIterFrameBuffer;
+
+ ComPtr<gfx::ITextureResource> gDepthTexture;
+ ComPtr<gfx::IResourceView> gDepthTextureView;
+
+ ComPtr<gfx::IResourceView> gIterImageDSV;
+
+ ComPtr<gfx::ITextureResource> gIterImage;
+ ComPtr<gfx::IResourceView> gIterImageSRV;
+ ComPtr<gfx::IResourceView> gIterImageRTV;
+
+ ComPtr<gfx::ITextureResource> gRefImage;
+ ComPtr<gfx::IResourceView> gRefImageSRV;
+ ComPtr<gfx::IResourceView> gRefImageRTV;
+
+ ComPtr<gfx::IBufferResource> gAccumulateBuffer;
+ ComPtr<gfx::IBufferResource> gReconstructBuffer;
+ ComPtr<gfx::IResourceView> gAccumulateBufferView;
+ ComPtr<gfx::IResourceView> gReconstructBufferView;
+
+ ClearValue kClearValue;
+ bool resetLearntTexture = false;
+
+ ComPtr<gfx::ITextureResource> createRenderTargetTexture(gfx::Format format, int w, int h, int levels)
+ {
+ gfx::ITextureResource::Desc textureDesc = {};
+ textureDesc.allowedStates.add(ResourceState::ShaderResource);
+ textureDesc.allowedStates.add(ResourceState::UnorderedAccess);
+ textureDesc.allowedStates.add(ResourceState::RenderTarget);
+ textureDesc.defaultState = ResourceState::RenderTarget;
+ textureDesc.format = format;
+ textureDesc.numMipLevels = levels;
+ textureDesc.type = gfx::IResource::Type::Texture2D;
+ textureDesc.size.width = w;
+ textureDesc.size.height = h;
+ textureDesc.size.depth = 1;
+ textureDesc.optimalClearValue = &kClearValue;
+ return gDevice->createTextureResource(textureDesc, nullptr);
+ }
+ ComPtr<gfx::ITextureResource> createDepthTexture()
+ {
+ gfx::ITextureResource::Desc textureDesc = {};
+ textureDesc.allowedStates.add(ResourceState::DepthWrite);
+ textureDesc.defaultState = ResourceState::DepthWrite;
+ textureDesc.format = gfx::Format::D32_FLOAT;
+ textureDesc.numMipLevels = 1;
+ textureDesc.type = gfx::IResource::Type::Texture2D;
+ textureDesc.size.width = windowWidth;
+ textureDesc.size.height = windowHeight;
+ textureDesc.size.depth = 1;
+ ClearValue clearValue = {};
+ textureDesc.optimalClearValue = &clearValue;
+ return gDevice->createTextureResource(textureDesc, nullptr);
+ }
+ ComPtr<gfx::IFramebuffer> createRenderTargetFramebuffer(IResourceView* tex)
+ {
+ IFramebuffer::Desc desc = {};
+ desc.layout = gFramebufferLayout.get();
+ desc.renderTargetCount = 1;
+ desc.renderTargetViews = &tex;
+ desc.depthStencilView = gDepthTextureView;
+ return gDevice->createFramebuffer(desc);
+ }
+ ComPtr<gfx::IResourceView> createRTV(ITextureResource* tex, Format f)
+ {
+ IResourceView::Desc rtvDesc = {};
+ rtvDesc.type = IResourceView::Type::RenderTarget;
+ rtvDesc.subresourceRange.mipLevelCount = 1;
+ rtvDesc.format = f;
+ rtvDesc.renderTarget.shape = gfx::IResource::Type::Texture2D;
+ return gDevice->createTextureView(tex, rtvDesc);
+ }
+ ComPtr<gfx::IResourceView> createDSV(ITextureResource* tex)
+ {
+ IResourceView::Desc dsvDesc = {};
+ dsvDesc.type = IResourceView::Type::DepthStencil;
+ dsvDesc.subresourceRange.mipLevelCount = 1;
+ dsvDesc.format = Format::D32_FLOAT;
+ dsvDesc.renderTarget.shape = gfx::IResource::Type::Texture2D;
+ return gDevice->createTextureView(tex, dsvDesc);
+ }
+ ComPtr<gfx::IResourceView> createSRV(ITextureResource* tex)
+ {
+ IResourceView::Desc rtvDesc = {};
+ rtvDesc.type = IResourceView::Type::ShaderResource;
+ return gDevice->createTextureView(tex, rtvDesc);
+ }
+ ComPtr<gfx::IPipelineState> createRenderPipelineState(
+ IInputLayout* inputLayout, IShaderProgram* program)
+ {
+ GraphicsPipelineStateDesc desc;
+ desc.inputLayout = inputLayout;
+ desc.program = program;
+ desc.rasterizer.cullMode = gfx::CullMode::None;
+ desc.framebufferLayout = gFramebufferLayout;
+ auto pipelineState = gDevice->createGraphicsPipelineState(desc);
+ return pipelineState;
+ }
+ ComPtr<gfx::IPipelineState> createComputePipelineState(IShaderProgram* program)
+ {
+ ComputePipelineStateDesc desc = {};
+ desc.program = program;
+ auto pipelineState = gDevice->createComputePipelineState(desc);
+ return pipelineState;
+ }
+ ComPtr<gfx::IResourceView> createUAV(IBufferResource* buffer)
+ {
+ IResourceView::Desc desc = {};
+ desc.type = IResourceView::Type::UnorderedAccess;
+ desc.bufferElementSize = 0;
+ return gDevice->createBufferView(buffer, nullptr, desc);
+ }
+ ComPtr<gfx::IResourceView> createUAV(ITextureResource* texture, int level)
+ {
+ IResourceView::Desc desc = {};
+ desc.type = IResourceView::Type::UnorderedAccess;
+ desc.subresourceRange.layerCount = 1;
+ desc.subresourceRange.mipLevel = level;
+ desc.subresourceRange.baseArrayLayer = 0;
+ return gDevice->createTextureView(texture,desc);
+ }
+ Slang::Result initialize()
+ {
+ initializeBase("autodiff-texture", 1024, 768);
+ srand(20421);
+
+ gWindow->events.keyPress = [this](platform::KeyEventArgs& e)
+ {
+ if (e.keyChar == 'R' || e.keyChar == 'r')
+ resetLearntTexture = true;
+ };
+
+ kClearValue.color.floatValues[0] = 0.3f;
+ kClearValue.color.floatValues[1] = 0.5f;
+ kClearValue.color.floatValues[2] = 0.7f;
+ kClearValue.color.floatValues[3] = 1.0f;
+
+ auto clientRect = gWindow->getClientRect();
+ windowWidth = clientRect.width;
+ windowHeight = clientRect.height;
+
+ InputElementDesc inputElements[] = {
+ {"POSITION", 0, Format::R32G32B32_FLOAT, offsetof(Vertex, position)}};
+ auto inputLayout = gDevice->createInputLayout(sizeof(Vertex), &inputElements[0], 1);
+ if (!inputLayout)
+ return SLANG_FAIL;
+
+ IBufferResource::Desc vertexBufferDesc;
+ vertexBufferDesc.type = IResource::Type::Buffer;
+ vertexBufferDesc.sizeInBytes = kVertexCount * sizeof(Vertex);
+ vertexBufferDesc.defaultState = ResourceState::VertexBuffer;
+ gVertexBuffer = gDevice->createBufferResource(vertexBufferDesc, &kVertexData[0]);
+ if (!gVertexBuffer)
+ return SLANG_FAIL;
+
+ {
+ ComPtr<IShaderProgram> shaderProgram;
+ SLANG_RETURN_ON_FAIL(
+ loadRenderProgram(gDevice, "train", "fragmentMain", shaderProgram.writeRef()));
+ gRefPipelineState = createRenderPipelineState(inputLayout, shaderProgram);
+ }
+ {
+ ComPtr<IShaderProgram> shaderProgram;
+ SLANG_RETURN_ON_FAIL(
+ loadRenderProgram(gDevice, "train", "diffFragmentMain", shaderProgram.writeRef()));
+ gIterPipelineState = createRenderPipelineState(inputLayout, shaderProgram);
+ }
+ {
+ ComPtr<IShaderProgram> shaderProgram;
+ SLANG_RETURN_ON_FAIL(
+ loadRenderProgram(gDevice, "draw-quad", "fragmentMain", shaderProgram.writeRef()));
+ gDrawQuadPipelineState = createRenderPipelineState(inputLayout, shaderProgram);
+ }
+ {
+ ComPtr<IShaderProgram> shaderProgram;
+ SLANG_RETURN_ON_FAIL(
+ loadComputeProgram(gDevice, "reconstruct", shaderProgram.writeRef()));
+ gReconstructPipelineState = createComputePipelineState(shaderProgram);
+ }
+ {
+ ComPtr<IShaderProgram> shaderProgram;
+ SLANG_RETURN_ON_FAIL(loadComputeProgram(gDevice, "convert", shaderProgram.writeRef()));
+ gConvertPipelineState = createComputePipelineState(shaderProgram);
+ }
+ {
+ ComPtr<IShaderProgram> shaderProgram;
+ SLANG_RETURN_ON_FAIL(loadComputeProgram(gDevice, "buildmip", shaderProgram.writeRef()));
+ gBuildMipPipelineState = createComputePipelineState(shaderProgram);
+ }
+ {
+ ComPtr<IShaderProgram> shaderProgram;
+ SLANG_RETURN_ON_FAIL(loadComputeProgram(gDevice, "learnmip", shaderProgram.writeRef()));
+ gLearnMipPipelineState = createComputePipelineState(shaderProgram);
+ }
+
+ gTexView = createTextureFromFile("checkerboard.jpg", textureWidth, textureHeight);
+ initMipOffsets(textureWidth, textureHeight);
+
+ gfx::IBufferResource::Desc bufferDesc = {};
+ bufferDesc.allowedStates.add(ResourceState::ShaderResource);
+ bufferDesc.allowedStates.add(ResourceState::UnorderedAccess);
+ bufferDesc.allowedStates.add(ResourceState::General);
+ bufferDesc.sizeInBytes = mipMapOffset.getLast() * sizeof(uint32_t);
+ bufferDesc.type = IResource::Type::Buffer;
+ gAccumulateBuffer = gDevice->createBufferResource(bufferDesc);
+ gReconstructBuffer = gDevice->createBufferResource(bufferDesc);
+
+ gAccumulateBufferView = createUAV(gAccumulateBuffer);
+ gReconstructBufferView = createUAV(gReconstructBuffer);
+
+ int mipCount = 1 + Math::Log2Ceil(Math::Max(textureWidth, textureHeight));
+ gLearningTexture = createRenderTargetTexture(
+ Format::R32G32B32A32_FLOAT,
+ textureWidth,
+ textureHeight,
+ mipCount);
+ gLearningTextureSRV = createSRV(gLearningTexture);
+ for (int i = 0; i < mipCount; i++)
+ gLearningTextureUAVs.add(createUAV(gLearningTexture, i));
+
+ gDiffTexture = createRenderTargetTexture(
+ Format::R32G32B32A32_FLOAT,
+ textureWidth,
+ textureHeight,
+ mipCount);
+ gDiffTextureSRV = createSRV(gDiffTexture);
+ for (int i = 0; i < mipCount; i++)
+ gDiffTextureUAVs.add(createUAV(gDiffTexture, i));
+
+ gfx::ISamplerState::Desc samplerDesc = {};
+ //samplerDesc.maxLOD = 0.0f;
+ gSampler = gDevice->createSamplerState(samplerDesc);
+
+ gDepthTexture = createDepthTexture();
+ gDepthTextureView = createDSV(gDepthTexture);
+
+ gRefImage = createRenderTargetTexture(Format::R8G8B8A8_UNORM, windowWidth, windowHeight, 1);
+ gRefImageRTV = createRTV(gRefImage, Format::R8G8B8A8_UNORM);
+ gRefImageSRV = createSRV(gRefImage);
+
+ gIterImage = createRenderTargetTexture(Format::R8G8B8A8_UNORM, windowWidth, windowHeight, 1);
+ gIterImageRTV = createRTV(gIterImage, Format::R8G8B8A8_UNORM);
+ gIterImageSRV = createSRV(gIterImage);
+
+ gRefFrameBuffer = createRenderTargetFramebuffer(gRefImageRTV);
+ gIterFrameBuffer = createRenderTargetFramebuffer(gIterImageRTV);
+
+ {
+ ComPtr<ICommandBuffer> commandBuffer = gTransientHeaps[0]->createCommandBuffer();
+ auto encoder = commandBuffer->encodeResourceCommands();
+ encoder->textureBarrier(gLearningTexture, ResourceState::RenderTarget, ResourceState::UnorderedAccess);
+ encoder->textureBarrier(gDiffTexture, ResourceState::RenderTarget, ResourceState::UnorderedAccess);
+ encoder->textureBarrier(gRefImage, ResourceState::RenderTarget, ResourceState::ShaderResource);
+ encoder->textureBarrier(gIterImage, ResourceState::RenderTarget, ResourceState::ShaderResource);
+ for (int i = 0; i < gLearningTextureUAVs.getCount(); i++)
+ {
+ ClearValue clearValue = {};
+ encoder->clearResourceView(gLearningTextureUAVs[i], &clearValue, ClearResourceViewFlags::None);
+ encoder->clearResourceView(gDiffTextureUAVs[i], &clearValue, ClearResourceViewFlags::None);
+ }
+ encoder->textureBarrier(gLearningTexture, ResourceState::UnorderedAccess, ResourceState::ShaderResource);
+
+ encoder->endEncoding();
+ commandBuffer->close();
+ gQueue->executeCommandBuffer(commandBuffer);
+ }
+
+ return SLANG_OK;
+ }
+
+ void initMipOffsets(int w, int h)
+ {
+ int layers = 1 + Math::Log2Ceil(Math::Max(w, h));
+ uint32_t offset = 0;
+ for (int i = 0; i < layers; i++)
+ {
+ auto lw = Math::Max(1, w >> i);
+ auto lh = Math::Max(1, h >> i);
+ mipMapOffset.add(offset);
+ offset += lw * lh * 4;
+ }
+ mipMapOffset.add(offset);
+ }
+
+ glm::mat4x4 getTransformMatrix()
+ {
+ float rotX = (rand() / (float)RAND_MAX) * 0.3f;
+ float rotY = (rand() / (float)RAND_MAX) * 0.2f;
+ glm::mat4x4 matProj = glm::perspectiveRH_ZO(
+ glm::radians(60.0f), (float)windowWidth / (float)windowHeight, 0.1f, 1000.0f);
+ auto identity = glm::mat4(1.0f);
+ auto translate = glm::translate(
+ identity,
+ glm::vec3(
+ -0.6f + 0.2f * (rand() / (float)RAND_MAX),
+ -0.6f + 0.2f * (rand() / (float)RAND_MAX),
+ -1.0f));
+ auto rot = glm::rotate(translate, -glm::pi<float>() * rotX, glm::vec3(1.0f, 0.0f, 0.0f));
+ rot = glm::rotate(rot, -glm::pi<float>() * rotY, glm::vec3(0.0f, 1.0f, 0.0f));
+ auto transformMatrix = matProj * rot;
+ transformMatrix = glm::transpose(transformMatrix);
+ return transformMatrix;
+ }
+
+ template <typename SetupPipelineFunc>
+ void renderImage(
+ int transientHeapIndex, IFramebuffer* fb, const SetupPipelineFunc& setupPipeline)
+ {
+ ComPtr<ICommandBuffer> commandBuffer =
+ gTransientHeaps[transientHeapIndex]->createCommandBuffer();
+ auto renderEncoder = commandBuffer->encodeRenderCommands(gRenderPass, fb);
+
+ gfx::Viewport viewport = {};
+ viewport.maxZ = 1.0f;
+ viewport.extentX = (float)windowWidth;
+ viewport.extentY = (float)windowHeight;
+ renderEncoder->setViewportAndScissor(viewport);
+
+ setupPipeline(renderEncoder);
+
+ renderEncoder->setVertexBuffer(0, gVertexBuffer);
+ renderEncoder->setPrimitiveTopology(PrimitiveTopology::TriangleStrip);
+
+ renderEncoder->draw(4);
+ renderEncoder->endEncoding();
+ commandBuffer->close();
+ gQueue->executeCommandBuffer(commandBuffer);
+ }
+
+ void renderReferenceImage(int transientHeapIndex, glm::mat4x4 transformMatrix)
+ {
+ {
+ ComPtr<ICommandBuffer> commandBuffer = gTransientHeaps[transientHeapIndex]->createCommandBuffer();
+ auto encoder = commandBuffer->encodeResourceCommands();
+ encoder->textureBarrier(gRefImage, ResourceState::ShaderResource, ResourceState::RenderTarget);
+ encoder->endEncoding();
+ commandBuffer->close();
+ gQueue->executeCommandBuffer(commandBuffer);
+ }
+
+ renderImage(
+ transientHeapIndex,
+ gRefFrameBuffer,
+ [&](IRenderCommandEncoder* encoder)
+ {
+ auto rootObject = encoder->bindPipeline(gRefPipelineState);
+ ShaderCursor rootCursor(rootObject);
+ rootCursor["Uniforms"]["modelViewProjection"].setData(
+ &transformMatrix, sizeof(float) * 16);
+ rootCursor["Uniforms"]["bwdTexture"]["texture"].setResource(gTexView);
+ rootCursor["Uniforms"]["sampler"].setSampler(gSampler);
+ rootCursor["Uniforms"]["mipOffset"].setData(mipMapOffset.getBuffer(), sizeof(uint32_t) * mipMapOffset.getCount());
+ rootCursor["Uniforms"]["texRef"].setResource(gTexView);
+ rootCursor["Uniforms"]["bwdTexture"]["accumulateBuffer"].setResource(gAccumulateBufferView);
+ });
+ }
+
+ virtual void renderFrame(int frameBufferIndex) override
+ {
+ static uint32_t frameCount = 0;
+ frameCount++;
+ auto transformMatrix = getTransformMatrix();
+ renderReferenceImage(frameBufferIndex, transformMatrix);
+
+ // Barriers.
+ {
+ ComPtr<ICommandBuffer> commandBuffer =
+ gTransientHeaps[frameBufferIndex]->createCommandBuffer();
+ auto resEncoder = commandBuffer->encodeResourceCommands();
+ ClearValue clearValue = {};
+ resEncoder->bufferBarrier(gAccumulateBuffer, ResourceState::Undefined, ResourceState::UnorderedAccess);
+ resEncoder->bufferBarrier(gReconstructBuffer, ResourceState::Undefined, ResourceState::UnorderedAccess);
+ resEncoder->textureBarrier(gRefImage, ResourceState::Present, ResourceState::ShaderResource);
+ resEncoder->textureBarrier(gIterImage, ResourceState::ShaderResource, ResourceState::RenderTarget);
+ resEncoder->clearResourceView(gAccumulateBufferView, &clearValue, ClearResourceViewFlags::None);
+ resEncoder->clearResourceView(gReconstructBufferView, &clearValue, ClearResourceViewFlags::None);
+ if (resetLearntTexture)
+ {
+ resEncoder->textureBarrier(gLearningTexture, ResourceState::ShaderResource, ResourceState::UnorderedAccess);
+ for (Index i =0; i <gLearningTextureUAVs.getCount(); i++)
+ resEncoder->clearResourceView(gLearningTextureUAVs[i], &clearValue, ClearResourceViewFlags::None);
+ resEncoder->textureBarrier(gLearningTexture, ResourceState::UnorderedAccess, ResourceState::ShaderResource);
+ resetLearntTexture = false;
+ }
+ resEncoder->endEncoding();
+ commandBuffer->close();
+ gQueue->executeCommandBuffer(commandBuffer);
+ }
+
+ // Render image using backward propagate shader to obtain texture-space gradients.
+ renderImage(
+ frameBufferIndex,
+ gIterFrameBuffer,
+ [&](IRenderCommandEncoder* encoder)
+ {
+ auto rootObject = encoder->bindPipeline(gIterPipelineState);
+ ShaderCursor rootCursor(rootObject);
+
+ rootCursor["Uniforms"]["modelViewProjection"].setData(
+ &transformMatrix, sizeof(float) * 16);
+ rootCursor["Uniforms"]["bwdTexture"]["texture"].setResource(gLearningTextureSRV);
+ rootCursor["Uniforms"]["sampler"].setSampler(gSampler);
+ rootCursor["Uniforms"]["mipOffset"].setData(mipMapOffset.getBuffer(), sizeof(uint32_t) * mipMapOffset.getCount());
+ rootCursor["Uniforms"]["texRef"].setResource(gRefImageSRV);
+ rootCursor["Uniforms"]["bwdTexture"]["accumulateBuffer"].setResource(gAccumulateBufferView);
+ rootCursor["Uniforms"]["bwdTexture"]["minLOD"].setData(5.0);
+
+ });
+
+ // Propagete gradients through mip map layers from top (lowest res) to bottom (highest res).
+ {
+ ComPtr<ICommandBuffer> commandBuffer =
+ gTransientHeaps[frameBufferIndex]->createCommandBuffer();
+ auto encoder = commandBuffer->encodeComputeCommands();
+ encoder->textureBarrier(gLearningTexture, ResourceState::ShaderResource, ResourceState::UnorderedAccess);
+ auto rootObject = encoder->bindPipeline(gReconstructPipelineState);
+ for (int i = (int)mipMapOffset.getCount() - 2; i >= 0; i--)
+ {
+ ShaderCursor rootCursor(rootObject);
+ rootCursor["Uniforms"]["mipOffset"].setData(mipMapOffset.getBuffer(), sizeof(uint32_t) * mipMapOffset.getCount());
+ rootCursor["Uniforms"]["dstLayer"].setData(i);
+ rootCursor["Uniforms"]["layerCount"].setData(mipMapOffset.getCount() - 1);
+ rootCursor["Uniforms"]["width"].setData(textureWidth);
+ rootCursor["Uniforms"]["height"].setData(textureHeight);
+ rootCursor["Uniforms"]["accumulateBuffer"].setResource(gAccumulateBufferView);
+ rootCursor["Uniforms"]["dstBuffer"].setResource(gReconstructBufferView);
+ encoder->dispatchCompute(
+ ((textureWidth >> i) + 15) / 16, ((textureHeight >> i) + 15) / 16, 1);
+ encoder->bufferBarrier(gReconstructBuffer, ResourceState::UnorderedAccess, ResourceState::UnorderedAccess);
+ }
+
+ // Convert bottom layer mip from buffer to texture.
+ rootObject = encoder->bindPipeline(gConvertPipelineState);
+ ShaderCursor rootCursor(rootObject);
+ rootCursor["Uniforms"]["mipOffset"].setData(mipMapOffset.getBuffer(), sizeof(uint32_t) * mipMapOffset.getCount());
+ rootCursor["Uniforms"]["dstLayer"].setData(0);
+ rootCursor["Uniforms"]["width"].setData(textureWidth);
+ rootCursor["Uniforms"]["height"].setData(textureHeight);
+ rootCursor["Uniforms"]["srcBuffer"].setResource(gReconstructBufferView);
+ rootCursor["Uniforms"]["dstTexture"].setResource(gDiffTextureUAVs[0]);
+ encoder->dispatchCompute(
+ (textureWidth + 15) / 16, (textureHeight + 15) / 16, 1);
+ encoder->textureBarrier(gDiffTexture, ResourceState::UnorderedAccess, ResourceState::UnorderedAccess);
+
+ // Build higher level mip map layers.
+ rootObject = encoder->bindPipeline(gBuildMipPipelineState);
+ for (int i = 1; i < (int)mipMapOffset.getCount() - 1; i++)
+ {
+ ShaderCursor rootCursor(rootObject);
+ rootCursor["Uniforms"]["dstWidth"].setData(textureWidth >> i);
+ rootCursor["Uniforms"]["dstHeight"].setData(textureHeight >> i);
+ rootCursor["Uniforms"]["srcTexture"].setResource(gDiffTextureUAVs[i-1]);
+ rootCursor["Uniforms"]["dstTexture"].setResource(gDiffTextureUAVs[i]);
+ encoder->dispatchCompute(
+ ((textureWidth >> i) + 15) / 16, ((textureHeight >> i) + 15) / 16, 1);
+ encoder->textureBarrier(gDiffTexture, ResourceState::UnorderedAccess, ResourceState::UnorderedAccess);
+ }
+
+ // Accumulate gradients to learnt texture.
+ rootObject = encoder->bindPipeline(gLearnMipPipelineState);
+ for (int i = 0; i < (int)mipMapOffset.getCount() - 1; i++)
+ {
+ ShaderCursor rootCursor(rootObject);
+ rootCursor["Uniforms"]["dstWidth"].setData(textureWidth >> i);
+ rootCursor["Uniforms"]["dstHeight"].setData(textureHeight >> i);
+ rootCursor["Uniforms"]["learningRate"].setData(0.1f);
+ rootCursor["Uniforms"]["srcTexture"].setResource(gDiffTextureUAVs[i]);
+ rootCursor["Uniforms"]["dstTexture"].setResource(gLearningTextureUAVs[i]);
+ encoder->dispatchCompute(
+ ((textureWidth >> i) + 15) / 16, ((textureHeight >> i) + 15) / 16, 1);
+ }
+ encoder->textureBarrier(gLearningTexture, ResourceState::UnorderedAccess, ResourceState::ShaderResource);
+ encoder->textureBarrier(gIterImage, ResourceState::Present, ResourceState::ShaderResource);
+
+ encoder->endEncoding();
+ commandBuffer->close();
+ gQueue->executeCommandBuffer(commandBuffer);
+ }
+
+ // Draw currently learnt texture.
+ {
+ ComPtr<ICommandBuffer> commandBuffer =
+ gTransientHeaps[frameBufferIndex]->createCommandBuffer();
+ auto renderEncoder = commandBuffer->encodeRenderCommands(gRenderPass, gFramebuffers[frameBufferIndex]);
+ drawTexturedQuad(renderEncoder, 0, 0, textureWidth, textureHeight, gLearningTextureSRV);
+ int refImageWidth = windowWidth - textureWidth - 10;
+ int refImageHeight = refImageWidth * windowHeight / windowWidth;
+ drawTexturedQuad(renderEncoder, textureWidth + 10, 0, refImageWidth, refImageHeight, gRefImageSRV);
+ drawTexturedQuad(renderEncoder, textureWidth + 10, refImageHeight + 10, refImageWidth, refImageHeight, gIterImageSRV);
+ renderEncoder->endEncoding();
+ commandBuffer->close();
+ gQueue->executeCommandBuffer(commandBuffer);
+ }
+
+ gSwapchain->present();
+ }
+
+ void drawTexturedQuad(IRenderCommandEncoder* renderEncoder, int x, int y, int w, int h, IResourceView* srv)
+ {
+ gfx::Viewport viewport = {};
+ viewport.maxZ = 1.0f;
+ viewport.extentX = (float)windowWidth;
+ viewport.extentY = (float)windowHeight;
+ renderEncoder->setViewportAndScissor(viewport);
+
+ auto root = renderEncoder->bindPipeline(gDrawQuadPipelineState);
+ ShaderCursor rootCursor(root);
+ rootCursor["Uniforms"]["x"].setData(x);
+ rootCursor["Uniforms"]["y"].setData(y);
+ rootCursor["Uniforms"]["width"].setData(w);
+ rootCursor["Uniforms"]["height"].setData(h);
+ rootCursor["Uniforms"]["viewWidth"].setData(windowWidth);
+ rootCursor["Uniforms"]["viewHeight"].setData(windowHeight);
+ rootCursor["Uniforms"]["texture"].setResource(srv);
+ rootCursor["Uniforms"]["sampler"].setSampler(gSampler);
+ renderEncoder->setVertexBuffer(0, gVertexBuffer);
+ renderEncoder->setPrimitiveTopology(PrimitiveTopology::TriangleStrip);
+ renderEncoder->draw(4);
+ }
+
+};
+
+PLATFORM_UI_MAIN(innerMain<AutoDiffTexture>)
diff --git a/examples/autodiff-texture/reconstruct.slang b/examples/autodiff-texture/reconstruct.slang
new file mode 100644
index 000000000..c123010e5
--- /dev/null
+++ b/examples/autodiff-texture/reconstruct.slang
@@ -0,0 +1,48 @@
+// A compute shader to propagate gradients from high level mip(low-res) to lower level mip (high-res).
+
+cbuffer Uniforms
+{
+ uint4 mipOffset[16];
+ uint dstLayer;
+ uint layerCount;
+ uint width;
+ uint height;
+ RWStructuredBuffer<int> accumulateBuffer;
+ RWStructuredBuffer<float> dstBuffer;
+}
+
+[shader("compute")]
+[numthreads(16, 16, 1)]
+void computeMain(uint3 threadIdx : SV_DispatchThreadID)
+{
+ uint x = threadIdx.x;
+ uint y = threadIdx.y;
+ uint dstW = width >> dstLayer;
+ uint dstH = height >> dstLayer;
+ if (x >= dstW) return;
+ if (y >= dstH) return;
+ uint dstOffset = mipOffset[dstLayer / 4][dstLayer % 4] + (y * dstW + x) * 4;
+ var dstVal = int4(accumulateBuffer[dstOffset], accumulateBuffer[dstOffset + 1], accumulateBuffer[dstOffset + 2], accumulateBuffer[dstOffset + 3]);
+ var newDstValToAdd = float3(0.0);
+ if (dstVal.w > 0)
+ newDstValToAdd = (float3)dstVal.xyz * float3(1.0 / (dstVal.w * 65536.0));
+
+ float4 existingVal = 0.0;
+
+ if (dstLayer < layerCount - 1 )
+ {
+ uint parentOffset = mipOffset[(dstLayer + 1) / 4][(dstLayer + 1) % 4];
+ uint parentW = dstW / 2;
+ uint parentPixelLoc = parentOffset + ((y / 2) * parentW + (x / 2)) * 4;
+ existingVal.x = dstBuffer[parentPixelLoc] * 0.25;
+ existingVal.y = dstBuffer[parentPixelLoc + 1] * 0.25;
+ existingVal.z = dstBuffer[parentPixelLoc + 2] * 0.25;
+ existingVal.w = 0.0;
+ }
+
+ var newDstVal = existingVal + float4(newDstValToAdd, 0.0);
+ dstBuffer[dstOffset] = newDstVal.x;
+ dstBuffer[dstOffset + 1] = newDstVal.y;
+ dstBuffer[dstOffset + 2] = newDstVal.z;
+ dstBuffer[dstOffset + 3] = 1.0;
+}
diff --git a/examples/autodiff-texture/train.slang b/examples/autodiff-texture/train.slang
new file mode 100644
index 000000000..16f3be35c
--- /dev/null
+++ b/examples/autodiff-texture/train.slang
@@ -0,0 +1,216 @@
+// shaders.slang
+
+struct BwdTexture
+{
+ RWStructuredBuffer<int> accumulateBuffer;
+ Texture2D texture;
+ float minLOD;
+ void writeDiffToTexel(uint offset, uint layerW, uint layerH, float x, float y, float3 val)
+ {
+ int4 uval = int4(int3(val * 65536), 1);
+ InterlockedAdd(accumulateBuffer[offset + ((uint)y * layerW + (uint)x) * 4 + 0], uval.x);
+ InterlockedAdd(accumulateBuffer[offset + ((uint)y * layerW + (uint)x) * 4 + 1], uval.y);
+ InterlockedAdd(accumulateBuffer[offset + ((uint)y * layerW + (uint)x) * 4 + 2], uval.z);
+ InterlockedAdd(accumulateBuffer[offset + ((uint)y * layerW + (uint)x) * 4 + 3], uval.w);
+ }
+
+ void broadcastDiffToLayer(uint lod, float3 diff, float2 uv, uint w, uint h)
+ {
+ var offset = mipOffset[lod / 4][lod % 4];
+ w >>= lod;
+ h >>= lod;
+ uv = uv - floor(uv);
+ float2 loc = uv * float2(w, h) - float2(0.5);
+ float x0 = floor(loc.x);
+ float y0 = floor(loc.y);
+ float fracX = loc.x - x0;
+ float fracY = loc.y - y0;
+ float x1 = x0 + 1;
+ float y1 = y0 + 1;
+ if (x0 < 0) x0 += w;
+ if (y0 < 0) y0 += h;
+ if (x1 >= w) x1 -= w;
+ if (y1 >= h) y1 -= h;
+ float weight0 = 1.0f - fracY;
+ float weight1 = fracY;
+ float weight00 = weight0 * (1.0f - fracX);
+ float weight01 = weight0 * fracX;
+ float weight10 = weight1 * (1.0f - fracX);
+ float weight11 = weight1 * fracX;
+
+ writeDiffToTexel(offset, w, h, x0, y0, weight00 * diff);
+ writeDiffToTexel(offset, w, h, x1, y0, weight01 * diff);
+ writeDiffToTexel(offset, w, h, x0, y1, weight10 * diff);
+ writeDiffToTexel(offset, w, h, x1, y1, weight11 * diff);
+ }
+
+ void sampleTexture_trilinear_bwd(uint w, uint h, uint levels, float2 uv, float2 dX, float2 dY, float4 dOut)
+ {
+ dX = dX * float2(w, h);
+ dY = dY * float2(w, h);
+
+ // Isotropic filter.
+ float lengthX = length(dX);
+ float lengthY = length(dY);
+ float LOD = log2(max(lengthX, lengthY));
+ float maxLOD = levels - 1;
+ float clampedLOD = max(minLOD, (min(maxLOD, LOD)));
+
+ float lodFrac = clampedLOD - floor(clampedLOD);
+ uint lod0 = (uint)floor(clampedLOD);
+ uint lod1 = min(levels - 1, lod0 + 1);
+ float weightLod0 = 1.0 - lodFrac;
+ float weightLod1 = lodFrac;
+ weightLod0 = 1.0;
+ broadcastDiffToLayer(lod0, dOut.xyz * weightLod0, uv, w, h);
+ broadcastDiffToLayer(lod1, dOut.xyz * weightLod1, uv, w, h);
+ }
+
+ float4 sampleTexture_linear(uint lod, float2 uv, uint w, uint h)
+ {
+ w >>= lod;
+ h >>= lod;
+ uv = uv - floor(uv);
+ float2 loc = uv * float2(w, h) - float2(0.5);
+ float x0 = floor(loc.x);
+ float y0 = floor(loc.y);
+ float fracX = loc.x - x0;
+ float fracY = loc.y - y0;
+ float x1 = x0 + 1;
+ float y1 = y0 + 1;
+ if (x0 < 0) x0 += w;
+ if (y0 < 0) y0 += h;
+ if (x1 >= w) x1 -= w;
+ if (y1 >= h) y1 -= h;
+ float weight0 = 1.0f - fracY;
+ float weight1 = fracY;
+ float weight00 = weight0 * (1.0f - fracX);
+ float weight01 = weight0 * fracX;
+ float weight10 = weight1 * (1.0f - fracX);
+ float weight11 = weight1 * fracX;
+ return texture.Load(int3(int(x0), int(y0), int(lod)), int2(0)) * weight00 +
+ texture.Load(int3(int(x1), int(y0), int(lod)), int2(0)) * weight01 +
+ texture.Load(int3(int(x0), int(y1), int(lod)), int2(0)) * weight10 +
+ texture.Load(int3(int(x1), int(y1), int(lod)), int2(0)) * weight11;
+ }
+
+ float4 sampleTexture_trilinear(uint w, uint h, uint levels, float2 uv, float2 dX, float2 dY)
+ {
+ dX = dX * float2(w, h);
+ dY = dY * float2(w, h);
+
+ // Isotropic filter.
+ float lengthX = length(dX);
+ float lengthY = length(dY);
+ float LOD = log2(max(lengthX, lengthY));
+ float maxLOD = levels - 1;
+ float clampedLOD = max(minLOD, (min(maxLOD, LOD)));
+
+ float lodFrac = clampedLOD - floor(clampedLOD);
+ uint lod0 = (uint)floor(clampedLOD);
+ uint lod1 = min(levels - 1, lod0 + 1);
+ float weightLod0 = 1.0 - lodFrac;
+ float weightLod1 = lodFrac;
+
+ let v0 = sampleTexture_linear(lod0, uv, w, h) * weightLod0;
+ let v1 = sampleTexture_linear(lod1, uv, w, h) * weightLod1;
+ return v0 + v1;
+ }
+
+ static float4 sample(BwdTexture t, SamplerState s, no_diff float2 uv)
+ {
+ return t.texture.Sample(s, uv);
+ }
+
+ [ForwardDerivativeOf(BwdTexture.sample)]
+ static float4 fwd_sample(BwdTexture t, SamplerState s, no_diff float2 uv)
+ {
+ return float4(0.0);
+ }
+
+ [BackwardDerivativeOf(BwdTexture.sample)]
+ static void bwd_sample(BwdTexture t, SamplerState s, no_diff float2 uv, float4 dOut)
+ {
+ float2 dX = ddx_coarse(uv);
+ float2 dY = ddy_coarse(uv);
+ uint w;
+ uint h;
+ uint levels;
+ t.texture.GetDimensions(0, w, h, levels);
+ t.sampleTexture_trilinear_bwd(w, h, levels, uv, dX, dY, dOut);
+ }
+}
+
+cbuffer Uniforms
+{
+ float4x4 modelViewProjection;
+ uint4 mipOffset[16];
+
+ Texture2D texRef;
+ SamplerState sampler;
+ BwdTexture bwdTexture;
+}
+
+struct AssembledVertex
+{
+ float3 position : POSITION;
+};
+
+struct Fragment
+{
+ float4 color;
+};
+
+struct VertexStageOutput
+{
+ float2 uv : UV;
+ float4 sv_position : SV_Position;
+};
+
+[BackwardDifferentiable]
+float4 shadeFragment(no_diff float2 uv)
+{
+ float3 color = BwdTexture.sample(bwdTexture, sampler, uv * 2).xyz;
+ return float4(color, 1.0);
+}
+
+[BackwardDifferentiable]
+float3 loss(no_diff float2 uv, no_diff float4 screenPos)
+{
+ float3 refColor = (no_diff texRef.Load(int3(int2(screenPos.xy), 0))).xyz;
+ float3 rs = shadeFragment(uv).xyz - refColor;
+ rs *= rs;
+ return rs;
+}
+
+[shader("vertex")]
+VertexStageOutput vertexMain(
+ AssembledVertex assembledVertex)
+{
+ VertexStageOutput output;
+
+ float3 position = assembledVertex.position;
+
+ output.uv = position.xy;
+ output.sv_position = mul(modelViewProjection, float4(position, 1.0));
+
+ return output;
+}
+
+float3 sqr(float3 v) { return v * v; }
+
+[shader("fragment")]
+float4 fragmentMain(
+ float2 uv : UV) : SV_Target
+{
+ return shadeFragment(uv);
+}
+
+[shader("fragment")]
+float4 diffFragmentMain(
+ float2 uv : UV,
+ float4 screenPos : SV_POSITION) : SV_Target
+{
+ __bwd_diff(loss)(uv, screenPos, float3(1.0));
+ return float4(loss(uv, screenPos), 1.0);
+}
diff --git a/examples/example-base/example-base.cpp b/examples/example-base/example-base.cpp
index 13df90d6d..b29e3cde5 100644
--- a/examples/example-base/example-base.cpp
+++ b/examples/example-base/example-base.cpp
@@ -6,6 +6,9 @@
#include <Windows.h>
#endif
+#define STB_IMAGE_IMPLEMENTATION
+#include "external/stb/stb_image.h"
+
using namespace Slang;
using namespace gfx;
@@ -88,7 +91,7 @@ Slang::Result WindowedAppBase::initializeBase(
renderTargetAccess.finalState = ResourceState::Present;
depthStencilAccess.loadOp = IRenderPassLayout::TargetLoadOp::Clear;
depthStencilAccess.storeOp = IRenderPassLayout::TargetStoreOp::Store;
- depthStencilAccess.initialState = ResourceState::Undefined;
+ depthStencilAccess.initialState = ResourceState::DepthWrite;
depthStencilAccess.finalState = ResourceState::DepthWrite;
renderPassDesc.renderTargetAccess = &renderTargetAccess;
renderPassDesc.depthStencilAccess = &depthStencilAccess;
@@ -121,7 +124,8 @@ void WindowedAppBase::createSwapchainFramebuffers()
depthBufferDesc.format = gfx::Format::D32_FLOAT;
depthBufferDesc.defaultState = ResourceState::DepthWrite;
depthBufferDesc.allowedStates = ResourceStateSet(ResourceState::DepthWrite);
-
+ ClearValue depthClearValue = {};
+ depthBufferDesc.optimalClearValue = &depthClearValue;
ComPtr<gfx::ITextureResource> depthBufferResource =
gDevice->createTextureResource(depthBufferDesc, nullptr);
ComPtr<gfx::ITextureResource> colorBuffer;
@@ -153,6 +157,72 @@ void WindowedAppBase::createSwapchainFramebuffers()
}
}
+ComPtr<gfx::IResourceView> WindowedAppBase::createTextureFromFile(String fileName, int& textureWidth, int& textureHeight)
+{
+ int channelsInFile = 0;
+ auto textureContent = stbi_load(fileName.getBuffer(), &textureWidth, &textureHeight, &channelsInFile, 4);
+ gfx::ITextureResource::Desc textureDesc = {};
+ textureDesc.allowedStates.add(ResourceState::ShaderResource);
+ textureDesc.format = gfx::Format::R8G8B8A8_UNORM;
+ textureDesc.numMipLevels = Math::Log2Ceil(Math::Min(textureWidth, textureHeight)) + 1;
+ textureDesc.type = gfx::IResource::Type::Texture2D;
+ textureDesc.size.width = textureWidth;
+ textureDesc.size.height = textureHeight;
+ textureDesc.size.depth = 1;
+ List<gfx::ITextureResource::SubresourceData> subresData;
+ List<List<uint32_t>> mipMapData;
+ mipMapData.setCount(textureDesc.numMipLevels);
+ subresData.setCount(textureDesc.numMipLevels);
+ mipMapData[0].setCount(textureWidth * textureHeight);
+ memcpy(mipMapData[0].getBuffer(), textureContent, textureWidth * textureHeight * 4);
+ stbi_image_free(textureContent);
+ subresData[0].data = mipMapData[0].getBuffer();
+ subresData[0].strideY = textureWidth * 4;
+ subresData[0].strideZ = textureWidth * textureHeight * 4;
+
+ // Build mipmaps.
+ struct RGBA { uint8_t v[4]; };
+ auto castToRGBA = [](uint32_t v) { RGBA result; memcpy(&result, &v, 4); return result; };
+ auto castToUint = [](RGBA v) { uint32_t result; memcpy(&result, &v, 4); return result; };
+
+ int lastMipWidth = textureWidth;
+ int lastMipHeight = textureHeight;
+ for (int m = 1; m < textureDesc.numMipLevels; m++)
+ {
+ auto lastMipmapData = mipMapData[m - 1].getBuffer();
+ int w = lastMipWidth / 2;
+ int h = lastMipHeight / 2;
+ mipMapData[m].setCount(w * h);
+ subresData[m].data = mipMapData[m].getBuffer();
+ subresData[m].strideY = w * 4;
+ subresData[m].strideZ = h * w * 4;
+ for (int x = 0; x < w; x++)
+ {
+ for (int y = 0; y < h; y++)
+ {
+ auto pix1 = castToRGBA(lastMipmapData[(y * 2) * lastMipWidth + (x * 2)]);
+ auto pix2 = castToRGBA(lastMipmapData[(y * 2) * lastMipWidth + (x * 2 + 1)]);
+ auto pix3 = castToRGBA(lastMipmapData[(y * 2 + 1) * lastMipWidth + (x * 2)]);
+ auto pix4 = castToRGBA(lastMipmapData[(y * 2 + 1) * lastMipWidth + (x * 2 + 1)]);
+ RGBA pix;
+ for (int c = 0; c < 4; c++)
+ {
+ pix.v[c] = (uint8_t)(((uint32_t)pix1.v[c] + pix2.v[c] + pix3.v[c] + pix4.v[c]) / 4);
+ }
+ mipMapData[m][y * w + x] = castToUint(pix);
+ }
+ }
+ lastMipWidth = w;
+ lastMipHeight = h;
+ }
+
+ auto texture = gDevice->createTextureResource(textureDesc, subresData.getBuffer());
+
+ gfx::IResourceView::Desc viewDesc = {};
+ viewDesc.type = gfx::IResourceView::Type::ShaderResource;
+ return gDevice->createTextureView(texture.get(), viewDesc);
+}
+
void WindowedAppBase::windowSizeChanged()
{
// Wait for the GPU to finish.
diff --git a/examples/example-base/example-base.h b/examples/example-base/example-base.h
index 49c992ca3..56eb46909 100644
--- a/examples/example-base/example-base.h
+++ b/examples/example-base/example-base.h
@@ -33,7 +33,7 @@ protected:
gfx::DeviceType deviceType = gfx::DeviceType::Default);
void createSwapchainFramebuffers();
void mainLoop();
-
+ Slang::ComPtr<gfx::IResourceView> createTextureFromFile(Slang::String fileName, int& textureWidth, int& textureHeight);
virtual void windowSizeChanged();
protected:
diff --git a/examples/heterogeneous-hello-world/README.md b/examples/heterogeneous-hello-world/README.md
deleted file mode 100644
index 709652922..000000000
--- a/examples/heterogeneous-hello-world/README.md
+++ /dev/null
@@ -1,4 +0,0 @@
-Slang "CPU Hello World Heterogeneous" Example
-===============================
-
-This example is a work-in-progress to illustrate how a heterogeneous programming example might work. It should NOT be used as a reference for working Slang code yet. \ No newline at end of file
diff --git a/examples/heterogeneous-hello-world/main.slang b/examples/heterogeneous-hello-world/main.slang
deleted file mode 100644
index b852ca1a1..000000000
--- a/examples/heterogeneous-hello-world/main.slang
+++ /dev/null
@@ -1,25 +0,0 @@
-// main.slang
-
-__target_intrinsic(cpp, "printf(\"%s\", ($0).getBuffer())")
-public void writeln(String text);
-
-[DllImport("User32")]
-int MessageBoxA(void* hwnd, String text, String caption, uint flags);
-
-[COM("111702C2-2FD7-46F9-A318-DCCEEC96357E")]
-interface IObject
-{
- int getValue(int value) throws int;
-}
-
-[DllImport("test-com")]
-IObject createObject();
-
-public __extern_cpp void main() throws int
-{
- writeln("hello world");
- MessageBoxA(nullptr, "hello world!", "example", 0);
- IObject object = createObject();
-
- int rs = try object.getValue(2);
-} \ No newline at end of file
diff --git a/premake5.lua b/premake5.lua
index 486f7366d..add7544f2 100644
--- a/premake5.lua
+++ b/premake5.lua
@@ -766,6 +766,8 @@ example "triangle"
example "ray-tracing"
example "ray-tracing-pipeline"
+example "autodiff-texture"
+
example "gpu-printing"
kind "ConsoleApp"
@@ -781,17 +783,7 @@ example "cpu-com-example"
example "cpu-hello-world"
kind "ConsoleApp"
-if enableExperimental then
- project "heterogeneous-hello-world"
- kind "ConsoleApp"
- links "slangc"
- location("build/" .. slangUtil.getBuildLocationName(targetInfo) .. "/heterogeneous-hello-world")
- prebuildcommands {
- "\"%{cfg.targetdir}/slangc\" \"%{wks.location:lower()}/examples/heterogeneous-hello-world/main.slang\" -o \"%{cfg.targetdir}/heterogeneous-hello-world.exe\""
- }
- files {"examples/heterogeneous-hello-world/*.slang"}
-end
-- Most of the other projects have more interesting configuration going
-- on, so let's walk through them in order of increasing complexity.
--
diff --git a/slang.sln b/slang.sln
index a10569253..8f682f3d3 100644
--- a/slang.sln
+++ b/slang.sln
@@ -37,6 +37,8 @@ Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "core", "build\visual-studio
EndProject
Project("{2150E333-8FDC-42A3-9474-1A3956D46DE8}") = "examples", "examples", "{EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}"
EndProject
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "autodiff-texture", "build\visual-studio\autodiff-texture\autodiff-texture.vcxproj", "{2F0A3884-1B72-FE51-84DD-12687064153E}"
+EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "cpu-com-example", "build\visual-studio\cpu-com-example\cpu-com-example.vcxproj", "{0996B38F-F512-A7D2-1E90-A7E60A6C4366}"
EndProject
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "cpu-hello-world", "build\visual-studio\cpu-hello-world\cpu-hello-world.vcxproj", "{4B47A364-37C4-96A7-6041-97BB4C1D333B}"
@@ -263,6 +265,18 @@ Global
{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|Win32.Build.0 = Release|Win32
{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|x64.ActiveCfg = Release|x64
{F9BE7957-8399-899E-0C49-E714FDDD4B65}.Release|x64.Build.0 = Release|x64
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Debug|aarch64.ActiveCfg = Debug aarch64|ARM64
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Debug|aarch64.Build.0 = Debug aarch64|ARM64
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Debug|Win32.ActiveCfg = Debug|Win32
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Debug|Win32.Build.0 = Debug|Win32
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Debug|x64.ActiveCfg = Debug|x64
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Debug|x64.Build.0 = Debug|x64
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Release|aarch64.ActiveCfg = Release aarch64|ARM64
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Release|aarch64.Build.0 = Release aarch64|ARM64
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Release|Win32.ActiveCfg = Release|Win32
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Release|Win32.Build.0 = Release|Win32
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Release|x64.ActiveCfg = Release|x64
+ {2F0A3884-1B72-FE51-84DD-12687064153E}.Release|x64.Build.0 = Release|x64
{0996B38F-F512-A7D2-1E90-A7E60A6C4366}.Debug|aarch64.ActiveCfg = Debug aarch64|ARM64
{0996B38F-F512-A7D2-1E90-A7E60A6C4366}.Debug|aarch64.Build.0 = Debug aarch64|ARM64
{0996B38F-F512-A7D2-1E90-A7E60A6C4366}.Debug|Win32.ActiveCfg = Debug|Win32
@@ -543,6 +557,7 @@ Global
{B2D63B45-92B0-40F7-B242-CCA4DFD64341} = {FD47AE19-69FD-260F-F2F1-20E65EA61D13}
{BE412850-4BB9-429A-877C-BFBC4B34186C} = {FD47AE19-69FD-260F-F2F1-20E65EA61D13}
{23149706-C12F-4329-B6AA-8266407C32D3} = {FD47AE19-69FD-260F-F2F1-20E65EA61D13}
+ {2F0A3884-1B72-FE51-84DD-12687064153E} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}
{0996B38F-F512-A7D2-1E90-A7E60A6C4366} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}
{4B47A364-37C4-96A7-6041-97BB4C1D333B} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}
{37BED5B5-23FA-D81F-8C0C-F1167867813A} = {EB5FC2C6-D72D-B6CC-C0C1-26F3AC2E9231}
diff --git a/source/slang/slang-check-decl.cpp b/source/slang/slang-check-decl.cpp
index 2613e6430..3779f48e3 100644
--- a/source/slang/slang-check-decl.cpp
+++ b/source/slang/slang-check-decl.cpp
@@ -6934,9 +6934,12 @@ namespace Slang
arg->type.isLeftValue = param->findModifier<OutModifier>() ? true : false;
arg->type.type = param->getType();
arg->loc = loc;
- if (auto pairType = visitor->getDifferentialPairType(param->getType()))
+ if (!param->findModifier<NoDiffModifier>())
{
- arg->type.type = pairType;
+ if (auto pairType = visitor->getDifferentialPairType(param->getType()))
+ {
+ arg->type.type = pairType;
+ }
}
imaginaryArguments.add(arg);
}
@@ -6958,18 +6961,26 @@ namespace Slang
arg->type.isLeftValue = param->findModifier<OutModifier>() ? true : false;
arg->type.type = param->getType();
arg->loc = loc;
- if (auto pairType = as<DifferentialPairType>(visitor->getDifferentialPairType(param->getType())))
+ bool isDiffParam = (!param->findModifier<NoDiffModifier>());
+ if (isDiffParam)
{
- arg->type.type = pairType;
- if (isOutParam(param))
+ if (auto pairType = as<DifferentialPairType>(visitor->getDifferentialPairType(param->getType())))
{
- // out T -> in T.Differential
- arg->type.isLeftValue = false;
- arg->type.type = visitor->tryGetDifferentialType(
- visitor->getASTBuilder(), pairType->getPrimalType());
+ arg->type.type = pairType;
+ if (isOutParam(param))
+ {
+ // out T -> in T.Differential
+ arg->type.isLeftValue = false;
+ arg->type.type = visitor->tryGetDifferentialType(
+ visitor->getASTBuilder(), pairType->getPrimalType());
+ }
+ }
+ else
+ {
+ isDiffParam = false;
}
}
- else
+ if (!isDiffParam)
{
if (isOutParam(param))
{
@@ -7010,7 +7021,7 @@ namespace Slang
HigherOrderInvokeExpr* higherOrderFuncExpr = visitor->getASTBuilder()->create<TDifferentiateExpr>();
higherOrderFuncExpr->baseFunction = derivativeOfAttr->funcExpr;
higherOrderFuncExpr->loc = derivativeOfAttr->loc;
- Expr* checkedHigherOrderFuncExpr = visitor->dispatchExpr(higherOrderFuncExpr, *visitor);
+ Expr* checkedHigherOrderFuncExpr = visitor->dispatchExpr(higherOrderFuncExpr, visitor->allowStaticReferenceToNonStaticMember());
if (!checkedHigherOrderFuncExpr)
{
visitor->getSink()->diagnose(derivativeOfAttr, Diagnostics::cannotResolveOriginalFunctionForDerivative);
diff --git a/source/slang/slang-ir-autodiff-fwd.cpp b/source/slang/slang-ir-autodiff-fwd.cpp
index cf45a83f5..ac4e3825a 100644
--- a/source/slang/slang-ir-autodiff-fwd.cpp
+++ b/source/slang/slang-ir-autodiff-fwd.cpp
@@ -525,10 +525,14 @@ InstPair ForwardDiffTranscriber::transcribeCall(IRBuilder* builder, IRCall* orig
return InstPair(primalCall, nullptr);
}
- auto calleeType = _getCalleeActualFuncType(diffCallee);
+ auto calleeType = _getCalleeActualFuncType(primalCallee);
SLANG_ASSERT(calleeType);
SLANG_RELEASE_ASSERT(calleeType->getParamCount() == origCall->getArgCount());
+ auto diffCalleeType = _getCalleeActualFuncType(diffCallee);
+ SLANG_ASSERT(diffCalleeType);
+ SLANG_RELEASE_ASSERT(diffCalleeType->getParamCount() == origCall->getArgCount());
+
auto placeholderCall = builder->emitCallInst(nullptr, builder->emitUndefined(builder->getTypeKind()), 0, nullptr);
builder->setInsertBefore(placeholderCall);
IRBuilder argBuilder = *builder;
@@ -545,8 +549,9 @@ InstPair ForwardDiffTranscriber::transcribeCall(IRBuilder* builder, IRCall* orig
auto origType = origCall->getArg(ii)->getDataType();
auto primalType = primalArg->getDataType();
- auto paramType = calleeType->getParamType(ii);
- if (!isNoDiffType(paramType))
+ auto originalParamType = calleeType->getParamType(ii);
+ auto diffParamType = diffCalleeType->getParamType(ii);
+ if (!isNoDiffType(originalParamType))
{
if (isNoDiffType(primalType))
{
@@ -561,7 +566,7 @@ InstPair ForwardDiffTranscriber::transcribeCall(IRBuilder* builder, IRCall* orig
auto pairValType = as<IRDifferentialPairType>(
pairPtrType ? pairPtrType->getValueType() : pairType);
auto diffType = differentiableTypeConformanceContext.getDifferentialTypeFromDiffPairType(&argBuilder, pairValType);
- if (auto ptrParamType = as<IRPtrTypeBase>(paramType))
+ if (auto ptrParamType = as<IRPtrTypeBase>(diffParamType))
{
// Create temp var to pass in/out arguments.
auto srcVar = argBuilder.emitVar(ptrParamType->getValueType());
diff --git a/tools/platform/windows/win-window.cpp b/tools/platform/windows/win-window.cpp
index 5e8c8524f..f0217e04c 100644
--- a/tools/platform/windows/win-window.cpp
+++ b/tools/platform/windows/win-window.cpp
@@ -342,7 +342,6 @@ public:
NULL, // menu
instance,
NULL);
-
if (handle)
Win32AppContext::windows[handle] = this;
}
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-02 12:59:37 +0000