blob: 95aeaef38212270a02f12aed95da700af50c6b02 [file] [log] [blame]
diagnostic(off, derivative_uniformity);
var<private> u_xlatu0 : vec2<u32>;
var<private> u_xlati4 : i32;
var<private> u_xlat4 : f32;
struct S {
m : array<u32>,
}
@group(0u) @binding(0u) var<storage, read> v : S;
struct tint_symbol {
value : array<u32, 1u>,
}
var<workgroup> TGSM0 : array<tint_symbol, 128u>;
var<private> u_xlatb1 : vec4<bool>;
var<private> u_xlati6 : i32;
var<private> u_xlatb0 : bool;
var<private> u_xlati0 : i32;
var<private> u_xlati2 : i32;
struct S_1 {
m_1 : array<u32>,
}
@group(0u) @binding(1u) var<storage, read_write> v_1 : S_1;
@compute @workgroup_size(128u, 1u, 1u)
fn main(@builtin(local_invocation_index) gl_LocalInvocationIndex : u32, @builtin(global_invocation_id) gl_GlobalInvocationID : vec3<u32>) {
var param : i32;
var param_1 : i32;
var param_2 : i32;
var param_3 : i32;
var u_xlat_precise_vec4 : vec4<f32>;
var u_xlat_precise_ivec4 : vec4<i32>;
var u_xlat_precise_bvec4 : vec4<bool>;
var u_xlat_precise_uvec4 : vec4<u32>;
u_xlatu0 = (vec2<u32>(gl_LocalInvocationIndex, gl_LocalInvocationIndex) & vec2<u32>(31u, 96u));
u_xlati4 = (bitcast<i32>(gl_GlobalInvocationID.x) << bitcast<u32>(2i));
u_xlat4 = bitcast<f32>(v.m[((u_xlati4 >> bitcast<u32>(2i)) + 0i)]);
u_xlati4 = bitcast<i32>(select(0u, 4294967295u, (bitcast<i32>(u_xlat4) == bitcast<i32>(gl_LocalInvocationIndex))));
param = 0i;
param_1 = u_xlati4;
param_2 = (bitcast<i32>(gl_LocalInvocationIndex) & 31i);
param_3 = 1i;
u_xlati4 = v_2(&(param), &(param_1), &(param_2), &(param_3));
TGSM0[gl_LocalInvocationIndex].value[0i] = bitcast<u32>(u_xlati4);
workgroupBarrier();
u_xlatb1 = (u_xlatu0.xxxx < vec4<u32>(16u, 8u, 4u, 2u));
if (u_xlatb1.x) {
u_xlati4 = bitcast<i32>(TGSM0[gl_LocalInvocationIndex].value[0i]);
u_xlati6 = (bitcast<i32>(gl_LocalInvocationIndex) + 16i);
u_xlati6 = bitcast<i32>(TGSM0[u_xlati6].value[0i]);
u_xlati4 = bitcast<i32>((bitcast<u32>(u_xlati6) | bitcast<u32>(u_xlati4)));
TGSM0[gl_LocalInvocationIndex].value[0i] = bitcast<u32>(u_xlati4);
}
workgroupBarrier();
if (u_xlatb1.y) {
u_xlati4 = bitcast<i32>(TGSM0[gl_LocalInvocationIndex].value[0i]);
u_xlati6 = (bitcast<i32>(gl_LocalInvocationIndex) + 8i);
u_xlati6 = bitcast<i32>(TGSM0[u_xlati6].value[0i]);
u_xlati4 = bitcast<i32>((bitcast<u32>(u_xlati6) | bitcast<u32>(u_xlati4)));
TGSM0[gl_LocalInvocationIndex].value[0i] = bitcast<u32>(u_xlati4);
}
workgroupBarrier();
if (u_xlatb1.z) {
u_xlati4 = bitcast<i32>(TGSM0[gl_LocalInvocationIndex].value[0i]);
u_xlati6 = (bitcast<i32>(gl_LocalInvocationIndex) + 4i);
u_xlati6 = bitcast<i32>(TGSM0[u_xlati6].value[0i]);
u_xlati4 = bitcast<i32>((bitcast<u32>(u_xlati6) | bitcast<u32>(u_xlati4)));
TGSM0[gl_LocalInvocationIndex].value[0i] = bitcast<u32>(u_xlati4);
}
workgroupBarrier();
if (u_xlatb1.w) {
u_xlati4 = bitcast<i32>(TGSM0[gl_LocalInvocationIndex].value[0i]);
u_xlati6 = (bitcast<i32>(gl_LocalInvocationIndex) + 2i);
u_xlati6 = bitcast<i32>(TGSM0[u_xlati6].value[0i]);
u_xlati4 = bitcast<i32>((bitcast<u32>(u_xlati6) | bitcast<u32>(u_xlati4)));
TGSM0[gl_LocalInvocationIndex].value[0i] = bitcast<u32>(u_xlati4);
}
workgroupBarrier();
u_xlatb0 = (u_xlatu0.x < 1u);
if (u_xlatb0) {
u_xlati0 = bitcast<i32>(TGSM0[gl_LocalInvocationIndex].value[0i]);
u_xlati4 = (bitcast<i32>(gl_LocalInvocationIndex) + 1i);
u_xlati4 = bitcast<i32>(TGSM0[u_xlati4].value[0i]);
u_xlati0 = bitcast<i32>((bitcast<u32>(u_xlati4) | bitcast<u32>(u_xlati0)));
TGSM0[gl_LocalInvocationIndex].value[0i] = bitcast<u32>(u_xlati0);
}
workgroupBarrier();
u_xlati0 = bitcast<i32>(TGSM0[u_xlatu0.y].value[0i]);
u_xlati0 = countOneBits(u_xlati0);
u_xlatb0 = (u_xlati0 == 32i);
u_xlati0 = select(0i, 1i, u_xlatb0);
u_xlati2 = (bitcast<i32>(gl_LocalInvocationIndex) << bitcast<u32>(2i));
let v_3 = (u_xlati2 >> bitcast<u32>(2i));
v_1.m_1[v_3] = bitcast<u32>(u_xlati0);
}
fn v_2(base : ptr<function, i32>, insert : ptr<function, i32>, offset : ptr<function, i32>, bits : ptr<function, i32>) -> i32 {
var mask : u32;
mask = (~((4294967295u << bitcast<u32>(*(bits)))) << bitcast<u32>(*(offset)));
return bitcast<i32>(((bitcast<u32>(*(base)) & ~(mask)) | ((bitcast<u32>(*(insert)) << bitcast<u32>(*(offset))) & mask)));
}