1// Single pass downsampling shader for creating the mip chain for an array texture2// Ported from https://github.com/GPUOpen-LibrariesAndSDKs/FidelityFX-SDK/blob/c16b1d286b5b438b75da159ab51ff426bacea3d1/sdk/include/FidelityFX/gpu/spd/ffx_spd.h4@group(0) @binding(0) var sampler_linear_clamp: sampler;5@group(0) @binding(1) var<uniform> constants: Constants;6#ifdef COMBINE_BIND_GROUP7@group(0) @binding(2) var mip_0: texture_2d_array<f32>;8@group(0) @binding(3) var mip_1: texture_storage_2d_array<rgba16float, write>;9@group(0) @binding(4) var mip_2: texture_storage_2d_array<rgba16float, write>;10@group(0) @binding(5) var mip_3: texture_storage_2d_array<rgba16float, write>;11@group(0) @binding(6) var mip_4: texture_storage_2d_array<rgba16float, write>;12@group(0) @binding(7) var mip_5: texture_storage_2d_array<rgba16float, write>;13@group(0) @binding(8) var mip_6: texture_storage_2d_array<rgba16float, read_write>;14@group(0) @binding(9) var mip_7: texture_storage_2d_array<rgba16float, write>;15@group(0) @binding(10) var mip_8: texture_storage_2d_array<rgba16float, write>;16@group(0) @binding(11) var mip_9: texture_storage_2d_array<rgba16float, write>;17@group(0) @binding(12) var mip_10: texture_storage_2d_array<rgba16float, write>;18@group(0) @binding(13) var mip_11: texture_storage_2d_array<rgba16float, write>;19@group(0) @binding(14) var mip_12: texture_storage_2d_array<rgba16float, write>;20#endif22#ifdef FIRST_PASS23@group(0) @binding(2) var mip_0: texture_2d_array<f32>;24@group(0) @binding(3) var mip_1: texture_storage_2d_array<rgba16float, write>;25@group(0) @binding(4) var mip_2: texture_storage_2d_array<rgba16float, write>;26@group(0) @binding(5) var mip_3: texture_storage_2d_array<rgba16float, write>;27@group(0) @binding(6) var mip_4: texture_storage_2d_array<rgba16float, write>;28@group(0) @binding(7) var mip_5: texture_storage_2d_array<rgba16float, write>;29@group(0) @binding(8) var mip_6: texture_storage_2d_array<rgba16float, write>;30#endif32#ifdef SECOND_PASS33@group(0) @binding(2) var mip_6: texture_2d_array<f32>;34@group(0) @binding(3) var mip_7: texture_storage_2d_array<rgba16float, write>;35@group(0) @binding(4) var mip_8: texture_storage_2d_array<rgba16float, write>;36@group(0) @binding(5) var mip_9: texture_storage_2d_array<rgba16float, write>;37@group(0) @binding(6) var mip_10: texture_storage_2d_array<rgba16float, write>;38@group(0) @binding(7) var mip_11: texture_storage_2d_array<rgba16float, write>;39@group(0) @binding(8) var mip_12: texture_storage_2d_array<rgba16float, write>;40#endif42struct Constants { mips: u32, inverse_input_size: vec2f }44var<workgroup> spd_intermediate_r: array<array<f32, 16>, 16>;45var<workgroup> spd_intermediate_g: array<array<f32, 16>, 16>;46var<workgroup> spd_intermediate_b: array<array<f32, 16>, 16>;47var<workgroup> spd_intermediate_a: array<array<f32, 16>, 16>;49@compute50@workgroup_size(256, 1, 1)51fn downsample_first(52 @builtin(workgroup_id) workgroup_id: vec3u,53 @builtin(local_invocation_index) local_invocation_index: u3254) {56 let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);57 let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);58 let y = sub_xy.y + 8u * (local_invocation_index >> 7u);60 spd_downsample_mips_0_1(x, y, workgroup_id.xy, local_invocation_index, constants.mips, workgroup_id.z);62 spd_downsample_next_four(x, y, workgroup_id.xy, local_invocation_index, 2u, constants.mips, workgroup_id.z);63}65// TODO: Once wgpu supports globallycoherent buffers, make it actually a single pass66@compute67@workgroup_size(256, 1, 1)68fn downsample_second(69 @builtin(workgroup_id) workgroup_id: vec3u,70 @builtin(local_invocation_index) local_invocation_index: u32,71) {72 let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);73 let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);74 let y = sub_xy.y + 8u * (local_invocation_index >> 7u);76 spd_downsample_mips_6_7(x, y, constants.mips, workgroup_id.z);78 spd_downsample_next_four(x, y, vec2(0u), local_invocation_index, 8u, constants.mips, workgroup_id.z);79}81fn spd_downsample_mips_0_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, mips: u32, slice: u32) {82 var v: array<vec4f, 4>;84 var tex = (workgroup_id * 64u) + vec2(x * 2u, y * 2u);85 var pix = (workgroup_id * 32u) + vec2(x, y);86 v[0] = spd_reduce_load_source_image(tex, slice);87 spd_store(pix, v[0], 0u, slice);89 tex = (workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u);90 pix = (workgroup_id * 32u) + vec2(x + 16u, y);91 v[1] = spd_reduce_load_source_image(tex, slice);92 spd_store(pix, v[1], 0u, slice);94 tex = (workgroup_id * 64u) + vec2(x * 2u, y * 2u + 32u);95 pix = (workgroup_id * 32u) + vec2(x, y + 16u);96 v[2] = spd_reduce_load_source_image(tex, slice);97 spd_store(pix, v[2], 0u, slice);99 tex = (workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u + 32u);100 pix = (workgroup_id * 32u) + vec2(x + 16u, y + 16u);101 v[3] = spd_reduce_load_source_image(tex, slice);102 spd_store(pix, v[3], 0u, slice);104 if mips <= 1u { return; }106#ifdef SUBGROUP_SUPPORT107 v[0] = spd_reduce_quad(v[0]);108 v[1] = spd_reduce_quad(v[1]);109 v[2] = spd_reduce_quad(v[2]);110 v[3] = spd_reduce_quad(v[3]);112 if local_invocation_index % 4u == 0u {113 spd_store((workgroup_id * 16u) + vec2(x / 2u, y / 2u), v[0], 1u, slice);114 spd_store_intermediate(x / 2u, y / 2u, v[0]);116 spd_store((workgroup_id * 16u) + vec2(x / 2u + 8u, y / 2u), v[1], 1u, slice);117 spd_store_intermediate(x / 2u + 8u, y / 2u, v[1]);119 spd_store((workgroup_id * 16u) + vec2(x / 2u, y / 2u + 8u), v[2], 1u, slice);120 spd_store_intermediate(x / 2u, y / 2u + 8u, v[2]);122 spd_store((workgroup_id * 16u) + vec2(x / 2u + 8u, y / 2u + 8u), v[3], 1u, slice);123 spd_store_intermediate(x / 2u + 8u, y / 2u + 8u, v[3]);124 }125#else126 for (var i = 0u; i < 4u; i++) {127 spd_store_intermediate(x, y, v[i]);128 workgroupBarrier();129 if local_invocation_index < 64u {130 v[i] = spd_reduce_intermediate(131 vec2(x * 2u + 0u, y * 2u + 0u),132 vec2(x * 2u + 1u, y * 2u + 0u),133 vec2(x * 2u + 0u, y * 2u + 1u),134 vec2(x * 2u + 1u, y * 2u + 1u),135 );136 spd_store(vec2(workgroup_id * 16) + vec2(x + (i % 2u) * 8u, y + (i / 2u) * 8u), v[i], 1u, slice);137 }138 workgroupBarrier();139 }141 if local_invocation_index < 64u {142 spd_store_intermediate(x + 0u, y + 0u, v[0]);143 spd_store_intermediate(x + 8u, y + 0u, v[1]);144 spd_store_intermediate(x + 0u, y + 8u, v[2]);145 spd_store_intermediate(x + 8u, y + 8u, v[3]);146 }147#endif148}150fn spd_downsample_next_four(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, mips: u32, slice: u32) {151 if mips <= base_mip { return; }152 workgroupBarrier();153 spd_downsample_mip_2(x, y, workgroup_id, local_invocation_index, base_mip, slice);155 if mips <= base_mip + 1u { return; }156 workgroupBarrier();157 spd_downsample_mip_3(x, y, workgroup_id, local_invocation_index, base_mip + 1u, slice);159 if mips <= base_mip + 2u { return; }160 workgroupBarrier();161 spd_downsample_mip_4(x, y, workgroup_id, local_invocation_index, base_mip + 2u, slice);163 if mips <= base_mip + 3u { return; }164 workgroupBarrier();165 spd_downsample_mip_5(x, y, workgroup_id, local_invocation_index, base_mip + 3u, slice);166}168fn spd_downsample_mip_2(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {169#ifdef SUBGROUP_SUPPORT170 var v = spd_load_intermediate(x, y);171 v = spd_reduce_quad(v);172 if local_invocation_index % 4u == 0u {173 spd_store((workgroup_id * 8u) + vec2(x / 2u, y / 2u), v, base_mip, slice);174 spd_store_intermediate(x + (y / 2u) % 2u, y, v);175 }176#else177 if local_invocation_index < 64u {178 let v = spd_reduce_intermediate(179 vec2(x * 2u + 0u, y * 2u + 0u),180 vec2(x * 2u + 1u, y * 2u + 0u),181 vec2(x * 2u + 0u, y * 2u + 1u),182 vec2(x * 2u + 1u, y * 2u + 1u),183 );184 spd_store((workgroup_id * 8u) + vec2(x, y), v, base_mip, slice);185 spd_store_intermediate(x * 2u + y % 2u, y * 2u, v);186 }187#endif188}190fn spd_downsample_mip_3(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {191#ifdef SUBGROUP_SUPPORT192 if local_invocation_index < 64u {193 var v = spd_load_intermediate(x * 2u + y % 2u, y * 2u);194 v = spd_reduce_quad(v);195 if local_invocation_index % 4u == 0u {196 spd_store((workgroup_id * 4u) + vec2(x / 2u, y / 2u), v, base_mip, slice);197 spd_store_intermediate(x * 2u + y / 2u, y * 2u, v);198 }199 }200#else201 if local_invocation_index < 16u {202 let v = spd_reduce_intermediate(203 vec2(x * 4u + 0u + 0u, y * 4u + 0u),204 vec2(x * 4u + 2u + 0u, y * 4u + 0u),205 vec2(x * 4u + 0u + 1u, y * 4u + 2u),206 vec2(x * 4u + 2u + 1u, y * 4u + 2u),207 );208 spd_store((workgroup_id * 4u) + vec2(x, y), v, base_mip, slice);209 spd_store_intermediate(x * 4u + y, y * 4u, v);210 }211#endif212}214fn spd_downsample_mip_4(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {215#ifdef SUBGROUP_SUPPORT216 if local_invocation_index < 16u {217 var v = spd_load_intermediate(x * 4u + y, y * 4u);218 v = spd_reduce_quad(v);219 if local_invocation_index % 4u == 0u {220 spd_store((workgroup_id * 2u) + vec2(x / 2u, y / 2u), v, base_mip, slice);221 spd_store_intermediate(x / 2u + y, 0u, v);222 }223 }224#else225 if local_invocation_index < 4u {226 let v = spd_reduce_intermediate(227 vec2(x * 8u + 0u + 0u + y * 2u, y * 8u + 0u),228 vec2(x * 8u + 4u + 0u + y * 2u, y * 8u + 0u),229 vec2(x * 8u + 0u + 1u + y * 2u, y * 8u + 4u),230 vec2(x * 8u + 4u + 1u + y * 2u, y * 8u + 4u),231 );232 spd_store((workgroup_id * 2u) + vec2(x, y), v, base_mip, slice);233 spd_store_intermediate(x + y * 2u, 0u, v);234 }235#endif236}238fn spd_downsample_mip_5(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {239#ifdef SUBGROUP_SUPPORT240 if local_invocation_index < 4u {241 var v = spd_load_intermediate(local_invocation_index, 0u);242 v = spd_reduce_quad(v);243 if local_invocation_index % 4u == 0u {244 spd_store(workgroup_id, v, base_mip, slice);245 }246 }247#else248 if local_invocation_index < 1u {249 let v = spd_reduce_intermediate(vec2(0u, 0u), vec2(1u, 0u), vec2(2u, 0u), vec2(3u, 0u));250 spd_store(workgroup_id, v, base_mip, slice);251 }252#endif253}255fn spd_downsample_mips_6_7(x: u32, y: u32, mips: u32, slice: u32) {256 var tex = vec2(x * 4u + 0u, y * 4u + 0u);257 var pix = vec2(x * 2u + 0u, y * 2u + 0u);258 let v0 = spd_reduce_load_4(259 vec2(x * 4u + 0u, y * 4u + 0u),260 vec2(x * 4u + 1u, y * 4u + 0u),261 vec2(x * 4u + 0u, y * 4u + 1u),262 vec2(x * 4u + 1u, y * 4u + 1u),263 slice264 );265 spd_store(pix, v0, 6u, slice);267 tex = vec2(x * 4u + 2u, y * 4u + 0u);268 pix = vec2(x * 2u + 1u, y * 2u + 0u);269 let v1 = spd_reduce_load_4(270 vec2(x * 4u + 2u, y * 4u + 0u),271 vec2(x * 4u + 3u, y * 4u + 0u),272 vec2(x * 4u + 2u, y * 4u + 1u),273 vec2(x * 4u + 3u, y * 4u + 1u),274 slice275 );276 spd_store(pix, v1, 6u, slice);278 tex = vec2(x * 4u + 0u, y * 4u + 2u);279 pix = vec2(x * 2u + 0u, y * 2u + 1u);280 let v2 = spd_reduce_load_4(281 vec2(x * 4u + 0u, y * 4u + 2u),282 vec2(x * 4u + 1u, y * 4u + 2u),283 vec2(x * 4u + 0u, y * 4u + 3u),284 vec2(x * 4u + 1u, y * 4u + 3u),285 slice286 );287 spd_store(pix, v2, 6u, slice);289 tex = vec2(x * 4u + 2u, y * 4u + 2u);290 pix = vec2(x * 2u + 1u, y * 2u + 1u);291 let v3 = spd_reduce_load_4(292 vec2(x * 4u + 2u, y * 4u + 2u),293 vec2(x * 4u + 3u, y * 4u + 2u),294 vec2(x * 4u + 2u, y * 4u + 3u),295 vec2(x * 4u + 3u, y * 4u + 3u),296 slice297 );298 spd_store(pix, v3, 6u, slice);300 if mips < 7u { return; }302 let v = spd_reduce_4(v0, v1, v2, v3);303 spd_store(vec2(x, y), v, 7u, slice);304 spd_store_intermediate(x, y, v);305}307fn remap_for_wave_reduction(a: u32) -> vec2u {308 // This function maps linear thread IDs to 2D coordinates in a special pattern309 // to ensure that neighboring threads process neighboring pixels310 // For example, this transforms linear thread IDs 0,1,2,3 into a 2×2 square312 // Extract bits to form the X and Y coordinates313 let x = insertBits(extractBits(a, 2u, 3u), a, 0u, 1u);314 let y = insertBits(extractBits(a, 3u, 3u), extractBits(a, 1u, 2u), 0u, 2u);316 return vec2u(x, y);317}319fn spd_reduce_load_source_image(uv: vec2u, slice: u32) -> vec4f {320 let texture_coord = (vec2f(uv) + 0.5) * constants.inverse_input_size;322 #ifdef COMBINE_BIND_GROUP323 let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, slice, 0.0);324 #endif325 #ifdef FIRST_PASS326 let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, slice, 0.0);327 #endif328 #ifdef SECOND_PASS329 let result = textureSampleLevel(mip_6, sampler_linear_clamp, texture_coord, slice, 0.0);330 #endif332#ifdef SRGB_CONVERSION333 return vec4(334 srgb_from_linear(result.r),335 srgb_from_linear(result.g),336 srgb_from_linear(result.b),337 result.a338 );339#else340 return result;341#endif343}345fn spd_store(pix: vec2u, value: vec4f, mip: u32, slice: u32) {346 if mip >= constants.mips { return; }347 switch mip {348 #ifdef COMBINE_BIND_GROUP349 case 0u: { textureStore(mip_1, pix, slice, value); }350 case 1u: { textureStore(mip_2, pix, slice, value); }351 case 2u: { textureStore(mip_3, pix, slice, value); }352 case 3u: { textureStore(mip_4, pix, slice, value); }353 case 4u: { textureStore(mip_5, pix, slice, value); }354 case 5u: { textureStore(mip_6, pix, slice, value); }355 case 6u: { textureStore(mip_7, pix, slice, value); }356 case 7u: { textureStore(mip_8, pix, slice, value); }357 case 8u: { textureStore(mip_9, pix, slice, value); }358 case 9u: { textureStore(mip_10, pix, slice, value); }359 case 10u: { textureStore(mip_11, pix, slice, value); }360 case 11u: { textureStore(mip_12, pix, slice, value); }361 #endif362 #ifdef FIRST_PASS363 case 0u: { textureStore(mip_1, pix, slice, value); }364 case 1u: { textureStore(mip_2, pix, slice, value); }365 case 2u: { textureStore(mip_3, pix, slice, value); }366 case 3u: { textureStore(mip_4, pix, slice, value); }367 case 4u: { textureStore(mip_5, pix, slice, value); }368 case 5u: { textureStore(mip_6, pix, slice, value); }369 #endif370 #ifdef SECOND_PASS371 case 6u: { textureStore(mip_7, pix, slice, value); }372 case 7u: { textureStore(mip_8, pix, slice, value); }373 case 8u: { textureStore(mip_9, pix, slice, value); }374 case 9u: { textureStore(mip_10, pix, slice, value); }375 case 10u: { textureStore(mip_11, pix, slice, value); }376 case 11u: { textureStore(mip_12, pix, slice, value); }377 #endif378 default: {}379 }380}382fn spd_store_intermediate(x: u32, y: u32, value: vec4f) {383 spd_intermediate_r[x][y] = value.x;384 spd_intermediate_g[x][y] = value.y;385 spd_intermediate_b[x][y] = value.z;386 spd_intermediate_a[x][y] = value.w;387}389fn spd_load_intermediate(x: u32, y: u32) -> vec4f {390 return vec4(spd_intermediate_r[x][y], spd_intermediate_g[x][y], spd_intermediate_b[x][y], spd_intermediate_a[x][y]);391}393fn spd_reduce_intermediate(i0: vec2u, i1: vec2u, i2: vec2u, i3: vec2u) -> vec4f {394 let v0 = spd_load_intermediate(i0.x, i0.y);395 let v1 = spd_load_intermediate(i1.x, i1.y);396 let v2 = spd_load_intermediate(i2.x, i2.y);397 let v3 = spd_load_intermediate(i3.x, i3.y);398 return spd_reduce_4(v0, v1, v2, v3);399}401fn spd_reduce_load_4(i0: vec2u, i1: vec2u, i2: vec2u, i3: vec2u, slice: u32) -> vec4f {402 #ifdef COMBINE_BIND_GROUP403 let v0 = textureLoad(mip_6, i0, slice);404 let v1 = textureLoad(mip_6, i1, slice);405 let v2 = textureLoad(mip_6, i2, slice);406 let v3 = textureLoad(mip_6, i3, slice);407 return spd_reduce_4(v0, v1, v2, v3);408 #endif409 #ifdef FIRST_PASS410 return vec4(0.0, 0.0, 0.0, 0.0);411 #endif412 #ifdef SECOND_PASS413 let v0 = textureLoad(mip_6, i0, slice, 0);414 let v1 = textureLoad(mip_6, i1, slice, 0);415 let v2 = textureLoad(mip_6, i2, slice, 0);416 let v3 = textureLoad(mip_6, i3, slice, 0);417 return spd_reduce_4(v0, v1, v2, v3);418 #endif419}421fn spd_reduce_4(v0: vec4f, v1: vec4f, v2: vec4f, v3: vec4f) -> vec4f {422 return (v0 + v1 + v2 + v3) * 0.25;423}425#ifdef SUBGROUP_SUPPORT426fn spd_reduce_quad(v: vec4f) -> vec4f {427 let v0 = v;428 let v1 = quadSwapX(v);429 let v2 = quadSwapY(v);430 let v3 = quadSwapDiagonal(v);431 return spd_reduce_4(v0, v1, v2, v3);432}433#endif435fn srgb_from_linear(value: f32) -> f32 {436 let j = vec3(0.0031308 * 12.92, 12.92, 1.0 / 2.4);437 let k = vec2(1.055, -0.055);438 return clamp(j.x, value * j.y, pow(value, j.z) * k.x + k.y);439}
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.