1diagnostic(off, derivative_uniformity); 2diagnostic(off, chromium.unreachable_code); 3struct CSIn { 4 @builtin(global_invocation_id) sk_GlobalInvocationID: vec3<u32>, 5}; 6struct inputs { 7 in_data: array<f32>, 8}; 9@group(0) @binding(0) var<storage, read> _storage0 : inputs; 10struct outputs { 11 out_data: array<f32>, 12}; 13@group(0) @binding(1) var<storage, read_write> _storage1 : outputs; 14var<workgroup> shared_data: array<f32, 512>; 15fn store_vIf(i: u32, value: f32) { 16 { 17 shared_data[i] = value; 18 } 19} 20fn _skslMain(_stageIn: CSIn) { 21 { 22 let id: u32 = _stageIn.sk_GlobalInvocationID.x; 23 var rd_id: u32; 24 var wr_id: u32; 25 var mask: u32; 26 let _skTemp2 = id * 2u; 27 let _skTemp3 = id * 2u; 28 shared_data[_skTemp2] = _storage0.in_data[_skTemp3]; 29 let _skTemp4 = id * 2u + 1u; 30 let _skTemp5 = id * 2u + 1u; 31 shared_data[_skTemp4] = _storage0.in_data[_skTemp5]; 32 workgroupBarrier(); 33 const steps: u32 = 9u; 34 { 35 var _0_step: u32 = 0u; 36 loop { 37 { 38 mask = (1u << _0_step) - 1u; 39 rd_id = ((id >> _0_step) << (_0_step + 1u)) + mask; 40 wr_id = (rd_id + 1u) + (id & mask); 41 store_vIf(wr_id, shared_data[wr_id] + shared_data[rd_id]); 42 workgroupBarrier(); 43 } 44 continuing { 45 _0_step = _0_step + u32(1); 46 break if _0_step >= steps; 47 } 48 } 49 } 50 let _skTemp6 = id * 2u; 51 let _skTemp7 = id * 2u; 52 _storage1.out_data[_skTemp6] = shared_data[_skTemp7]; 53 let _skTemp8 = id * 2u + 1u; 54 let _skTemp9 = id * 2u + 1u; 55 _storage1.out_data[_skTemp8] = shared_data[_skTemp9]; 56 } 57} 58@compute @workgroup_size(256, 1, 1) fn main(_stageIn: CSIn) { 59 _skslMain(_stageIn); 60} 61