yum-mirror/slang

Making it easier to work with shaders

git clone https://git.yummers.dev/yum-mirror/slang

Matthew MoultonImprove documentation and example formatting consistency (#4299)78d34f3b3

master
1.6 KiB48 linesraw
1// A compute shader to propagate gradients from high level mip(low-res) to lower level mip (high-res).
2
3cbuffer Uniforms
4{
5    uint4 mipOffset[16];
6    uint dstLayer;
7    uint layerCount;
8    uint width;
9    uint height;
10    RWStructuredBuffer<int> accumulateBuffer;
11    RWStructuredBuffer<float> dstBuffer;
12}
13
14[shader("compute")]
15[numthreads(16, 16, 1)]
16void computeMain(uint3 threadIdx : SV_DispatchThreadID)
17{
18    uint x = threadIdx.x;
19    uint y = threadIdx.y;
20    uint dstW = width >> dstLayer;
21    uint dstH = height >> dstLayer;
22    if (x >= dstW) return;
23    if (y >= dstH) return;
24    uint dstOffset = mipOffset[dstLayer / 4][dstLayer % 4] + (y * dstW + x) * 4;
25    var dstVal = int4(accumulateBuffer[dstOffset], accumulateBuffer[dstOffset + 1], accumulateBuffer[dstOffset + 2], accumulateBuffer[dstOffset + 3]);
26    var newDstValToAdd = float3(0.0);
27    if (dstVal.w > 0)
28        newDstValToAdd = (float3)dstVal.xyz * float3(1.0 / (dstVal.w * 65536.0));
29
30    float4 existingVal = 0.0;
31
32    if (dstLayer < layerCount - 1)
33    {
34        uint parentOffset = mipOffset[(dstLayer + 1) / 4][(dstLayer + 1) % 4];
35        uint parentW = dstW / 2;
36        uint parentPixelLoc = parentOffset + ((y / 2) * parentW + (x / 2)) * 4;
37        existingVal.x = dstBuffer[parentPixelLoc] * 0.25;
38        existingVal.y = dstBuffer[parentPixelLoc + 1] * 0.25;
39        existingVal.z = dstBuffer[parentPixelLoc + 2] * 0.25;
40        existingVal.w = 0.0;
41    }
42
43    var newDstVal = existingVal + float4(newDstValToAdd, 0.0);
44    dstBuffer[dstOffset] = newDstVal.x;
45    dstBuffer[dstOffset + 1] = newDstVal.y;
46    dstBuffer[dstOffset + 2] = newDstVal.z;
47    dstBuffer[dstOffset + 3] = 1.0;
48}