summaryrefslogtreecommitdiff
path: root/examples/autodiff-texture/reconstruct.slang
diff options
context:
space:
mode:
authorYong He <yonghe@outlook.com>2023-03-21 15:44:21 -0700
committerGitHub <noreply@github.com>2023-03-21 15:44:21 -0700
commit96caba75e8dfbb879eff12cbe1a4c148a259f684 (patch)
tree1c7b2f25484ac22c738e006334d4df559bb733a5 /examples/autodiff-texture/reconstruct.slang
parent7f11f883d0781952f002b3aa3222a3aa0040f18a (diff)
Add texture tri-linear autodiff example. (#2715)
* Add quad texture example. * delete output image * remove irrelavent files * update project files * fix * Update example. * Fix. * remove out-texture --------- Co-authored-by: Yong He <yhe@nvidia.com>
Diffstat (limited to 'examples/autodiff-texture/reconstruct.slang')
-rw-r--r--examples/autodiff-texture/reconstruct.slang48
1 files changed, 48 insertions, 0 deletions
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;
+}