diff options
| author | Yong He <yonghe@outlook.com> | 2023-03-21 15:44:21 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2023-03-21 15:44:21 -0700 |
| commit | 96caba75e8dfbb879eff12cbe1a4c148a259f684 (patch) | |
| tree | 1c7b2f25484ac22c738e006334d4df559bb733a5 /examples/autodiff-texture/reconstruct.slang | |
| parent | 7f11f883d0781952f002b3aa3222a3aa0040f18a (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.slang | 48 |
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; +} |
