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<push_constant> 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#L52833// TODO:34// * Subgroup support35// * True single pass downsampling37var<workgroup> intermediate_memory: array<array<f32, 16>, 16>;39@compute40@workgroup_size(256, 1, 1)41fn downsample_depth_first(42 @builtin(workgroup_id) workgroup_id: vec3u,43 @builtin(local_invocation_index) local_invocation_index: u32,44) {45 let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);46 let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);47 let y = sub_xy.y + 8u * (local_invocation_index >> 7u);49 downsample_mips_0_and_1(x, y, workgroup_id.xy, local_invocation_index);51 downsample_mips_2_to_5(x, y, workgroup_id.xy, local_invocation_index);52}54@compute55@workgroup_size(256, 1, 1)56fn downsample_depth_second(@builtin(local_invocation_index) local_invocation_index: u32) {57 let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);58 let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);59 let y = sub_xy.y + 8u * (local_invocation_index >> 7u);61 downsample_mips_6_and_7(x, y);63 downsample_mips_8_to_11(x, y, local_invocation_index);64}66fn downsample_mips_0_and_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {67 var v: vec4f;69 var tex = vec2(workgroup_id * 64u) + vec2(x * 2u, y * 2u);70 var pix = vec2(workgroup_id * 32u) + vec2(x, y);71 v[0] = reduce_load_mip_0(tex);72 textureStore(mip_1, pix, vec4(v[0]));74 tex = vec2(workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u);75 pix = vec2(workgroup_id * 32u) + vec2(x + 16u, y);76 v[1] = reduce_load_mip_0(tex);77 textureStore(mip_1, pix, vec4(v[1]));79 tex = vec2(workgroup_id * 64u) + vec2(x * 2u, y * 2u + 32u);80 pix = vec2(workgroup_id * 32u) + vec2(x, y + 16u);81 v[2] = reduce_load_mip_0(tex);82 textureStore(mip_1, pix, vec4(v[2]));84 tex = vec2(workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u + 32u);85 pix = vec2(workgroup_id * 32u) + vec2(x + 16u, y + 16u);86 v[3] = reduce_load_mip_0(tex);87 textureStore(mip_1, pix, vec4(v[3]));89 if constants.max_mip_level <= 1u { return; }91 for (var i = 0u; i < 4u; i++) {92 intermediate_memory[x][y] = v[i];93 workgroupBarrier();94 if local_invocation_index < 64u {95 v[i] = reduce_4(vec4(96 intermediate_memory[x * 2u + 0u][y * 2u + 0u],97 intermediate_memory[x * 2u + 1u][y * 2u + 0u],98 intermediate_memory[x * 2u + 0u][y * 2u + 1u],99 intermediate_memory[x * 2u + 1u][y * 2u + 1u],100 ));101 pix = (workgroup_id * 16u) + vec2(102 x + (i % 2u) * 8u,103 y + (i / 2u) * 8u,104 );105 textureStore(mip_2, pix, vec4(v[i]));106 }107 workgroupBarrier();108 }110 if local_invocation_index < 64u {111 intermediate_memory[x + 0u][y + 0u] = v[0];112 intermediate_memory[x + 8u][y + 0u] = v[1];113 intermediate_memory[x + 0u][y + 8u] = v[2];114 intermediate_memory[x + 8u][y + 8u] = v[3];115 }116}118fn downsample_mips_2_to_5(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {119 if constants.max_mip_level <= 2u { return; }120 workgroupBarrier();121 downsample_mip_2(x, y, workgroup_id, local_invocation_index);123 if constants.max_mip_level <= 3u { return; }124 workgroupBarrier();125 downsample_mip_3(x, y, workgroup_id, local_invocation_index);127 if constants.max_mip_level <= 4u { return; }128 workgroupBarrier();129 downsample_mip_4(x, y, workgroup_id, local_invocation_index);131 if constants.max_mip_level <= 5u { return; }132 workgroupBarrier();133 downsample_mip_5(workgroup_id, local_invocation_index);134}136fn downsample_mip_2(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {137 if local_invocation_index < 64u {138 let v = reduce_4(vec4(139 intermediate_memory[x * 2u + 0u][y * 2u + 0u],140 intermediate_memory[x * 2u + 1u][y * 2u + 0u],141 intermediate_memory[x * 2u + 0u][y * 2u + 1u],142 intermediate_memory[x * 2u + 1u][y * 2u + 1u],143 ));144 textureStore(mip_3, (workgroup_id * 8u) + vec2(x, y), vec4(v));145 intermediate_memory[x * 2u + y % 2u][y * 2u] = v;146 }147}149fn downsample_mip_3(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {150 if local_invocation_index < 16u {151 let v = reduce_4(vec4(152 intermediate_memory[x * 4u + 0u + 0u][y * 4u + 0u],153 intermediate_memory[x * 4u + 2u + 0u][y * 4u + 0u],154 intermediate_memory[x * 4u + 0u + 1u][y * 4u + 2u],155 intermediate_memory[x * 4u + 2u + 1u][y * 4u + 2u],156 ));157 textureStore(mip_4, (workgroup_id * 4u) + vec2(x, y), vec4(v));158 intermediate_memory[x * 4u + y][y * 4u] = v;159 }160}162fn downsample_mip_4(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32) {163 if local_invocation_index < 4u {164 let v = reduce_4(vec4(165 intermediate_memory[x * 8u + 0u + 0u + y * 2u][y * 8u + 0u],166 intermediate_memory[x * 8u + 4u + 0u + y * 2u][y * 8u + 0u],167 intermediate_memory[x * 8u + 0u + 1u + y * 2u][y * 8u + 4u],168 intermediate_memory[x * 8u + 4u + 1u + y * 2u][y * 8u + 4u],169 ));170 textureStore(mip_5, (workgroup_id * 2u) + vec2(x, y), vec4(v));171 intermediate_memory[x + y * 2u][0u] = v;172 }173}175fn downsample_mip_5(workgroup_id: vec2u, local_invocation_index: u32) {176 if local_invocation_index < 1u {177 let v = reduce_4(vec4(178 intermediate_memory[0u][0u],179 intermediate_memory[1u][0u],180 intermediate_memory[2u][0u],181 intermediate_memory[3u][0u],182 ));183 textureStore(mip_6, workgroup_id, vec4(v));184 }185}187fn downsample_mips_6_and_7(x: u32, y: u32) {188 var v: vec4f;190 var tex = vec2(x * 4u + 0u, y * 4u + 0u);191 var pix = vec2(x * 2u + 0u, y * 2u + 0u);192 v[0] = reduce_load_mip_6(tex);193 textureStore(mip_7, pix, vec4(v[0]));195 tex = vec2(x * 4u + 2u, y * 4u + 0u);196 pix = vec2(x * 2u + 1u, y * 2u + 0u);197 v[1] = reduce_load_mip_6(tex);198 textureStore(mip_7, pix, vec4(v[1]));200 tex = vec2(x * 4u + 0u, y * 4u + 2u);201 pix = vec2(x * 2u + 0u, y * 2u + 1u);202 v[2] = reduce_load_mip_6(tex);203 textureStore(mip_7, pix, vec4(v[2]));205 tex = vec2(x * 4u + 2u, y * 4u + 2u);206 pix = vec2(x * 2u + 1u, y * 2u + 1u);207 v[3] = reduce_load_mip_6(tex);208 textureStore(mip_7, pix, vec4(v[3]));210 if constants.max_mip_level <= 7u { return; }212 let vr = reduce_4(v);213 textureStore(mip_8, vec2(x, y), vec4(vr));214 intermediate_memory[x][y] = vr;215}217fn downsample_mips_8_to_11(x: u32, y: u32, local_invocation_index: u32) {218 if constants.max_mip_level <= 8u { return; }219 workgroupBarrier();220 downsample_mip_8(x, y, local_invocation_index);222 if constants.max_mip_level <= 9u { return; }223 workgroupBarrier();224 downsample_mip_9(x, y, local_invocation_index);226 if constants.max_mip_level <= 10u { return; }227 workgroupBarrier();228 downsample_mip_10(x, y, local_invocation_index);230 if constants.max_mip_level <= 11u { return; }231 workgroupBarrier();232 downsample_mip_11(local_invocation_index);233}235fn downsample_mip_8(x: u32, y: u32, local_invocation_index: u32) {236 if local_invocation_index < 64u {237 let v = reduce_4(vec4(238 intermediate_memory[x * 2u + 0u][y * 2u + 0u],239 intermediate_memory[x * 2u + 1u][y * 2u + 0u],240 intermediate_memory[x * 2u + 0u][y * 2u + 1u],241 intermediate_memory[x * 2u + 1u][y * 2u + 1u],242 ));243 textureStore(mip_9, vec2(x, y), vec4(v));244 intermediate_memory[x * 2u + y % 2u][y * 2u] = v;245 }246}248fn downsample_mip_9(x: u32, y: u32, local_invocation_index: u32) {249 if local_invocation_index < 16u {250 let v = reduce_4(vec4(251 intermediate_memory[x * 4u + 0u + 0u][y * 4u + 0u],252 intermediate_memory[x * 4u + 2u + 0u][y * 4u + 0u],253 intermediate_memory[x * 4u + 0u + 1u][y * 4u + 2u],254 intermediate_memory[x * 4u + 2u + 1u][y * 4u + 2u],255 ));256 textureStore(mip_10, vec2(x, y), vec4(v));257 intermediate_memory[x * 4u + y][y * 4u] = v;258 }259}261fn downsample_mip_10(x: u32, y: u32, local_invocation_index: u32) {262 if local_invocation_index < 4u {263 let v = reduce_4(vec4(264 intermediate_memory[x * 8u + 0u + 0u + y * 2u][y * 8u + 0u],265 intermediate_memory[x * 8u + 4u + 0u + y * 2u][y * 8u + 0u],266 intermediate_memory[x * 8u + 0u + 1u + y * 2u][y * 8u + 4u],267 intermediate_memory[x * 8u + 4u + 1u + y * 2u][y * 8u + 4u],268 ));269 textureStore(mip_11, vec2(x, y), vec4(v));270 intermediate_memory[x + y * 2u][0u] = v;271 }272}274fn downsample_mip_11(local_invocation_index: u32) {275 if local_invocation_index < 1u {276 let v = reduce_4(vec4(277 intermediate_memory[0u][0u],278 intermediate_memory[1u][0u],279 intermediate_memory[2u][0u],280 intermediate_memory[3u][0u],281 ));282 textureStore(mip_12, vec2(0u, 0u), vec4(v));283 }284}286fn remap_for_wave_reduction(a: u32) -> vec2u {287 return vec2(288 insertBits(extractBits(a, 2u, 3u), a, 0u, 1u),289 insertBits(extractBits(a, 3u, 3u), extractBits(a, 1u, 2u), 0u, 2u),290 );291}293fn reduce_load_mip_0(tex: vec2u) -> f32 {294 let a = load_mip_0(tex.x, tex.y);295 let b = load_mip_0(tex.x + 1u, tex.y);296 let c = load_mip_0(tex.x, tex.y + 1u);297 let d = load_mip_0(tex.x + 1u, tex.y + 1u);298 return reduce_4(vec4(a, b, c, d));299}301fn reduce_load_mip_6(tex: vec2u) -> f32 {302 return reduce_4(vec4(303 textureLoad(mip_6, tex + vec2(0u, 0u)).r,304 textureLoad(mip_6, tex + vec2(0u, 1u)).r,305 textureLoad(mip_6, tex + vec2(1u, 0u)).r,306 textureLoad(mip_6, tex + vec2(1u, 1u)).r,307 ));308}310fn load_mip_0(x: u32, y: u32) -> f32 {311#ifdef MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT312 let visibility = textureLoad(mip_0, vec2(x, y)).r;313 return bitcast<f32>(u32(visibility >> 32u));314#else // MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT315#ifdef MESHLET316 let visibility = textureLoad(mip_0, vec2(x, y)).r;317 return bitcast<f32>(visibility);318#else // MESHLET319 // Downsample the top level.320#ifdef MULTISAMPLE321 // The top level is multisampled, so we need to loop over all the samples322 // and reduce them to 1.323 var result = textureLoad(mip_0, vec2(x, y), 0);324 let sample_count = i32(textureNumSamples(mip_0));325 for (var sample = 1; sample < sample_count; sample += 1) {326 result = min(result, textureLoad(mip_0, vec2(x, y), sample));327 }328 return result;329#else // MULTISAMPLE330 return textureLoad(mip_0, vec2(x, y), 0);331#endif // MULTISAMPLE332#endif // MESHLET333#endif // MESHLET_VISIBILITY_BUFFER_RASTER_PASS_OUTPUT334}336fn reduce_4(v: vec4f) -> f32 {337 return min(min(v.x, v.y), min(v.z, v.w));338}
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.