1#ifdef MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT2@group(0) @binding(0) var mip_0: texture_storage_2d<r64uint, read>;3#else4#ifdef MESHLET5@group(0) @binding(0) var mip_0: texture_storage_2d<r32uint, read>;6#else // MESHLET7#ifdef MULTISAMPLE8@group(0) @binding(0) var mip_0: texture_depth_multisampled_2d;9#else // MULTISAMPLE10@group(0) @binding(0) var mip_0: texture_depth_2d;11#endif // MULTISAMPLE12#endif // MESHLET13#endif // MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT14@group(0) @binding(1) var mip_1: texture_storage_2d<r32float, write>;15@group(0) @binding(2) var mip_2: texture_storage_2d<r32float, write>;16@group(0) @binding(3) var mip_3: texture_storage_2d<r32float, write>;17@group(0) @binding(4) var mip_4: texture_storage_2d<r32float, write>;18@group(0) @binding(5) var mip_5: texture_storage_2d<r32float, write>;19@group(0) @binding(6) var mip_6: texture_storage_2d<r32float, read_write>;20@group(0) @binding(7) var mip_7: texture_storage_2d<r32float, write>;21@group(0) @binding(8) var mip_8: texture_storage_2d<r32float, write>;22@group(0) @binding(9) var mip_9: texture_storage_2d<r32float, write>;23@group(0) @binding(10) var mip_10: texture_storage_2d<r32float, write>;24@group(0) @binding(11) var mip_11: texture_storage_2d<r32float, write>;25@group(0) @binding(12) var mip_12: texture_storage_2d<r32float, write>;26@group(0) @binding(13) var samplr: sampler;27struct Constants { max_mip_level: u32 }28var<immediate> constants: Constants;30/// Generates a hierarchical depth buffer.31/// Based on FidelityFX SPD v2.1 https://github.com/GPUOpen-LibrariesAndSDKs/FidelityFX-SDK/blob/d7531ae47d8b36a5d4025663e731a47a38be882f/sdk/include/FidelityFX/gpu/spd/ffx_spd.h#L52832///33/// `mip_0` may be of any size, but `mip_1` and down must have side lengths that34/// are powers of two.36// TODO:37// * Subgroup support38// * True single pass downsampling40var<workgroup> intermediate_memory: array<array<f32, 16>, 16>;42@compute43@workgroup_size(256, 1, 1)44fn downsample_depth_first(45 @builtin(workgroup_id) workgroup_id: vec3u,46 @builtin(local_invocation_index) local_invocation_index: u32,47) {48 let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);49 let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);50 let y = sub_xy.y + 8u * (local_invocation_index >> 7u);52 downsample_mips_0_and_1(x, y, workgroup_id.xy, local_invocation_index);54 downsample_mips_2_to_5(x, y, workgroup_id.xy, local_invocation_index);55}57@compute58@workgroup_size(256, 1, 1)59fn downsample_depth_second(@builtin(local_invocation_index) local_invocation_index: u32) {60 let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);61 let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);62 let y = sub_xy.y + 8u * (local_invocation_index >> 7u);64 downsample_mips_6_and_7(x, y);66 downsample_mips_8_to_11(x, y, local_invocation_index);67}69fn downsample_mips_0_and_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {70 var v: vec4f;72 var tex = vec2(workgroup_id * 64u) + vec2(x * 2u, y * 2u);73 var pix = vec2(workgroup_id * 32u) + vec2(x, y);74 v[0] = reduce_load_mip_0(tex);75 textureStore(mip_1, pix, vec4(v[0]));77 tex = vec2(workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u);78 pix = vec2(workgroup_id * 32u) + vec2(x + 16u, y);79 v[1] = reduce_load_mip_0(tex);80 textureStore(mip_1, pix, vec4(v[1]));82 tex = vec2(workgroup_id * 64u) + vec2(x * 2u, y * 2u + 32u);83 pix = vec2(workgroup_id * 32u) + vec2(x, y + 16u);84 v[2] = reduce_load_mip_0(tex);85 textureStore(mip_1, pix, vec4(v[2]));87 tex = vec2(workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u + 32u);88 pix = vec2(workgroup_id * 32u) + vec2(x + 16u, y + 16u);89 v[3] = reduce_load_mip_0(tex);90 textureStore(mip_1, pix, vec4(v[3]));92 if constants.max_mip_level <= 1u { return; }94 for (var i = 0u; i < 4u; i++) {95 intermediate_memory[x][y] = v[i];96 workgroupBarrier();97 if local_invocation_index < 64u {98 v[i] = reduce_4(vec4(99 intermediate_memory[x * 2u + 0u][y * 2u + 0u],100 intermediate_memory[x * 2u + 1u][y * 2u + 0u],101 intermediate_memory[x * 2u + 0u][y * 2u + 1u],102 intermediate_memory[x * 2u + 1u][y * 2u + 1u],103 ));104 pix = (workgroup_id * 16u) + vec2(105 x + (i % 2u) * 8u,106 y + (i / 2u) * 8u,107 );108 textureStore(mip_2, pix, vec4(v[i]));109 }110 workgroupBarrier();111 }113 if local_invocation_index < 64u {114 intermediate_memory[x + 0u][y + 0u] = v[0];115 intermediate_memory[x + 8u][y + 0u] = v[1];116 intermediate_memory[x + 0u][y + 8u] = v[2];117 intermediate_memory[x + 8u][y + 8u] = v[3];118 }119}121fn downsample_mips_2_to_5(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {122 if constants.max_mip_level <= 2u { return; }123 workgroupBarrier();124 downsample_mip_2(x, y, workgroup_id, local_invocation_index);126 if constants.max_mip_level <= 3u { return; }127 workgroupBarrier();128 downsample_mip_3(x, y, workgroup_id, local_invocation_index);130 if constants.max_mip_level <= 4u { return; }131 workgroupBarrier();132 downsample_mip_4(x, y, workgroup_id, local_invocation_index);134 if constants.max_mip_level <= 5u { return; }135 workgroupBarrier();136 downsample_mip_5(workgroup_id, local_invocation_index);137}139fn downsample_mip_2(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {140 if local_invocation_index < 64u {141 let v = reduce_4(vec4(142 intermediate_memory[x * 2u + 0u][y * 2u + 0u],143 intermediate_memory[x * 2u + 1u][y * 2u + 0u],144 intermediate_memory[x * 2u + 0u][y * 2u + 1u],145 intermediate_memory[x * 2u + 1u][y * 2u + 1u],146 ));147 textureStore(mip_3, (workgroup_id * 8u) + vec2(x, y), vec4(v));148 intermediate_memory[x * 2u + y % 2u][y * 2u] = v;149 }150}152fn downsample_mip_3(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {153 if local_invocation_index < 16u {154 let v = reduce_4(vec4(155 intermediate_memory[x * 4u + 0u + 0u][y * 4u + 0u],156 intermediate_memory[x * 4u + 2u + 0u][y * 4u + 0u],157 intermediate_memory[x * 4u + 0u + 1u][y * 4u + 2u],158 intermediate_memory[x * 4u + 2u + 1u][y * 4u + 2u],159 ));160 textureStore(mip_4, (workgroup_id * 4u) + vec2(x, y), vec4(v));161 intermediate_memory[x * 4u + y][y * 4u] = v;162 }163}165fn downsample_mip_4(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {166 if local_invocation_index < 4u {167 let v = reduce_4(vec4(168 intermediate_memory[x * 8u + 0u + 0u + y * 2u][y * 8u + 0u],169 intermediate_memory[x * 8u + 4u + 0u + y * 2u][y * 8u + 0u],170 intermediate_memory[x * 8u + 0u + 1u + y * 2u][y * 8u + 4u],171 intermediate_memory[x * 8u + 4u + 1u + y * 2u][y * 8u + 4u],172 ));173 textureStore(mip_5, (workgroup_id * 2u) + vec2(x, y), vec4(v));174 intermediate_memory[x + y * 2u][0u] = v;175 }176}178fn downsample_mip_5(workgroup_id: vec2u, local_invocation_index: u32) {179 if local_invocation_index < 1u {180 let v = reduce_4(vec4(181 intermediate_memory[0u][0u],182 intermediate_memory[1u][0u],183 intermediate_memory[2u][0u],184 intermediate_memory[3u][0u],185 ));186 textureStore(mip_6, workgroup_id, vec4(v));187 }188}190fn downsample_mips_6_and_7(x: u32, y: u32) {191 var v: vec4f;193 var tex = vec2(x * 4u + 0u, y * 4u + 0u);194 var pix = vec2(x * 2u + 0u, y * 2u + 0u);195 v[0] = reduce_load_mip_6(tex);196 textureStore(mip_7, pix, vec4(v[0]));198 tex = vec2(x * 4u + 2u, y * 4u + 0u);199 pix = vec2(x * 2u + 1u, y * 2u + 0u);200 v[1] = reduce_load_mip_6(tex);201 textureStore(mip_7, pix, vec4(v[1]));203 tex = vec2(x * 4u + 0u, y * 4u + 2u);204 pix = vec2(x * 2u + 0u, y * 2u + 1u);205 v[2] = reduce_load_mip_6(tex);206 textureStore(mip_7, pix, vec4(v[2]));208 tex = vec2(x * 4u + 2u, y * 4u + 2u);209 pix = vec2(x * 2u + 1u, y * 2u + 1u);210 v[3] = reduce_load_mip_6(tex);211 textureStore(mip_7, pix, vec4(v[3]));213 if constants.max_mip_level <= 7u { return; }215 let vr = reduce_4(v);216 textureStore(mip_8, vec2(x, y), vec4(vr));217 intermediate_memory[x][y] = vr;218}220fn downsample_mips_8_to_11(x: u32, y: u32, local_invocation_index: u32) {221 if constants.max_mip_level <= 8u { return; }222 workgroupBarrier();223 downsample_mip_8(x, y, local_invocation_index);225 if constants.max_mip_level <= 9u { return; }226 workgroupBarrier();227 downsample_mip_9(x, y, local_invocation_index);229 if constants.max_mip_level <= 10u { return; }230 workgroupBarrier();231 downsample_mip_10(x, y, local_invocation_index);233 if constants.max_mip_level <= 11u { return; }234 workgroupBarrier();235 downsample_mip_11(local_invocation_index);236}238fn downsample_mip_8(x: u32, y: u32, local_invocation_index: u32) {239 if local_invocation_index < 64u {240 let v = reduce_4(vec4(241 intermediate_memory[x * 2u + 0u][y * 2u + 0u],242 intermediate_memory[x * 2u + 1u][y * 2u + 0u],243 intermediate_memory[x * 2u + 0u][y * 2u + 1u],244 intermediate_memory[x * 2u + 1u][y * 2u + 1u],245 ));246 textureStore(mip_9, vec2(x, y), vec4(v));247 intermediate_memory[x * 2u + y % 2u][y * 2u] = v;248 }249}251fn downsample_mip_9(x: u32, y: u32, local_invocation_index: u32) {252 if local_invocation_index < 16u {253 let v = reduce_4(vec4(254 intermediate_memory[x * 4u + 0u + 0u][y * 4u + 0u],255 intermediate_memory[x * 4u + 2u + 0u][y * 4u + 0u],256 intermediate_memory[x * 4u + 0u + 1u][y * 4u + 2u],257 intermediate_memory[x * 4u + 2u + 1u][y * 4u + 2u],258 ));259 textureStore(mip_10, vec2(x, y), vec4(v));260 intermediate_memory[x * 4u + y][y * 4u] = v;261 }262}264fn downsample_mip_10(x: u32, y: u32, local_invocation_index: u32) {265 if local_invocation_index < 4u {266 let v = reduce_4(vec4(267 intermediate_memory[x * 8u + 0u + 0u + y * 2u][y * 8u + 0u],268 intermediate_memory[x * 8u + 4u + 0u + y * 2u][y * 8u + 0u],269 intermediate_memory[x * 8u + 0u + 1u + y * 2u][y * 8u + 4u],270 intermediate_memory[x * 8u + 4u + 1u + y * 2u][y * 8u + 4u],271 ));272 textureStore(mip_11, vec2(x, y), vec4(v));273 intermediate_memory[x + y * 2u][0u] = v;274 }275}277fn downsample_mip_11(local_invocation_index: u32) {278 if local_invocation_index < 1u {279 let v = reduce_4(vec4(280 intermediate_memory[0u][0u],281 intermediate_memory[1u][0u],282 intermediate_memory[2u][0u],283 intermediate_memory[3u][0u],284 ));285 textureStore(mip_12, vec2(0u, 0u), vec4(v));286 }287}289fn remap_for_wave_reduction(a: u32) -> vec2u {290 return vec2(291 insertBits(extractBits(a, 2u, 3u), a, 0u, 1u),292 insertBits(extractBits(a, 3u, 3u), extractBits(a, 1u, 2u), 0u, 2u),293 );294}296fn reduce_load_mip_0(tex: vec2u) -> f32 {297 let a = load_mip_0(tex.x, tex.y);298 let b = load_mip_0(tex.x + 1u, tex.y);299 let c = load_mip_0(tex.x, tex.y + 1u);300 let d = load_mip_0(tex.x + 1u, tex.y + 1u);301 return reduce_4(vec4(a, b, c, d));302}304fn reduce_load_mip_6(tex: vec2u) -> f32 {305 return reduce_4(vec4(306 textureLoad(mip_6, tex + vec2(0u, 0u)).r,307 textureLoad(mip_6, tex + vec2(0u, 1u)).r,308 textureLoad(mip_6, tex + vec2(1u, 0u)).r,309 textureLoad(mip_6, tex + vec2(1u, 1u)).r,310 ));311}313// Loads the top mip level at virtual position (x, y).314//315// This is the value that *would be* returned from sampling a scaled depth316// buffer with side lengths rounded up to the next power of two, without317// actually constructing such a depth buffer.318//319// See the comments in `ViewDepthPyramid::downsample_depth` for more320// information.321fn load_mip_0(x: u32, y: u32) -> f32 {322 let actual_size = textureDimensions(mip_0).xy;323 let virtual_size = vec2<u32>(324 next_power_of_two(actual_size.x),325 next_power_of_two(actual_size.y)326 );327 let virtual_uv = (vec2<f32>(f32(x), f32(y)) + 0.5) / vec2<f32>(virtual_size);328#ifdef MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT329 let virtual_st = virtual_uv * vec2<f32>(actual_size);330 let visibility = load_mip_0_meshlet(virtual_st, 32u);331 return reduce_4(visibility);332#else // MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT333#ifdef MESHLET334 let virtual_st = virtual_uv * vec2<f32>(actual_size);335 let visibility = load_mip_0_meshlet(virtual_st, 0u);336 return reduce_4(visibility);337#else // MESHLET338 // Downsample the top level.339#ifdef MULTISAMPLE340 // The top level is multisampled, so we need to loop over all the samples341 // and reduce them to 1.342 let virtual_st = virtual_uv * vec2<f32>(actual_size);343 var result = load_mip_0_single_sample(virtual_st, 0);344 let sample_count = i32(textureNumSamples(mip_0));345 for (var sample = 1; sample < sample_count; sample += 1) {346 result = min(result, load_mip_0_single_sample(virtual_st, sample));347 }348 return result;349#else // MULTISAMPLE350 return reduce_4(textureGather(mip_0, samplr, virtual_uv));351#endif // MULTISAMPLE352#endif // MESHLET353#endif // MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT354}356#ifdef MESHLET357// Loads a single 2×2 square of texels at the given position from the source358// image and returns all four (like `textureGather` does).359//360// `st` should be in texels, not in the [0, 1] range like UVs. That is, `st` is361// `uv * textureDimensions(mip_0).xy`.362fn load_mip_0_meshlet(st: vec2<f32>, shift: u32) -> vec4<f32> {363 let st0 = vec2<u32>(floor(st - 0.5));364 let st1 = st0 + 1u;365 return vec4<f32>(366 bitcast<f32>(u32(textureLoad(mip_0, vec2<u32>(st0.x, st0.y)).r) >> shift),367 bitcast<f32>(u32(textureLoad(mip_0, vec2<u32>(st0.x, st1.y)).r) >> shift),368 bitcast<f32>(u32(textureLoad(mip_0, vec2<u32>(st1.x, st0.y)).r) >> shift),369 bitcast<f32>(u32(textureLoad(mip_0, vec2<u32>(st1.x, st1.y)).r) >> shift)370 );371}372#endif // MESHLET374#ifdef MULTISAMPLE375// Loads a single 2×2 square of texels at the given position from the source376// image, reduces them, and returns the result.377//378// `st` should be in texels, not in the [0, 1] range like UVs. That is, `st` is379// `uv * textureDimensions(mip_0).xy`.380fn load_mip_0_single_sample(st: vec2<f32>, sample: i32) -> f32 {381 let st0 = vec2<u32>(floor(st - 0.5));382 let st1 = st0 + 1u;383 let v = vec4<f32>(384 textureLoad(mip_0, vec2<u32>(st0.x, st0.y), sample),385 textureLoad(mip_0, vec2<u32>(st0.x, st1.y), sample),386 textureLoad(mip_0, vec2<u32>(st1.x, st0.y), sample),387 textureLoad(mip_0, vec2<u32>(st1.x, st1.y), sample)388 );389 return reduce_4(v);390}391#endif // MULTISAMPLE393fn reduce_4(v: vec4f) -> f32 {394 return min(min(v.x, v.y), min(v.z, v.w));395}397// Returns the next power of two of x.398//399// If x is itself a power of two, this still returns the *next* power of two.400// This is different from Rust's `next_power_of_two` function.401fn next_power_of_two(x: u32) -> u32 {402 return 1u << (32u - countLeadingZeros(x));403}
Search across packages, versions, files, declarations, and comments.
Use multiple terms to narrow results (AND), and add @vertex,
@fragment, or @compute to filter by shader stage.