yum-mirror/slang
Making it easier to work with shaders
git clone https://git.yummers.dev/yum-mirror/slang
78d34f3b3
master
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}