1//! Single pass downsampling shader for creating the mip chain for a 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;7@if(TEXTURE_FORMAT_RGBA8UNORM) {8alias MipWrite2d = texture_storage_2d<rgba8unorm, write>;9alias MipReadWrite2d = texture_storage_2d<rgba8unorm, read_write>;10alias MipWrite2dArray = texture_storage_2d_array<rgba8unorm, write>;11alias MipReadWrite2dArray = texture_storage_2d_array<rgba8unorm, read_write>;12}14@if(TEXTURE_FORMAT_RGBA8SNORM) {15alias MipWrite2d = texture_storage_2d<rgba8snorm, write>;16alias MipReadWrite2d = texture_storage_2d<rgba8snorm, read_write>;17alias MipWrite2dArray = texture_storage_2d_array<rgba8snorm, write>;18alias MipReadWrite2dArray = texture_storage_2d_array<rgba8snorm, read_write>;19}21@if(TEXTURE_FORMAT_RGBA16UNORM) {22alias MipWrite2d = texture_storage_2d<rgba16unorm, write>;23alias MipReadWrite2d = texture_storage_2d<rgba16unorm, read_write>;24alias MipWrite2dArray = texture_storage_2d_array<rgba16unorm, write>;25alias MipReadWrite2dArray = texture_storage_2d_array<rgba16unorm, read_write>;26}28@if(TEXTURE_FORMAT_RGBA16SNORM) {29alias MipWrite2d = texture_storage_2d<rgba16snorm, write>;30alias MipReadWrite2d = texture_storage_2d<rgba16snorm, read_write>;31alias MipWrite2dArray = texture_storage_2d_array<rgba16snorm, write>;32alias MipReadWrite2dArray = texture_storage_2d_array<rgba16snorm, read_write>;33}35@if(TEXTURE_FORMAT_RGBA16FLOAT) {36alias MipWrite2d = texture_storage_2d<rgba16float, write>;37alias MipReadWrite2d = texture_storage_2d<rgba16float, read_write>;38alias MipWrite2dArray = texture_storage_2d_array<rgba16float, write>;39alias MipReadWrite2dArray = texture_storage_2d_array<rgba16float, read_write>;40}42@if(TEXTURE_FORMAT_RG8UNORM) {43alias MipWrite2d = texture_storage_2d<rg8unorm, write>;44alias MipReadWrite2d = texture_storage_2d<rg8unorm, read_write>;45alias MipWrite2dArray = texture_storage_2d_array<rg8unorm, write>;46alias MipReadWrite2dArray = texture_storage_2d_array<rg8unorm, read_write>;47}49@if(TEXTURE_FORMAT_RG8SNORM) {50alias MipWrite2d = texture_storage_2d<rg8snorm, write>;51alias MipReadWrite2d = texture_storage_2d<rg8snorm, read_write>;52alias MipWrite2dArray = texture_storage_2d_array<rg8snorm, write>;53alias MipReadWrite2dArray = texture_storage_2d_array<rg8snorm, read_write>;54}56@if(TEXTURE_FORMAT_RG16UNORM) {57alias MipWrite2d = texture_storage_2d<rg16unorm, write>;58alias MipReadWrite2d = texture_storage_2d<rg16unorm, read_write>;59alias MipWrite2dArray = texture_storage_2d_array<rg16unorm, write>;60alias MipReadWrite2dArray = texture_storage_2d_array<rg16unorm, read_write>;61}63@if(TEXTURE_FORMAT_RG16SNORM) {64alias MipWrite2d = texture_storage_2d<rg16snorm, write>;65alias MipReadWrite2d = texture_storage_2d<rg16snorm, read_write>;66alias MipWrite2dArray = texture_storage_2d_array<rg16snorm, write>;67alias MipReadWrite2dArray = texture_storage_2d_array<rg16snorm, read_write>;68}70@if(TEXTURE_FORMAT_RG16FLOAT) {71alias MipWrite2d = texture_storage_2d<rg16float, write>;72alias MipReadWrite2d = texture_storage_2d<rg16float, read_write>;73alias MipWrite2dArray = texture_storage_2d_array<rg16float, write>;74alias MipReadWrite2dArray = texture_storage_2d_array<rg16float, read_write>;75}77@if(TEXTURE_FORMAT_R32FLOAT) {78alias MipWrite2d = texture_storage_2d<r32float, write>;79alias MipReadWrite2d = texture_storage_2d<r32float, read_write>;80alias MipWrite2dArray = texture_storage_2d_array<r32float, write>;81alias MipReadWrite2dArray = texture_storage_2d_array<r32float, read_write>;82}84@if(TEXTURE_FORMAT_RG32FLOAT) {85alias MipWrite2d = texture_storage_2d<rg32float, write>;86alias MipReadWrite2d = texture_storage_2d<rg32float, read_write>;87alias MipWrite2dArray = texture_storage_2d_array<rg32float, write>;88alias MipReadWrite2dArray = texture_storage_2d_array<rg32float, read_write>;89}91@if(TEXTURE_FORMAT_RGBA32FLOAT) {92alias MipWrite2d = texture_storage_2d<rgba32float, write>;93alias MipReadWrite2d = texture_storage_2d<rgba32float, read_write>;94alias MipWrite2dArray = texture_storage_2d_array<rgba32float, write>;95alias MipReadWrite2dArray = texture_storage_2d_array<rgba32float, read_write>;96}98@if(TEXTURE_FORMAT_BGRA8UNORM) {99alias MipWrite2d = texture_storage_2d<bgra8unorm, write>;100alias MipReadWrite2d = texture_storage_2d<bgra8unorm, read_write>;101alias MipWrite2dArray = texture_storage_2d_array<bgra8unorm, write>;102alias MipReadWrite2dArray = texture_storage_2d_array<bgra8unorm, read_write>;103}105@if(TEXTURE_FORMAT_R8UNORM) {106alias MipWrite2d = texture_storage_2d<r8unorm, write>;107alias MipReadWrite2d = texture_storage_2d<r8unorm, read_write>;108alias MipWrite2dArray = texture_storage_2d_array<r8unorm, write>;109alias MipReadWrite2dArray = texture_storage_2d_array<r8unorm, read_write>;110}112@if(TEXTURE_FORMAT_R8SNORM) {113alias MipWrite2d = texture_storage_2d<r8snorm, write>;114alias MipReadWrite2d = texture_storage_2d<r8snorm, read_write>;115alias MipWrite2dArray = texture_storage_2d_array<r8snorm, write>;116alias MipReadWrite2dArray = texture_storage_2d_array<r8snorm, read_write>;117}119@if(TEXTURE_FORMAT_R16UNORM) {120alias MipWrite2d = texture_storage_2d<r16unorm, write>;121alias MipReadWrite2d = texture_storage_2d<r16unorm, read_write>;122alias MipWrite2dArray = texture_storage_2d_array<r16unorm, write>;123alias MipReadWrite2dArray = texture_storage_2d_array<r16unorm, read_write>;124}126@if(TEXTURE_FORMAT_R16SNORM) {127alias MipWrite2d = texture_storage_2d<r16snorm, write>;128alias MipReadWrite2d = texture_storage_2d<r16snorm, read_write>;129alias MipWrite2dArray = texture_storage_2d_array<r16snorm, write>;130alias MipReadWrite2dArray = texture_storage_2d_array<r16snorm, read_write>;131}133@if(TEXTURE_FORMAT_R16FLOAT) {134alias MipWrite2d = texture_storage_2d<r16float, write>;135alias MipReadWrite2d = texture_storage_2d<r16float, read_write>;136alias MipWrite2dArray = texture_storage_2d_array<r16float, write>;137alias MipReadWrite2dArray = texture_storage_2d_array<r16float, read_write>;138}140@if(TEXTURE_FORMAT_RGB10A2UNORM) {141alias MipWrite2d = texture_storage_2d<rgb10a2unorm, write>;142alias MipReadWrite2d = texture_storage_2d<rgb10a2unorm, read_write>;143alias MipWrite2dArray = texture_storage_2d_array<rgb10a2unorm, write>;144alias MipReadWrite2dArray = texture_storage_2d_array<rgb10a2unorm, read_write>;145}147@if(TEXTURE_FORMAT_RG11B10UFLOAT) {148alias MipWrite2d = texture_storage_2d<rg11b10ufloat, write>;149alias MipReadWrite2d = texture_storage_2d<rg11b10ufloat, read_write>;150alias MipWrite2dArray = texture_storage_2d_array<rg11b10ufloat, write>;151alias MipReadWrite2dArray = texture_storage_2d_array<rg11b10ufloat, read_write>;152}154@if(ARRAY_TEXTURE) {155alias SourceTexture = texture_2d_array<f32>;156alias MipWrite = MipWrite2dArray;157alias MipReadWrite = MipReadWrite2dArray;158} @else {159alias SourceTexture = texture_2d<f32>;160alias MipWrite = MipWrite2d;161alias MipReadWrite = MipReadWrite2d;162}164@if(COMBINE_BIND_GROUP || SPLIT_BIND_GROUP_FIRST) {165@group(0) @binding(2) var mip_0: SourceTexture;166@group(0) @binding(3) var mip_1: MipWrite;167@group(0) @binding(4) var mip_2: MipWrite;168@group(0) @binding(5) var mip_3: MipWrite;169@group(0) @binding(6) var mip_4: MipWrite;170@group(0) @binding(7) var mip_5: MipWrite;171}173@if(COMBINE_BIND_GROUP) {174@group(0) @binding(8) var mip_6: MipReadWrite;175@group(0) @binding(9) var mip_7: MipWrite;176@group(0) @binding(10) var mip_8: MipWrite;177@group(0) @binding(11) var mip_9: MipWrite;178@group(0) @binding(12) var mip_10: MipWrite;179@group(0) @binding(13) var mip_11: MipWrite;180@group(0) @binding(14) var mip_12: MipWrite;181}183@if(SPLIT_BIND_GROUP_FIRST) {184@group(0) @binding(8) var mip_6: MipWrite;185}187@if(SPLIT_BIND_GROUP_SECOND) {188@group(0) @binding(2) var mip_6: SourceTexture;189@group(0) @binding(3) var mip_7: MipWrite;190@group(0) @binding(4) var mip_8: MipWrite;191@group(0) @binding(5) var mip_9: MipWrite;192@group(0) @binding(6) var mip_10: MipWrite;193@group(0) @binding(7) var mip_11: MipWrite;194@group(0) @binding(8) var mip_12: MipWrite;195}197struct Constants { mips: u32, inverse_input_size: vec2f }199var<workgroup> spd_intermediate_r: array<array<f32, 16>, 16>;200var<workgroup> spd_intermediate_g: array<array<f32, 16>, 16>;201var<workgroup> spd_intermediate_b: array<array<f32, 16>, 16>;202var<workgroup> spd_intermediate_a: array<array<f32, 16>, 16>;204@compute205@workgroup_size(256, 1, 1)206fn downsample_first(207 @builtin(workgroup_id) workgroup_id: vec3u,208 @builtin(local_invocation_index) local_invocation_index: u32209) {211 let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);212 let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);213 let y = sub_xy.y + 8u * (local_invocation_index >> 7u);215 spd_downsample_mips_0_1(x, y, workgroup_id.xy, local_invocation_index, constants.mips, workgroup_id.z);217 spd_downsample_next_four(x, y, workgroup_id.xy, local_invocation_index, 2u, constants.mips, workgroup_id.z);218}220// TODO: Once wgpu supports globallycoherent buffers, make it actually a single pass221@compute222@workgroup_size(256, 1, 1)223fn downsample_second(224 @builtin(workgroup_id) workgroup_id: vec3u,225 @builtin(local_invocation_index) local_invocation_index: u32,226) {227 let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);228 let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);229 let y = sub_xy.y + 8u * (local_invocation_index >> 7u);231 spd_downsample_mips_6_7(x, y, constants.mips, workgroup_id.z);233 spd_downsample_next_four(x, y, vec2(0u), local_invocation_index, 8u, constants.mips, workgroup_id.z);234}236fn spd_downsample_mips_0_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, mips: u32, slice: u32) {237 var v: array<vec4f, 4>;239 var tex = (workgroup_id * 64u) + vec2(x * 2u, y * 2u);240 var pix = (workgroup_id * 32u) + vec2(x, y);241 v[0] = spd_reduce_load_source_image(tex, slice);242 spd_store(pix, v[0], 0u, slice);244 tex = (workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u);245 pix = (workgroup_id * 32u) + vec2(x + 16u, y);246 v[1] = spd_reduce_load_source_image(tex, slice);247 spd_store(pix, v[1], 0u, slice);249 tex = (workgroup_id * 64u) + vec2(x * 2u, y * 2u + 32u);250 pix = (workgroup_id * 32u) + vec2(x, y + 16u);251 v[2] = spd_reduce_load_source_image(tex, slice);252 spd_store(pix, v[2], 0u, slice);254 tex = (workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u + 32u);255 pix = (workgroup_id * 32u) + vec2(x + 16u, y + 16u);256 v[3] = spd_reduce_load_source_image(tex, slice);257 spd_store(pix, v[3], 0u, slice);259 if mips <= 1u { return; }261@if(SUBGROUP_SUPPORT) {262 v[0] = spd_reduce_quad(v[0]);263 v[1] = spd_reduce_quad(v[1]);264 v[2] = spd_reduce_quad(v[2]);265 v[3] = spd_reduce_quad(v[3]);267 if local_invocation_index % 4u == 0u {268 spd_store((workgroup_id * 16u) + vec2(x / 2u, y / 2u), v[0], 1u, slice);269 spd_store_intermediate(x / 2u, y / 2u, v[0]);271 spd_store((workgroup_id * 16u) + vec2(x / 2u + 8u, y / 2u), v[1], 1u, slice);272 spd_store_intermediate(x / 2u + 8u, y / 2u, v[1]);274 spd_store((workgroup_id * 16u) + vec2(x / 2u, y / 2u + 8u), v[2], 1u, slice);275 spd_store_intermediate(x / 2u, y / 2u + 8u, v[2]);277 spd_store((workgroup_id * 16u) + vec2(x / 2u + 8u, y / 2u + 8u), v[3], 1u, slice);278 spd_store_intermediate(x / 2u + 8u, y / 2u + 8u, v[3]);279 }280}281@else {282 for (var i = 0u; i < 4u; i++) {283 spd_store_intermediate(x, y, v[i]);284 workgroupBarrier();285 if local_invocation_index < 64u {286 v[i] = spd_reduce_intermediate(287 vec2(x * 2u + 0u, y * 2u + 0u),288 vec2(x * 2u + 1u, y * 2u + 0u),289 vec2(x * 2u + 0u, y * 2u + 1u),290 vec2(x * 2u + 1u, y * 2u + 1u),291 );292 spd_store(vec2(workgroup_id * 16) + vec2(x + (i % 2u) * 8u, y + (i / 2u) * 8u), v[i], 1u, slice);293 }294 workgroupBarrier();295 }297 if local_invocation_index < 64u {298 spd_store_intermediate(x + 0u, y + 0u, v[0]);299 spd_store_intermediate(x + 8u, y + 0u, v[1]);300 spd_store_intermediate(x + 0u, y + 8u, v[2]);301 spd_store_intermediate(x + 8u, y + 8u, v[3]);302 }303}304}306fn spd_downsample_next_four(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, mips: u32, slice: u32) {307 if mips <= base_mip { return; }308 workgroupBarrier();309 spd_downsample_mip_2(x, y, workgroup_id, local_invocation_index, base_mip, slice);311 if mips <= base_mip + 1u { return; }312 workgroupBarrier();313 spd_downsample_mip_3(x, y, workgroup_id, local_invocation_index, base_mip + 1u, slice);315 if mips <= base_mip + 2u { return; }316 workgroupBarrier();317 spd_downsample_mip_4(x, y, workgroup_id, local_invocation_index, base_mip + 2u, slice);319 if mips <= base_mip + 3u { return; }320 workgroupBarrier();321 spd_downsample_mip_5(x, y, workgroup_id, local_invocation_index, base_mip + 3u, slice);322}324fn spd_downsample_mip_2(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {325@if(SUBGROUP_SUPPORT) {326 var v = spd_load_intermediate(x, y);327 v = spd_reduce_quad(v);328 if local_invocation_index % 4u == 0u {329 spd_store((workgroup_id * 8u) + vec2(x / 2u, y / 2u), v, base_mip, slice);330 spd_store_intermediate(x + (y / 2u) % 2u, y, v);331 }332}333@else {334 if local_invocation_index < 64u {335 let v = spd_reduce_intermediate(336 vec2(x * 2u + 0u, y * 2u + 0u),337 vec2(x * 2u + 1u, y * 2u + 0u),338 vec2(x * 2u + 0u, y * 2u + 1u),339 vec2(x * 2u + 1u, y * 2u + 1u),340 );341 spd_store((workgroup_id * 8u) + vec2(x, y), v, base_mip, slice);342 spd_store_intermediate(x * 2u + y % 2u, y * 2u, v);343 }344}345}347fn spd_downsample_mip_3(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {348@if(SUBGROUP_SUPPORT)349 if local_invocation_index < 64u {350 var v = spd_load_intermediate(x * 2u + y % 2u, y * 2u);351 v = spd_reduce_quad(v);352 if local_invocation_index % 4u == 0u {353 spd_store((workgroup_id * 4u) + vec2(x / 2u, y / 2u), v, base_mip, slice);354 spd_store_intermediate(x * 2u + y / 2u, y * 2u, v);355 }356 }357@else // SUBGROUP_SUPPORT358 if local_invocation_index < 16u {359 let v = spd_reduce_intermediate(360 vec2(x * 4u + 0u + 0u, y * 4u + 0u),361 vec2(x * 4u + 2u + 0u, y * 4u + 0u),362 vec2(x * 4u + 0u + 1u, y * 4u + 2u),363 vec2(x * 4u + 2u + 1u, y * 4u + 2u),364 );365 spd_store((workgroup_id * 4u) + vec2(x, y), v, base_mip, slice);366 spd_store_intermediate(x * 4u + y, y * 4u, v);367 }368}370fn spd_downsample_mip_4(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {371@if(SUBGROUP_SUPPORT)372 if local_invocation_index < 16u {373 var v = spd_load_intermediate(x * 4u + y, y * 4u);374 v = spd_reduce_quad(v);375 if local_invocation_index % 4u == 0u {376 spd_store((workgroup_id * 2u) + vec2(x / 2u, y / 2u), v, base_mip, slice);377 spd_store_intermediate(x / 2u + y, 0u, v);378 }379 }380@else // SUBGROUP_SUPPORT381 if local_invocation_index < 4u {382 let v = spd_reduce_intermediate(383 vec2(x * 8u + 0u + 0u + y * 2u, y * 8u + 0u),384 vec2(x * 8u + 4u + 0u + y * 2u, y * 8u + 0u),385 vec2(x * 8u + 0u + 1u + y * 2u, y * 8u + 4u),386 vec2(x * 8u + 4u + 1u + y * 2u, y * 8u + 4u),387 );388 spd_store((workgroup_id * 2u) + vec2(x, y), v, base_mip, slice);389 spd_store_intermediate(x + y * 2u, 0u, v);390 }391}393fn spd_downsample_mip_5(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32) {394@if(SUBGROUP_SUPPORT)395 if local_invocation_index < 4u {396 var v = spd_load_intermediate(local_invocation_index, 0u);397 v = spd_reduce_quad(v);398 if local_invocation_index % 4u == 0u {399 spd_store(workgroup_id, v, base_mip, slice);400 }401 }402@else // SUBGROUP_SUPPORT403 if local_invocation_index < 1u {404 let v = spd_reduce_intermediate(vec2(0u, 0u), vec2(1u, 0u), vec2(2u, 0u), vec2(3u, 0u));405 spd_store(workgroup_id, v, base_mip, slice);406 }407}409fn spd_downsample_mips_6_7(x: u32, y: u32, mips: u32, slice: u32) {410 var tex = vec2(x * 4u + 0u, y * 4u + 0u);411 var pix = vec2(x * 2u + 0u, y * 2u + 0u);412 let v0 = spd_reduce_load_4(413 vec2(x * 4u + 0u, y * 4u + 0u),414 vec2(x * 4u + 1u, y * 4u + 0u),415 vec2(x * 4u + 0u, y * 4u + 1u),416 vec2(x * 4u + 1u, y * 4u + 1u),417 slice418 );419 spd_store(pix, v0, 6u, slice);421 tex = vec2(x * 4u + 2u, y * 4u + 0u);422 pix = vec2(x * 2u + 1u, y * 2u + 0u);423 let v1 = spd_reduce_load_4(424 vec2(x * 4u + 2u, y * 4u + 0u),425 vec2(x * 4u + 3u, y * 4u + 0u),426 vec2(x * 4u + 2u, y * 4u + 1u),427 vec2(x * 4u + 3u, y * 4u + 1u),428 slice429 );430 spd_store(pix, v1, 6u, slice);432 tex = vec2(x * 4u + 0u, y * 4u + 2u);433 pix = vec2(x * 2u + 0u, y * 2u + 1u);434 let v2 = spd_reduce_load_4(435 vec2(x * 4u + 0u, y * 4u + 2u),436 vec2(x * 4u + 1u, y * 4u + 2u),437 vec2(x * 4u + 0u, y * 4u + 3u),438 vec2(x * 4u + 1u, y * 4u + 3u),439 slice440 );441 spd_store(pix, v2, 6u, slice);443 tex = vec2(x * 4u + 2u, y * 4u + 2u);444 pix = vec2(x * 2u + 1u, y * 2u + 1u);445 let v3 = spd_reduce_load_4(446 vec2(x * 4u + 2u, y * 4u + 2u),447 vec2(x * 4u + 3u, y * 4u + 2u),448 vec2(x * 4u + 2u, y * 4u + 3u),449 vec2(x * 4u + 3u, y * 4u + 3u),450 slice451 );452 spd_store(pix, v3, 6u, slice);454 if mips < 7u { return; }456 let v = spd_reduce_4(v0, v1, v2, v3);457 spd_store(vec2(x, y), v, 7u, slice);458 spd_store_intermediate(x, y, v);459}461fn remap_for_wave_reduction(a: u32) -> vec2u {462 // This function maps linear thread IDs to 2D coordinates in a special pattern463 // to ensure that neighboring threads process neighboring pixels464 // For example, this transforms linear thread IDs 0,1,2,3 into a 2×2 square466 // Extract bits to form the X and Y coordinates467 let x = insertBits(extractBits(a, 2u, 3u), a, 0u, 1u);468 let y = insertBits(extractBits(a, 3u, 3u), extractBits(a, 1u, 2u), 0u, 2u);470 return vec2u(x, y);471}473@if(ARRAY_TEXTURE)474fn spd_reduce_load_source_image(uv: vec2u, slice: u32) -> vec4f {475 let texture_coord = (vec2f(uv) + 0.5) * constants.inverse_input_size;477 @if(COMBINE_BIND_GROUP)478 let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, slice, 0.0);479 @if(SPLIT_BIND_GROUP_FIRST)480 let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, slice, 0.0);481 @if(SPLIT_BIND_GROUP_SECOND)482 let result = textureSampleLevel(mip_6, sampler_linear_clamp, texture_coord, slice, 0.0);484@if(SRGB_CONVERSION)485 return vec4(486 srgb_from_linear(result.r),487 srgb_from_linear(result.g),488 srgb_from_linear(result.b),489 result.a490 );491@else // SRGB_CONVERSION492 return result;494}496@else // ARRAY_TEXTURE497fn spd_reduce_load_source_image(uv: vec2u, slice: u32) -> vec4f {498 let texture_coord = (vec2f(uv) + 0.5) * constants.inverse_input_size;500 @if(COMBINE_BIND_GROUP)501 let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, 0.0);502 @if(SPLIT_BIND_GROUP_FIRST)503 let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, 0.0);504 @if(SPLIT_BIND_GROUP_SECOND)505 let result = textureSampleLevel(mip_6, sampler_linear_clamp, texture_coord, 0.0);507@if(SRGB_CONVERSION)508 return vec4(509 srgb_from_linear(result.r),510 srgb_from_linear(result.g),511 srgb_from_linear(result.b),512 result.a513 );514@else // SRGB_CONVERSION515 return result;517}519@if(ARRAY_TEXTURE)520fn spd_store(pix: vec2u, value: vec4f, mip: u32, slice: u32) {521 if mip >= constants.mips { return; }522 switch mip {523 @if(COMBINE_BIND_GROUP)524 case 0u: { textureStore(mip_1, pix, slice, value); }525 @if(COMBINE_BIND_GROUP)526 case 1u: { textureStore(mip_2, pix, slice, value); }527 @if(COMBINE_BIND_GROUP)528 case 2u: { textureStore(mip_3, pix, slice, value); }529 @if(COMBINE_BIND_GROUP)530 case 3u: { textureStore(mip_4, pix, slice, value); }531 @if(COMBINE_BIND_GROUP)532 case 4u: { textureStore(mip_5, pix, slice, value); }533 @if(COMBINE_BIND_GROUP)534 case 5u: { textureStore(mip_6, pix, slice, value); }535 @if(COMBINE_BIND_GROUP)536 case 6u: { textureStore(mip_7, pix, slice, value); }537 @if(COMBINE_BIND_GROUP)538 case 7u: { textureStore(mip_8, pix, slice, value); }539 @if(COMBINE_BIND_GROUP)540 case 8u: { textureStore(mip_9, pix, slice, value); }541 @if(COMBINE_BIND_GROUP)542 case 9u: { textureStore(mip_10, pix, slice, value); }543 @if(COMBINE_BIND_GROUP)544 case 10u: { textureStore(mip_11, pix, slice, value); }545 @if(COMBINE_BIND_GROUP)546 case 11u: { textureStore(mip_12, pix, slice, value); }547 @if(SPLIT_BIND_GROUP_FIRST)548 case 0u: { textureStore(mip_1, pix, slice, value); }549 @if(SPLIT_BIND_GROUP_FIRST)550 case 1u: { textureStore(mip_2, pix, slice, value); }551 @if(SPLIT_BIND_GROUP_FIRST)552 case 2u: { textureStore(mip_3, pix, slice, value); }553 @if(SPLIT_BIND_GROUP_FIRST)554 case 3u: { textureStore(mip_4, pix, slice, value); }555 @if(SPLIT_BIND_GROUP_FIRST)556 case 4u: { textureStore(mip_5, pix, slice, value); }557 @if(SPLIT_BIND_GROUP_FIRST)558 case 5u: { textureStore(mip_6, pix, slice, value); }559 @if(SPLIT_BIND_GROUP_SECOND)560 case 6u: { textureStore(mip_7, pix, slice, value); }561 @if(SPLIT_BIND_GROUP_SECOND)562 case 7u: { textureStore(mip_8, pix, slice, value); }563 @if(SPLIT_BIND_GROUP_SECOND)564 case 8u: { textureStore(mip_9, pix, slice, value); }565 @if(SPLIT_BIND_GROUP_SECOND)566 case 9u: { textureStore(mip_10, pix, slice, value); }567 @if(SPLIT_BIND_GROUP_SECOND)568 case 10u: { textureStore(mip_11, pix, slice, value); }569 @if(SPLIT_BIND_GROUP_SECOND)570 case 11u: { textureStore(mip_12, pix, slice, value); }571 default: {}572 }573}574@else // ARRAY_TEXTURE575fn spd_store(pix: vec2u, value: vec4f, mip: u32, slice: u32) {576 if mip >= constants.mips { return; }577 switch mip {578 @if(COMBINE_BIND_GROUP)579 case 0u: { textureStore(mip_1, pix, value); }580 @if(COMBINE_BIND_GROUP)581 case 1u: { textureStore(mip_2, pix, value); }582 @if(COMBINE_BIND_GROUP)583 case 2u: { textureStore(mip_3, pix, value); }584 @if(COMBINE_BIND_GROUP)585 case 3u: { textureStore(mip_4, pix, value); }586 @if(COMBINE_BIND_GROUP)587 case 4u: { textureStore(mip_5, pix, value); }588 @if(COMBINE_BIND_GROUP)589 case 5u: { textureStore(mip_6, pix, value); }590 @if(COMBINE_BIND_GROUP)591 case 6u: { textureStore(mip_7, pix, value); }592 @if(COMBINE_BIND_GROUP)593 case 7u: { textureStore(mip_8, pix, value); }594 @if(COMBINE_BIND_GROUP)595 case 8u: { textureStore(mip_9, pix, value); }596 @if(COMBINE_BIND_GROUP)597 case 9u: { textureStore(mip_10, pix, value); }598 @if(COMBINE_BIND_GROUP)599 case 10u: { textureStore(mip_11, pix, value); }600 @if(COMBINE_BIND_GROUP)601 case 11u: { textureStore(mip_12, pix, value); }602 @if(SPLIT_BIND_GROUP_FIRST)603 case 0u: { textureStore(mip_1, pix, value); }604 @if(SPLIT_BIND_GROUP_FIRST)605 case 1u: { textureStore(mip_2, pix, value); }606 @if(SPLIT_BIND_GROUP_FIRST)607 case 2u: { textureStore(mip_3, pix, value); }608 @if(SPLIT_BIND_GROUP_FIRST)609 case 3u: { textureStore(mip_4, pix, value); }610 @if(SPLIT_BIND_GROUP_FIRST)611 case 4u: { textureStore(mip_5, pix, value); }612 @if(SPLIT_BIND_GROUP_FIRST)613 case 5u: { textureStore(mip_6, pix, value); }614 @if(SPLIT_BIND_GROUP_SECOND)615 case 6u: { textureStore(mip_7, pix, value); }616 @if(SPLIT_BIND_GROUP_SECOND)617 case 7u: { textureStore(mip_8, pix, value); }618 @if(SPLIT_BIND_GROUP_SECOND)619 case 8u: { textureStore(mip_9, pix, value); }620 @if(SPLIT_BIND_GROUP_SECOND)621 case 9u: { textureStore(mip_10, pix, value); }622 @if(SPLIT_BIND_GROUP_SECOND)623 case 10u: { textureStore(mip_11, pix, value); }624 @if(SPLIT_BIND_GROUP_SECOND)625 case 11u: { textureStore(mip_12, pix, value); }626 default: {}627 }628}630fn spd_store_intermediate(x: u32, y: u32, value: vec4f) {631 spd_intermediate_r[x][y] = value.x;632 spd_intermediate_g[x][y] = value.y;633 spd_intermediate_b[x][y] = value.z;634 spd_intermediate_a[x][y] = value.w;635}637fn spd_load_intermediate(x: u32, y: u32) -> vec4f {638 return vec4(spd_intermediate_r[x][y], spd_intermediate_g[x][y], spd_intermediate_b[x][y], spd_intermediate_a[x][y]);639}641fn spd_reduce_intermediate(i0: vec2u, i1: vec2u, i2: vec2u, i3: vec2u) -> vec4f {642 let v0 = spd_load_intermediate(i0.x, i0.y);643 let v1 = spd_load_intermediate(i1.x, i1.y);644 let v2 = spd_load_intermediate(i2.x, i2.y);645 let v3 = spd_load_intermediate(i3.x, i3.y);646 return spd_reduce_4(v0, v1, v2, v3);647}649@if(ARRAY_TEXTURE)650fn spd_reduce_load_4(i0: vec2u, i1: vec2u, i2: vec2u, i3: vec2u, slice: u32) -> vec4f {651 @if(COMBINE_BIND_GROUP) {652 let v0 = textureLoad(mip_6, i0, slice);653 let v1 = textureLoad(mip_6, i1, slice);654 let v2 = textureLoad(mip_6, i2, slice);655 let v3 = textureLoad(mip_6, i3, slice);656 return spd_reduce_4(v0, v1, v2, v3);657 }658 @if(SPLIT_BIND_GROUP_FIRST)659 return vec4(0.0, 0.0, 0.0, 0.0);660 @if(SPLIT_BIND_GROUP_SECOND) {661 let v0 = textureLoad(mip_6, i0, slice, 0);662 let v1 = textureLoad(mip_6, i1, slice, 0);663 let v2 = textureLoad(mip_6, i2, slice, 0);664 let v3 = textureLoad(mip_6, i3, slice, 0);665 return spd_reduce_4(v0, v1, v2, v3);666 }667}668@else // ARRAY_TEXTURE669fn spd_reduce_load_4(i0: vec2u, i1: vec2u, i2: vec2u, i3: vec2u, slice: u32) -> vec4f {670 @if(COMBINE_BIND_GROUP) {671 let v0 = textureLoad(mip_6, i0);672 let v1 = textureLoad(mip_6, i1);673 let v2 = textureLoad(mip_6, i2);674 let v3 = textureLoad(mip_6, i3);675 return spd_reduce_4(v0, v1, v2, v3);676 }677 @if(SPLIT_BIND_GROUP_FIRST)678 return vec4(0.0, 0.0, 0.0, 0.0);679 @if(SPLIT_BIND_GROUP_SECOND) {680 let v0 = textureLoad(mip_6, i0, 0);681 let v1 = textureLoad(mip_6, i1, 0);682 let v2 = textureLoad(mip_6, i2, 0);683 let v3 = textureLoad(mip_6, i3, 0);684 return spd_reduce_4(v0, v1, v2, v3);685 }686}688fn spd_reduce_4(v0: vec4f, v1: vec4f, v2: vec4f, v3: vec4f) -> vec4f {689 return (v0 + v1 + v2 + v3) * 0.25;690}692@if(SUBGROUP_SUPPORT)693fn spd_reduce_quad(v: vec4f) -> vec4f {694 let v0 = v;695 let v1 = quadSwapX(v);696 let v2 = quadSwapY(v);697 let v3 = quadSwapDiagonal(v);698 return spd_reduce_4(v0, v1, v2, v3);699}701fn srgb_from_linear(value: f32) -> f32 {702 let j = vec3(0.0031308 * 12.92, 12.92, 1.0 / 2.4);703 let k = vec2(1.055, -0.055);704 return clamp(j.x, value * j.y, pow(value, j.z) * k.x + k.y);705}
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.