1// Single pass downsampling shader for creating the mip chain for a texture
2// Ported from https://github.com/GPUOpen-LibrariesAndSDKs/FidelityFX-SDK/blob/c16b1d286b5b438b75da159ab51ff426bacea3d1/sdk/include/FidelityFX/gpu/spd/ffx_spd.h
3
4@group(0) @binding(0) var sampler_linear_clamp: sampler;
5@group(0) @binding(1) var<uniform> constants: Constants;
6
7#ifdef ARRAY_TEXTURE
8
9#ifdef COMBINE_BIND_GROUP
10@group(0) @binding(2) var mip_0: texture_2d_array<f32>;
11@group(0) @binding(3) var mip_1: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
12@group(0) @binding(4) var mip_2: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
13@group(0) @binding(5) var mip_3: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
14@group(0) @binding(6) var mip_4: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
15@group(0) @binding(7) var mip_5: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
16@group(0) @binding(8) var mip_6: texture_storage_2d_array<##TEXTURE_FORMAT##, read_write>;
17@group(0) @binding(9) var mip_7: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
18@group(0) @binding(10) var mip_8: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
19@group(0) @binding(11) var mip_9: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
20@group(0) @binding(12) var mip_10: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
21@group(0) @binding(13) var mip_11: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
22@group(0) @binding(14) var mip_12: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
23#endif  // COMBINE_BIND_GROUP
24
25#ifdef FIRST_PASS
26@group(0) @binding(2) var mip_0: texture_2d_array<f32>;
27@group(0) @binding(3) var mip_1: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
28@group(0) @binding(4) var mip_2: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
29@group(0) @binding(5) var mip_3: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
30@group(0) @binding(6) var mip_4: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
31@group(0) @binding(7) var mip_5: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
32@group(0) @binding(8) var mip_6: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
33#endif  // FIRST_PASS
34
35#ifdef SECOND_PASS
36@group(0) @binding(2) var mip_6: texture_2d_array<f32>;
37@group(0) @binding(3) var mip_7: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
38@group(0) @binding(4) var mip_8: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
39@group(0) @binding(5) var mip_9: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
40@group(0) @binding(6) var mip_10: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
41@group(0) @binding(7) var mip_11: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
42@group(0) @binding(8) var mip_12: texture_storage_2d_array<##TEXTURE_FORMAT##, write>;
43#endif  // SECOND_PASS
44
45#else   // ARRAY_TEXTURE
46
47#ifdef COMBINE_BIND_GROUP
48@group(0) @binding(2) var mip_0: texture_2d<f32>;
49@group(0) @binding(3) var mip_1: texture_storage_2d<##TEXTURE_FORMAT##, write>;
50@group(0) @binding(4) var mip_2: texture_storage_2d<##TEXTURE_FORMAT##, write>;
51@group(0) @binding(5) var mip_3: texture_storage_2d<##TEXTURE_FORMAT##, write>;
52@group(0) @binding(6) var mip_4: texture_storage_2d<##TEXTURE_FORMAT##, write>;
53@group(0) @binding(7) var mip_5: texture_storage_2d<##TEXTURE_FORMAT##, write>;
54@group(0) @binding(8) var mip_6: texture_storage_2d<##TEXTURE_FORMAT##, read_write>;
55@group(0) @binding(9) var mip_7: texture_storage_2d<##TEXTURE_FORMAT##, write>;
56@group(0) @binding(10) var mip_8: texture_storage_2d<##TEXTURE_FORMAT##, write>;
57@group(0) @binding(11) var mip_9: texture_storage_2d<##TEXTURE_FORMAT##, write>;
58@group(0) @binding(12) var mip_10: texture_storage_2d<##TEXTURE_FORMAT##, write>;
59@group(0) @binding(13) var mip_11: texture_storage_2d<##TEXTURE_FORMAT##, write>;
60@group(0) @binding(14) var mip_12: texture_storage_2d<##TEXTURE_FORMAT##, write>;
61#endif  // COMBINE_BIND_GROUP
62
63#ifdef FIRST_PASS
64@group(0) @binding(2) var mip_0: texture_2d<f32>;
65@group(0) @binding(3) var mip_1: texture_storage_2d<##TEXTURE_FORMAT##, write>;
66@group(0) @binding(4) var mip_2: texture_storage_2d<##TEXTURE_FORMAT##, write>;
67@group(0) @binding(5) var mip_3: texture_storage_2d<##TEXTURE_FORMAT##, write>;
68@group(0) @binding(6) var mip_4: texture_storage_2d<##TEXTURE_FORMAT##, write>;
69@group(0) @binding(7) var mip_5: texture_storage_2d<##TEXTURE_FORMAT##, write>;
70@group(0) @binding(8) var mip_6: texture_storage_2d<##TEXTURE_FORMAT##, write>;
71#endif  // FIRST_PASS
72
73#ifdef SECOND_PASS
74@group(0) @binding(2) var mip_6: texture_2d<f32>;
75@group(0) @binding(3) var mip_7: texture_storage_2d<##TEXTURE_FORMAT##, write>;
76@group(0) @binding(4) var mip_8: texture_storage_2d<##TEXTURE_FORMAT##, write>;
77@group(0) @binding(5) var mip_9: texture_storage_2d<##TEXTURE_FORMAT##, write>;
78@group(0) @binding(6) var mip_10: texture_storage_2d<##TEXTURE_FORMAT##, write>;
79@group(0) @binding(7) var mip_11: texture_storage_2d<##TEXTURE_FORMAT##, write>;
80@group(0) @binding(8) var mip_12: texture_storage_2d<##TEXTURE_FORMAT##, write>;
81#endif  // SECOND_PASS
82
83#endif  // ARRAY_TEXTURE
84
85struct Constants { mips: u32, inverse_input_size: vec2f }
86
87var<workgroup> spd_intermediate_r: array<array<f32, 16>, 16>;
88var<workgroup> spd_intermediate_g: array<array<f32, 16>, 16>;
89var<workgroup> spd_intermediate_b: array<array<f32, 16>, 16>;
90var<workgroup> spd_intermediate_a: array<array<f32, 16>, 16>;
91
92@compute
93@workgroup_size(256, 1, 1)
94fn downsample_first(
95    @builtin(workgroup_id) workgroup_id: vec3u,
96    @builtin(local_invocation_index) local_invocation_index: u32
97) {
98
99    let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);
100    let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);
101    let y = sub_xy.y + 8u * (local_invocation_index >> 7u);
102
103    spd_downsample_mips_0_1(x, y, workgroup_id.xy, local_invocation_index, constants.mips, workgroup_id.z);
104
105    spd_downsample_next_four(x, y, workgroup_id.xy, local_invocation_index, 2u, constants.mips, workgroup_id.z);
106}
107
108// TODO: Once wgpu supports globallycoherent buffers, make it actually a single pass
109@compute
110@workgroup_size(256, 1, 1)
111fn downsample_second(
112    @builtin(workgroup_id) workgroup_id: vec3u,
113    @builtin(local_invocation_index) local_invocation_index: u32,
114) {
115    let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);
116    let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);
117    let y = sub_xy.y + 8u * (local_invocation_index >> 7u);
118
119    spd_downsample_mips_6_7(x, y, constants.mips, workgroup_id.z);
120
121    spd_downsample_next_four(x, y, vec2(0u), local_invocation_index, 8u, constants.mips, workgroup_id.z);
122}
123
124fn spd_downsample_mips_0_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, mips: u32, slice: u32) {
125    var v: array<vec4f, 4>;
126
127    var tex = (workgroup_id * 64u) + vec2(x * 2u, y * 2u);
128    var pix = (workgroup_id * 32u) + vec2(x, y);
129    v[0] = spd_reduce_load_source_image(tex, slice);
130    spd_store(pix, v[0], 0u, slice);
131
132    tex = (workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u);
133    pix = (workgroup_id * 32u) + vec2(x + 16u, y);
134    v[1] = spd_reduce_load_source_image(tex, slice);
135    spd_store(pix, v[1], 0u, slice);
136
137    tex = (workgroup_id * 64u) + vec2(x * 2u, y * 2u + 32u);
138    pix = (workgroup_id * 32u) + vec2(x, y + 16u);
139    v[2] = spd_reduce_load_source_image(tex, slice);
140    spd_store(pix, v[2], 0u, slice);
141
142    tex = (workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u + 32u);
143    pix = (workgroup_id * 32u) + vec2(x + 16u, y + 16u);
144    v[3] = spd_reduce_load_source_image(tex, slice);
145    spd_store(pix, v[3], 0u, slice);
146
147    if mips <= 1u { return; }
148
149#ifdef SUBGROUP_SUPPORT
150    v[0] = spd_reduce_quad(v[0]);
151    v[1] = spd_reduce_quad(v[1]);
152    v[2] = spd_reduce_quad(v[2]);
153    v[3] = spd_reduce_quad(v[3]);
154
155    if local_invocation_index % 4u == 0u {
156        spd_store((workgroup_id * 16u) + vec2(x / 2u, y / 2u), v[0], 1u, slice);
157        spd_store_intermediate(x / 2u, y / 2u, v[0]);
158
159        spd_store((workgroup_id * 16u) + vec2(x / 2u + 8u, y / 2u), v[1], 1u, slice);
160        spd_store_intermediate(x / 2u + 8u, y / 2u, v[1]);
161
162        spd_store((workgroup_id * 16u) + vec2(x / 2u, y / 2u + 8u), v[2], 1u, slice);
163        spd_store_intermediate(x / 2u, y / 2u + 8u, v[2]);
164
165        spd_store((workgroup_id * 16u) + vec2(x / 2u + 8u, y / 2u + 8u), v[3], 1u, slice);
166        spd_store_intermediate(x / 2u + 8u, y / 2u + 8u, v[3]);
167    }
168#else   // SUBGROUP_SUPPORT
169    for (var i = 0u; i < 4u; i++) {
170        spd_store_intermediate(x, y, v[i]);
171        workgroupBarrier();
172        if local_invocation_index < 64u {
173            v[i] = spd_reduce_intermediate(
174                vec2(x * 2u + 0u, y * 2u + 0u),
175                vec2(x * 2u + 1u, y * 2u + 0u),
176                vec2(x * 2u + 0u, y * 2u + 1u),
177                vec2(x * 2u + 1u, y * 2u + 1u),
178            );
179            spd_store(vec2(workgroup_id * 16) + vec2(x + (i % 2u) * 8u, y + (i / 2u) * 8u), v[i], 1u, slice);
180        }
181        workgroupBarrier();
182    }
183
184    if local_invocation_index < 64u {
185        spd_store_intermediate(x + 0u, y + 0u, v[0]);
186        spd_store_intermediate(x + 8u, y + 0u, v[1]);
187        spd_store_intermediate(x + 0u, y + 8u, v[2]);
188        spd_store_intermediate(x + 8u, y + 8u, v[3]);
189    }
190#endif  // SUBGROUP_SUPPORT
191}
192
193fn spd_downsample_next_four(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, mips: u32, slice: u32) {
194    if mips <= base_mip { return; }
195    workgroupBarrier();
196    spd_downsample_mip_2(x, y, workgroup_id, local_invocation_index, base_mip, slice);
197
198    if mips <= base_mip + 1u { return; }
199    workgroupBarrier();
200    spd_downsample_mip_3(x, y, workgroup_id, local_invocation_index, base_mip + 1u, slice);
201
202    if mips <= base_mip + 2u { return; }
203    workgroupBarrier();
204    spd_downsample_mip_4(x, y, workgroup_id, local_invocation_index, base_mip + 2u, slice);
205
206    if mips <= base_mip + 3u { return; }
207    workgroupBarrier();
208    spd_downsample_mip_5(x, y, workgroup_id, local_invocation_index, base_mip + 3u, slice);
209}
210
211fn spd_downsample_mip_2(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {
212#ifdef SUBGROUP_SUPPORT
213    var v = spd_load_intermediate(x, y);
214    v = spd_reduce_quad(v);
215    if local_invocation_index % 4u == 0u {
216        spd_store((workgroup_id * 8u) + vec2(x / 2u, y / 2u), v, base_mip, slice);
217        spd_store_intermediate(x + (y / 2u) % 2u, y, v);
218    }
219#else   // SUBGROUP_SUPPORT
220    if local_invocation_index < 64u {
221        let v = spd_reduce_intermediate(
222            vec2(x * 2u + 0u, y * 2u + 0u),
223            vec2(x * 2u + 1u, y * 2u + 0u),
224            vec2(x * 2u + 0u, y * 2u + 1u),
225            vec2(x * 2u + 1u, y * 2u + 1u),
226        );
227        spd_store((workgroup_id * 8u) + vec2(x, y), v, base_mip, slice);
228        spd_store_intermediate(x * 2u + y % 2u, y * 2u, v);
229    }
230#endif  // SUBGROUP_SUPPORT
231}
232
233fn spd_downsample_mip_3(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {
234#ifdef SUBGROUP_SUPPORT
235    if local_invocation_index < 64u {
236        var v = spd_load_intermediate(x * 2u + y % 2u, y * 2u);
237        v = spd_reduce_quad(v);
238        if local_invocation_index % 4u == 0u {
239            spd_store((workgroup_id * 4u) + vec2(x / 2u, y / 2u), v, base_mip, slice);
240            spd_store_intermediate(x * 2u + y / 2u, y * 2u, v);
241        }
242    }
243#else   // SUBGROUP_SUPPORT
244    if local_invocation_index < 16u {
245        let v = spd_reduce_intermediate(
246            vec2(x * 4u + 0u + 0u, y * 4u + 0u),
247            vec2(x * 4u + 2u + 0u, y * 4u + 0u),
248            vec2(x * 4u + 0u + 1u, y * 4u + 2u),
249            vec2(x * 4u + 2u + 1u, y * 4u + 2u),
250        );
251        spd_store((workgroup_id * 4u) + vec2(x, y), v, base_mip, slice);
252        spd_store_intermediate(x * 4u + y, y * 4u, v);
253    }
254#endif  // SUBGROUP_SUPPORT
255}
256
257fn spd_downsample_mip_4(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {
258#ifdef SUBGROUP_SUPPORT
259    if local_invocation_index < 16u {
260        var v = spd_load_intermediate(x * 4u + y, y * 4u);
261        v = spd_reduce_quad(v);
262        if local_invocation_index % 4u == 0u {
263            spd_store((workgroup_id * 2u) + vec2(x / 2u, y / 2u), v, base_mip, slice);
264            spd_store_intermediate(x / 2u + y, 0u, v);
265        }
266    }
267#else   // SUBGROUP_SUPPORT
268    if local_invocation_index < 4u {
269        let v = spd_reduce_intermediate(
270            vec2(x * 8u + 0u + 0u + y * 2u, y * 8u + 0u),
271            vec2(x * 8u + 4u + 0u + y * 2u, y * 8u + 0u),
272            vec2(x * 8u + 0u + 1u + y * 2u, y * 8u + 4u),
273            vec2(x * 8u + 4u + 1u + y * 2u, y * 8u + 4u),
274        );
275        spd_store((workgroup_id * 2u) + vec2(x, y), v, base_mip, slice);
276        spd_store_intermediate(x + y * 2u, 0u, v);
277    }
278#endif  // SUBGROUP_SUPPORT
279}
280
281fn spd_downsample_mip_5(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {
282#ifdef SUBGROUP_SUPPORT
283    if local_invocation_index < 4u {
284        var v = spd_load_intermediate(local_invocation_index, 0u);
285        v = spd_reduce_quad(v);
286        if local_invocation_index % 4u == 0u {
287            spd_store(workgroup_id, v, base_mip, slice);
288        }
289    }
290#else   // SUBGROUP_SUPPORT
291    if local_invocation_index < 1u {
292        let v = spd_reduce_intermediate(vec2(0u, 0u), vec2(1u, 0u), vec2(2u, 0u), vec2(3u, 0u));
293        spd_store(workgroup_id, v, base_mip, slice);
294    }
295#endif  // SUBGROUP_SUPPORT
296}
297
298fn spd_downsample_mips_6_7(x: u32, y: u32, mips: u32, slice: u32) {
299    var tex = vec2(x * 4u + 0u, y * 4u + 0u);
300    var pix = vec2(x * 2u + 0u, y * 2u + 0u);
301    let v0 = spd_reduce_load_4(
302        vec2(x * 4u + 0u, y * 4u + 0u),
303        vec2(x * 4u + 1u, y * 4u + 0u),
304        vec2(x * 4u + 0u, y * 4u + 1u),
305        vec2(x * 4u + 1u, y * 4u + 1u),
306        slice
307    );
308    spd_store(pix, v0, 6u, slice);
309
310    tex = vec2(x * 4u + 2u, y * 4u + 0u);
311    pix = vec2(x * 2u + 1u, y * 2u + 0u);
312    let v1 = spd_reduce_load_4(
313        vec2(x * 4u + 2u, y * 4u + 0u),
314        vec2(x * 4u + 3u, y * 4u + 0u),
315        vec2(x * 4u + 2u, y * 4u + 1u),
316        vec2(x * 4u + 3u, y * 4u + 1u),
317        slice
318    );
319    spd_store(pix, v1, 6u, slice);
320
321    tex = vec2(x * 4u + 0u, y * 4u + 2u);
322    pix = vec2(x * 2u + 0u, y * 2u + 1u);
323    let v2 = spd_reduce_load_4(
324        vec2(x * 4u + 0u, y * 4u + 2u),
325        vec2(x * 4u + 1u, y * 4u + 2u),
326        vec2(x * 4u + 0u, y * 4u + 3u),
327        vec2(x * 4u + 1u, y * 4u + 3u),
328        slice
329    );
330    spd_store(pix, v2, 6u, slice);
331
332    tex = vec2(x * 4u + 2u, y * 4u + 2u);
333    pix = vec2(x * 2u + 1u, y * 2u + 1u);
334    let v3 = spd_reduce_load_4(
335        vec2(x * 4u + 2u, y * 4u + 2u),
336        vec2(x * 4u + 3u, y * 4u + 2u),
337        vec2(x * 4u + 2u, y * 4u + 3u),
338        vec2(x * 4u + 3u, y * 4u + 3u),
339        slice
340    );
341    spd_store(pix, v3, 6u, slice);
342
343    if mips < 7u { return; }
344
345    let v = spd_reduce_4(v0, v1, v2, v3);
346    spd_store(vec2(x, y), v, 7u, slice);
347    spd_store_intermediate(x, y, v);
348}
349
350fn remap_for_wave_reduction(a: u32) -> vec2u {
351    // This function maps linear thread IDs to 2D coordinates in a special pattern
352    // to ensure that neighboring threads process neighboring pixels
353    // For example, this transforms linear thread IDs 0,1,2,3 into a 2×2 square
354    
355    // Extract bits to form the X and Y coordinates
356    let x = insertBits(extractBits(a, 2u, 3u), a, 0u, 1u);
357    let y = insertBits(extractBits(a, 3u, 3u), extractBits(a, 1u, 2u), 0u, 2u);
358    
359    return vec2u(x, y);
360}
361
362#ifdef ARRAY_TEXTURE
363
364fn spd_reduce_load_source_image(uv: vec2u, slice: u32) -> vec4f {
365    let texture_coord = (vec2f(uv) + 0.5) * constants.inverse_input_size;
366
367    #ifdef COMBINE_BIND_GROUP
368    let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, slice, 0.0);
369    #endif  // COMBINE_BIND_GROUP
370    #ifdef FIRST_PASS
371    let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, slice, 0.0);
372    #endif  // FIRST_PASS
373    #ifdef SECOND_PASS
374    let result = textureSampleLevel(mip_6, sampler_linear_clamp, texture_coord, slice, 0.0);
375    #endif  // SECOND_PASS
376
377#ifdef SRGB_CONVERSION
378    return vec4(
379        srgb_from_linear(result.r),
380        srgb_from_linear(result.g),
381        srgb_from_linear(result.b),
382        result.a
383    );
384#else   // SRGB_CONVERSION
385    return result;
386#endif  // SRGB_CONVERSION
387
388}
389
390#else   // ARRAY_TEXTURE
391
392fn spd_reduce_load_source_image(uv: vec2u, slice: u32) -> vec4f {
393    let texture_coord = (vec2f(uv) + 0.5) * constants.inverse_input_size;
394
395    #ifdef COMBINE_BIND_GROUP
396    let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, 0.0);
397    #endif  // COMBINE_BIND_GROUP
398    #ifdef FIRST_PASS
399    let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, 0.0);
400    #endif  // FIRST_PASS
401    #ifdef SECOND_PASS
402    let result = textureSampleLevel(mip_6, sampler_linear_clamp, texture_coord, 0.0);
403    #endif  // SECOND_PASS
404
405#ifdef SRGB_CONVERSION
406    return vec4(
407        srgb_from_linear(result.r),
408        srgb_from_linear(result.g),
409        srgb_from_linear(result.b),
410        result.a
411    );
412#else   // SRGB_CONVERSION
413    return result;
414#endif  // SRGB_CONVERSION
415
416}
417
418#endif  // ARRAY_TEXTURE
419
420#ifdef ARRAY_TEXTURE
421fn spd_store(pix: vec2u, value: vec4f, mip: u32, slice: u32) {
422    if mip >= constants.mips { return; }
423    switch mip {
424        #ifdef COMBINE_BIND_GROUP
425        case 0u: { textureStore(mip_1, pix, slice, value); }
426        case 1u: { textureStore(mip_2, pix, slice, value); }
427        case 2u: { textureStore(mip_3, pix, slice, value); }
428        case 3u: { textureStore(mip_4, pix, slice, value); }
429        case 4u: { textureStore(mip_5, pix, slice, value); }
430        case 5u: { textureStore(mip_6, pix, slice, value); }
431        case 6u: { textureStore(mip_7, pix, slice, value); }
432        case 7u: { textureStore(mip_8, pix, slice, value); }
433        case 8u: { textureStore(mip_9, pix, slice, value); }
434        case 9u: { textureStore(mip_10, pix, slice, value); }
435        case 10u: { textureStore(mip_11, pix, slice, value); }
436        case 11u: { textureStore(mip_12, pix, slice, value); }
437        #endif  // COMBINE_BIND_GROUP
438        #ifdef FIRST_PASS
439        case 0u: { textureStore(mip_1, pix, slice, value); }
440        case 1u: { textureStore(mip_2, pix, slice, value); }
441        case 2u: { textureStore(mip_3, pix, slice, value); }
442        case 3u: { textureStore(mip_4, pix, slice, value); }
443        case 4u: { textureStore(mip_5, pix, slice, value); }
444        case 5u: { textureStore(mip_6, pix, slice, value); }
445        #endif  // FIRST_PASS
446        #ifdef SECOND_PASS
447        case 6u: { textureStore(mip_7, pix, slice, value); }
448        case 7u: { textureStore(mip_8, pix, slice, value); }
449        case 8u: { textureStore(mip_9, pix, slice, value); }
450        case 9u: { textureStore(mip_10, pix, slice, value); }
451        case 10u: { textureStore(mip_11, pix, slice, value); }
452        case 11u: { textureStore(mip_12, pix, slice, value); }
453        #endif  // SECOND_PASS
454        default: {}
455    }
456}
457#else   // ARRAY_TEXTURE
458fn spd_store(pix: vec2u, value: vec4f, mip: u32, slice: u32) {
459    if mip >= constants.mips { return; }
460    switch mip {
461        #ifdef COMBINE_BIND_GROUP
462        case 0u: { textureStore(mip_1, pix, value); }
463        case 1u: { textureStore(mip_2, pix, value); }
464        case 2u: { textureStore(mip_3, pix, value); }
465        case 3u: { textureStore(mip_4, pix, value); }
466        case 4u: { textureStore(mip_5, pix, value); }
467        case 5u: { textureStore(mip_6, pix, value); }
468        case 6u: { textureStore(mip_7, pix, value); }
469        case 7u: { textureStore(mip_8, pix, value); }
470        case 8u: { textureStore(mip_9, pix, value); }
471        case 9u: { textureStore(mip_10, pix, value); }
472        case 10u: { textureStore(mip_11, pix, value); }
473        case 11u: { textureStore(mip_12, pix, value); }
474        #endif  // COMBINE_BIND_GROUP
475        #ifdef FIRST_PASS
476        case 0u: { textureStore(mip_1, pix, value); }
477        case 1u: { textureStore(mip_2, pix, value); }
478        case 2u: { textureStore(mip_3, pix, value); }
479        case 3u: { textureStore(mip_4, pix, value); }
480        case 4u: { textureStore(mip_5, pix, value); }
481        case 5u: { textureStore(mip_6, pix, value); }
482        #endif  // FIRST_PASS
483        #ifdef SECOND_PASS
484        case 6u: { textureStore(mip_7, pix, value); }
485        case 7u: { textureStore(mip_8, pix, value); }
486        case 8u: { textureStore(mip_9, pix, value); }
487        case 9u: { textureStore(mip_10, pix, value); }
488        case 10u: { textureStore(mip_11, pix, value); }
489        case 11u: { textureStore(mip_12, pix, value); }
490        #endif  // SECOND_PASS
491        default: {}
492    }
493}
494#endif  // ARRAY_TEXTURE
495
496fn spd_store_intermediate(x: u32, y: u32, value: vec4f) {
497    spd_intermediate_r[x][y] = value.x;
498    spd_intermediate_g[x][y] = value.y;
499    spd_intermediate_b[x][y] = value.z;
500    spd_intermediate_a[x][y] = value.w;
501}
502
503fn spd_load_intermediate(x: u32, y: u32) -> vec4f {
504    return vec4(spd_intermediate_r[x][y], spd_intermediate_g[x][y], spd_intermediate_b[x][y], spd_intermediate_a[x][y]);
505}
506
507fn spd_reduce_intermediate(i0: vec2u, i1: vec2u, i2: vec2u, i3: vec2u) -> vec4f {
508    let v0 = spd_load_intermediate(i0.x, i0.y);
509    let v1 = spd_load_intermediate(i1.x, i1.y);
510    let v2 = spd_load_intermediate(i2.x, i2.y);
511    let v3 = spd_load_intermediate(i3.x, i3.y);
512    return spd_reduce_4(v0, v1, v2, v3);
513}
514
515#ifdef ARRAY_TEXTURE
516fn spd_reduce_load_4(i0: vec2u, i1: vec2u, i2: vec2u, i3: vec2u, slice: u32) -> vec4f {
517    #ifdef COMBINE_BIND_GROUP
518    let v0 = textureLoad(mip_6, i0, slice);
519    let v1 = textureLoad(mip_6, i1, slice);
520    let v2 = textureLoad(mip_6, i2, slice);
521    let v3 = textureLoad(mip_6, i3, slice);
522    return spd_reduce_4(v0, v1, v2, v3);
523    #endif  // COMBINE_BIND_GROUP
524    #ifdef FIRST_PASS
525    return vec4(0.0, 0.0, 0.0, 0.0);
526    #endif  // FIRST_PASS
527    #ifdef SECOND_PASS
528    let v0 = textureLoad(mip_6, i0, slice, 0);
529    let v1 = textureLoad(mip_6, i1, slice, 0);
530    let v2 = textureLoad(mip_6, i2, slice, 0);
531    let v3 = textureLoad(mip_6, i3, slice, 0);
532    return spd_reduce_4(v0, v1, v2, v3);
533    #endif  // SECOND_PASS
534}
535#else   // ARRAY_TEXTURE
536fn spd_reduce_load_4(i0: vec2u, i1: vec2u, i2: vec2u, i3: vec2u, slice: u32) -> vec4f {
537    #ifdef COMBINE_BIND_GROUP
538    let v0 = textureLoad(mip_6, i0);
539    let v1 = textureLoad(mip_6, i1);
540    let v2 = textureLoad(mip_6, i2);
541    let v3 = textureLoad(mip_6, i3);
542    return spd_reduce_4(v0, v1, v2, v3);
543    #endif  // COMBINE_BIND_GROUP
544    #ifdef FIRST_PASS
545    return vec4(0.0, 0.0, 0.0, 0.0);
546    #endif  // FIRST_PASS
547    #ifdef SECOND_PASS
548    let v0 = textureLoad(mip_6, i0, 0);
549    let v1 = textureLoad(mip_6, i1, 0);
550    let v2 = textureLoad(mip_6, i2, 0);
551    let v3 = textureLoad(mip_6, i3, 0);
552    return spd_reduce_4(v0, v1, v2, v3);
553    #endif  // SECOND_PASS
554}
555#endif  // ARRAY_TEXTURE
556
557fn spd_reduce_4(v0: vec4f, v1: vec4f, v2: vec4f, v3: vec4f) -> vec4f {
558    return (v0 + v1 + v2 + v3) * 0.25;
559}
560
561#ifdef SUBGROUP_SUPPORT
562fn spd_reduce_quad(v: vec4f) -> vec4f {
563    let v0 = v;
564    let v1 = quadSwapX(v);
565    let v2 = quadSwapY(v);
566    let v3 = quadSwapDiagonal(v);
567    return spd_reduce_4(v0, v1, v2, v3);
568}
569#endif  // SUBGROUP_SUPPORT
570
571fn srgb_from_linear(value: f32) -> f32 {
572    let j = vec3(0.0031308 * 12.92, 12.92, 1.0 / 2.4);
573    let k = vec2(1.055, -0.055);
574    return clamp(j.x, value * j.y, pow(value, j.z) * k.x + k.y);
575}