1//! Generates a hierarchical depth buffer.2//! Based on FidelityFX SPD v2.1 https://github.com/GPUOpen-LibrariesAndSDKs/FidelityFX-SDK/blob/d7531ae47d8b36a5d4025663e731a47a38be882f/sdk/include/FidelityFX/gpu/spd/ffx_spd.h#L5283//!4//! `mip_0` may be of any size, but `mip_1` and down must have side lengths that5//! are powers of two.7@if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT)8@group(0) @binding(0) var mip_0: texture_storage_2d<r64uint, read>;9@elif(MESHLET)10@group(0) @binding(0) var mip_0: texture_storage_2d<r32uint, read>;11@elif(MULTISAMPLE)12@group(0) @binding(0) var mip_0: texture_depth_multisampled_2d;13@else14@group(0) @binding(0) var mip_0: texture_depth_2d;15@group(0) @binding(1) var mip_1: texture_storage_2d<r32float, write>;16@group(0) @binding(2) var mip_2: texture_storage_2d<r32float, write>;17@group(0) @binding(3) var mip_3: texture_storage_2d<r32float, write>;18@group(0) @binding(4) var mip_4: texture_storage_2d<r32float, write>;19@group(0) @binding(5) var mip_5: texture_storage_2d<r32float, write>;20@group(0) @binding(6) var mip_6: texture_storage_2d<r32float, read_write>;21@group(0) @binding(7) var mip_7: texture_storage_2d<r32float, write>;22@group(0) @binding(8) var mip_8: texture_storage_2d<r32float, write>;23@group(0) @binding(9) var mip_9: texture_storage_2d<r32float, write>;24@group(0) @binding(10) var mip_10: texture_storage_2d<r32float, write>;25@group(0) @binding(11) var mip_11: texture_storage_2d<r32float, write>;26@group(0) @binding(12) var mip_12: texture_storage_2d<r32float, write>;27@group(0) @binding(13) var samplr: sampler;28struct Constants { max_mip_level: u32 }29var<immediate> constants: Constants;31// TODO:32// * Subgroup support33// * True single pass downsampling35var<workgroup> intermediate_memory: array<array<f32, 16>, 16>;37@compute38@workgroup_size(256, 1, 1)39fn downsample_depth_first(40 @builtin(workgroup_id) workgroup_id: vec3u,41 @builtin(local_invocation_index) local_invocation_index: u32,42) {43 let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);44 let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);45 let y = sub_xy.y + 8u * (local_invocation_index >> 7u);47 downsample_mips_0_and_1(x, y, workgroup_id.xy, local_invocation_index);49 downsample_mips_2_to_5(x, y, workgroup_id.xy, local_invocation_index);50}52@compute53@workgroup_size(256, 1, 1)54fn downsample_depth_second(@builtin(local_invocation_index) local_invocation_index: u32) {55 let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);56 let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);57 let y = sub_xy.y + 8u * (local_invocation_index >> 7u);59 downsample_mips_6_and_7(x, y);61 downsample_mips_8_to_11(x, y, local_invocation_index);62}64fn downsample_mips_0_and_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {65 var v: vec4f;67 var tex = vec2(workgroup_id * 64u) + vec2(x * 2u, y * 2u);68 var pix = vec2(workgroup_id * 32u) + vec2(x, y);69 v[0] = reduce_load_mip_0(tex);70 textureStore(mip_1, pix, vec4(v[0]));72 tex = vec2(workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u);73 pix = vec2(workgroup_id * 32u) + vec2(x + 16u, y);74 v[1] = reduce_load_mip_0(tex);75 textureStore(mip_1, pix, vec4(v[1]));77 tex = vec2(workgroup_id * 64u) + vec2(x * 2u, y * 2u + 32u);78 pix = vec2(workgroup_id * 32u) + vec2(x, y + 16u);79 v[2] = reduce_load_mip_0(tex);80 textureStore(mip_1, pix, vec4(v[2]));82 tex = vec2(workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u + 32u);83 pix = vec2(workgroup_id * 32u) + vec2(x + 16u, y + 16u);84 v[3] = reduce_load_mip_0(tex);85 textureStore(mip_1, pix, vec4(v[3]));87 if constants.max_mip_level <= 1u { return; }89 for (var i = 0u; i < 4u; i++) {90 intermediate_memory[x][y] = v[i];91 workgroupBarrier();92 if local_invocation_index < 64u {93 v[i] = reduce_4(vec4(94 intermediate_memory[x * 2u + 0u][y * 2u + 0u],95 intermediate_memory[x * 2u + 1u][y * 2u + 0u],96 intermediate_memory[x * 2u + 0u][y * 2u + 1u],97 intermediate_memory[x * 2u + 1u][y * 2u + 1u],98 ));99 pix = (workgroup_id * 16u) + vec2(100 x + (i % 2u) * 8u,101 y + (i / 2u) * 8u,102 );103 textureStore(mip_2, pix, vec4(v[i]));104 }105 workgroupBarrier();106 }108 if local_invocation_index < 64u {109 intermediate_memory[x + 0u][y + 0u] = v[0];110 intermediate_memory[x + 8u][y + 0u] = v[1];111 intermediate_memory[x + 0u][y + 8u] = v[2];112 intermediate_memory[x + 8u][y + 8u] = v[3];113 }114}116fn downsample_mips_2_to_5(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {117 if constants.max_mip_level <= 2u { return; }118 workgroupBarrier();119 downsample_mip_2(x, y, workgroup_id, local_invocation_index);121 if constants.max_mip_level <= 3u { return; }122 workgroupBarrier();123 downsample_mip_3(x, y, workgroup_id, local_invocation_index);125 if constants.max_mip_level <= 4u { return; }126 workgroupBarrier();127 downsample_mip_4(x, y, workgroup_id, local_invocation_index);129 if constants.max_mip_level <= 5u { return; }130 workgroupBarrier();131 downsample_mip_5(workgroup_id, local_invocation_index);132}134fn downsample_mip_2(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {135 if local_invocation_index < 64u {136 let v = reduce_4(vec4(137 intermediate_memory[x * 2u + 0u][y * 2u + 0u],138 intermediate_memory[x * 2u + 1u][y * 2u + 0u],139 intermediate_memory[x * 2u + 0u][y * 2u + 1u],140 intermediate_memory[x * 2u + 1u][y * 2u + 1u],141 ));142 textureStore(mip_3, (workgroup_id * 8u) + vec2(x, y), vec4(v));143 intermediate_memory[x * 2u + y % 2u][y * 2u] = v;144 }145}147fn downsample_mip_3(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {148 if local_invocation_index < 16u {149 let v = reduce_4(vec4(150 intermediate_memory[x * 4u + 0u + 0u][y * 4u + 0u],151 intermediate_memory[x * 4u + 2u + 0u][y * 4u + 0u],152 intermediate_memory[x * 4u + 0u + 1u][y * 4u + 2u],153 intermediate_memory[x * 4u + 2u + 1u][y * 4u + 2u],154 ));155 textureStore(mip_4, (workgroup_id * 4u) + vec2(x, y), vec4(v));156 intermediate_memory[x * 4u + y][y * 4u] = v;157 }158}160fn downsample_mip_4(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {161 if local_invocation_index < 4u {162 let v = reduce_4(vec4(163 intermediate_memory[x * 8u + 0u + 0u + y * 2u][y * 8u + 0u],164 intermediate_memory[x * 8u + 4u + 0u + y * 2u][y * 8u + 0u],165 intermediate_memory[x * 8u + 0u + 1u + y * 2u][y * 8u + 4u],166 intermediate_memory[x * 8u + 4u + 1u + y * 2u][y * 8u + 4u],167 ));168 textureStore(mip_5, (workgroup_id * 2u) + vec2(x, y), vec4(v));169 intermediate_memory[x + y * 2u][0u] = v;170 }171}173fn downsample_mip_5(workgroup_id: vec2u, local_invocation_index: u32) {174 if local_invocation_index < 1u {175 let v = reduce_4(vec4(176 intermediate_memory[0u][0u],177 intermediate_memory[1u][0u],178 intermediate_memory[2u][0u],179 intermediate_memory[3u][0u],180 ));181 textureStore(mip_6, workgroup_id, vec4(v));182 }183}185fn downsample_mips_6_and_7(x: u32, y: u32) {186 var v: vec4f;188 var tex = vec2(x * 4u + 0u, y * 4u + 0u);189 var pix = vec2(x * 2u + 0u, y * 2u + 0u);190 v[0] = reduce_load_mip_6(tex);191 textureStore(mip_7, pix, vec4(v[0]));193 tex = vec2(x * 4u + 2u, y * 4u + 0u);194 pix = vec2(x * 2u + 1u, y * 2u + 0u);195 v[1] = reduce_load_mip_6(tex);196 textureStore(mip_7, pix, vec4(v[1]));198 tex = vec2(x * 4u + 0u, y * 4u + 2u);199 pix = vec2(x * 2u + 0u, y * 2u + 1u);200 v[2] = reduce_load_mip_6(tex);201 textureStore(mip_7, pix, vec4(v[2]));203 tex = vec2(x * 4u + 2u, y * 4u + 2u);204 pix = vec2(x * 2u + 1u, y * 2u + 1u);205 v[3] = reduce_load_mip_6(tex);206 textureStore(mip_7, pix, vec4(v[3]));208 if constants.max_mip_level <= 7u { return; }210 let vr = reduce_4(v);211 textureStore(mip_8, vec2(x, y), vec4(vr));212 intermediate_memory[x][y] = vr;213}215fn downsample_mips_8_to_11(x: u32, y: u32, local_invocation_index: u32) {216 if constants.max_mip_level <= 8u { return; }217 workgroupBarrier();218 downsample_mip_8(x, y, local_invocation_index);220 if constants.max_mip_level <= 9u { return; }221 workgroupBarrier();222 downsample_mip_9(x, y, local_invocation_index);224 if constants.max_mip_level <= 10u { return; }225 workgroupBarrier();226 downsample_mip_10(x, y, local_invocation_index);228 if constants.max_mip_level <= 11u { return; }229 workgroupBarrier();230 downsample_mip_11(local_invocation_index);231}233fn downsample_mip_8(x: u32, y: u32, local_invocation_index: u32) {234 if local_invocation_index < 64u {235 let v = reduce_4(vec4(236 intermediate_memory[x * 2u + 0u][y * 2u + 0u],237 intermediate_memory[x * 2u + 1u][y * 2u + 0u],238 intermediate_memory[x * 2u + 0u][y * 2u + 1u],239 intermediate_memory[x * 2u + 1u][y * 2u + 1u],240 ));241 textureStore(mip_9, vec2(x, y), vec4(v));242 intermediate_memory[x * 2u + y % 2u][y * 2u] = v;243 }244}246fn downsample_mip_9(x: u32, y: u32, local_invocation_index: u32) {247 if local_invocation_index < 16u {248 let v = reduce_4(vec4(249 intermediate_memory[x * 4u + 0u + 0u][y * 4u + 0u],250 intermediate_memory[x * 4u + 2u + 0u][y * 4u + 0u],251 intermediate_memory[x * 4u + 0u + 1u][y * 4u + 2u],252 intermediate_memory[x * 4u + 2u + 1u][y * 4u + 2u],253 ));254 textureStore(mip_10, vec2(x, y), vec4(v));255 intermediate_memory[x * 4u + y][y * 4u] = v;256 }257}259fn downsample_mip_10(x: u32, y: u32, local_invocation_index: u32) {260 if local_invocation_index < 4u {261 let v = reduce_4(vec4(262 intermediate_memory[x * 8u + 0u + 0u + y * 2u][y * 8u + 0u],263 intermediate_memory[x * 8u + 4u + 0u + y * 2u][y * 8u + 0u],264 intermediate_memory[x * 8u + 0u + 1u + y * 2u][y * 8u + 4u],265 intermediate_memory[x * 8u + 4u + 1u + y * 2u][y * 8u + 4u],266 ));267 textureStore(mip_11, vec2(x, y), vec4(v));268 intermediate_memory[x + y * 2u][0u] = v;269 }270}272fn downsample_mip_11(local_invocation_index: u32) {273 if local_invocation_index < 1u {274 let v = reduce_4(vec4(275 intermediate_memory[0u][0u],276 intermediate_memory[1u][0u],277 intermediate_memory[2u][0u],278 intermediate_memory[3u][0u],279 ));280 textureStore(mip_12, vec2(0u, 0u), vec4(v));281 }282}284fn remap_for_wave_reduction(a: u32) -> vec2u {285 return vec2(286 insertBits(extractBits(a, 2u, 3u), a, 0u, 1u),287 insertBits(extractBits(a, 3u, 3u), extractBits(a, 1u, 2u), 0u, 2u),288 );289}291fn reduce_load_mip_0(tex: vec2u) -> f32 {292 let a = load_mip_0(tex.x, tex.y);293 let b = load_mip_0(tex.x + 1u, tex.y);294 let c = load_mip_0(tex.x, tex.y + 1u);295 let d = load_mip_0(tex.x + 1u, tex.y + 1u);296 return reduce_4(vec4(a, b, c, d));297}299fn reduce_load_mip_6(tex: vec2u) -> f32 {300 return reduce_4(vec4(301 textureLoad(mip_6, tex + vec2(0u, 0u)).r,302 textureLoad(mip_6, tex + vec2(0u, 1u)).r,303 textureLoad(mip_6, tex + vec2(1u, 0u)).r,304 textureLoad(mip_6, tex + vec2(1u, 1u)).r,305 ));306}308/// Loads the top mip level at virtual position (x, y).309///310/// This is the value that *would be* returned from sampling a scaled depth311/// buffer with side lengths rounded up to the next power of two, without312/// actually constructing such a depth buffer.313///314/// See the comments in `ViewDepthPyramid::downsample_depth` for more315/// information.316fn load_mip_0(x: u32, y: u32) -> f32 {317 let actual_size = textureDimensions(mip_0).xy;318 let virtual_size = vec2<u32>(319 next_power_of_two(actual_size.x),320 next_power_of_two(actual_size.y)321 );322 let virtual_uv = (vec2<f32>(f32(x), f32(y)) + 0.5) / vec2<f32>(virtual_size);323 @if(MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT) {324 let virtual_st = virtual_uv * vec2<f32>(actual_size);325 let visibility = load_mip_0_meshlet(virtual_st, 32u);326 return reduce_4(visibility);327 } @elif(MESHLET) {328 let virtual_st = virtual_uv * vec2<f32>(actual_size);329 let visibility = load_mip_0_meshlet(virtual_st, 0u);330 return reduce_4(visibility);331 } @elif(MULTISAMPLE) {332 // Downsample the top level.333 // The top level is multisampled, so we need to loop over all the samples334 // and reduce them to 1.335 let virtual_st = virtual_uv * vec2<f32>(actual_size);336 var result = load_mip_0_single_sample(virtual_st, 0);337 let sample_count = i32(textureNumSamples(mip_0));338 for (var sample = 1; sample < sample_count; sample += 1) {339 result = min(result, load_mip_0_single_sample(virtual_st, sample));340 }341 return result;342 } @else {343 return reduce_4(textureGather(mip_0, samplr, virtual_uv));344 }345}347/// Loads a single 2×2 square of texels at the given position from the source348/// image and returns all four (like `textureGather` does).349///350/// `st` should be in texels, not in the [0, 1] range like UVs. That is, `st` is351/// `uv * textureDimensions(mip_0).xy`.352@if(MESHLET)353fn load_mip_0_meshlet(st: vec2<f32>, shift: u32) -> vec4<f32> {354 let st0 = vec2<u32>(floor(st - 0.5));355 let st1 = st0 + 1u;356 return vec4<f32>(357 bitcast<f32>(u32(textureLoad(mip_0, vec2<u32>(st0.x, st0.y)).r) >> shift),358 bitcast<f32>(u32(textureLoad(mip_0, vec2<u32>(st0.x, st1.y)).r) >> shift),359 bitcast<f32>(u32(textureLoad(mip_0, vec2<u32>(st1.x, st0.y)).r) >> shift),360 bitcast<f32>(u32(textureLoad(mip_0, vec2<u32>(st1.x, st1.y)).r) >> shift)361 );362}364/// Loads a single 2×2 square of texels at the given position from the source365/// image, reduces them, and returns the result.366///367/// `st` should be in texels, not in the [0, 1] range like UVs. That is, `st` is368/// `uv * textureDimensions(mip_0).xy`.369@if(MULTISAMPLE)370fn load_mip_0_single_sample(st: vec2<f32>, sample: i32) -> f32 {371 let st0 = vec2<u32>(floor(st - 0.5));372 let st1 = st0 + 1u;373 let v = vec4<f32>(374 textureLoad(mip_0, vec2<u32>(st0.x, st0.y), sample),375 textureLoad(mip_0, vec2<u32>(st0.x, st1.y), sample),376 textureLoad(mip_0, vec2<u32>(st1.x, st0.y), sample),377 textureLoad(mip_0, vec2<u32>(st1.x, st1.y), sample)378 );379 return reduce_4(v);380}382fn reduce_4(v: vec4f) -> f32 {383 return min(min(v.x, v.y), min(v.z, v.w));384}386/// Returns the next power of two of x.387///388/// If x is itself a power of two, this still returns the *next* power of two.389/// This is different from Rust's `next_power_of_two` function.390fn next_power_of_two(x: u32) -> u32 {391 return 1u << (32u - countLeadingZeros(x));392}
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.