Skip to content

Commit

Permalink
[naga wgsl-out] Include the i suffix on i32 literals. (#4863)
Browse files Browse the repository at this point in the history 
Without the suffix, `Expression::Literal(Literal::I32)` expressions
get written without any suffix on the decimal number, meaning that
they get re-parsed as AbstractInt values. In theory, this should
always be fine, but since we don't actually support abstract types yet
in all the places we should, having them appear in the output causes
validation problems.
  • Loading branch information
@jimblandy
jimblandy authored Dec 12, 2023
1 parent 4b6a2d1 commit 6dc9cca
Show file tree
Hide file tree
Showing 41 changed files with 357 additions and 355 deletions.
2 changes: 2 additions & 0 deletions CHANGELOG.md
View file
Original file line number Diff line number Diff line change
Expand Up @@ -156,6 +156,8 @@ Passing an owned value `window` to `Surface` will return a `Surface<'static>`. S

- Emit and init `struct` member padding always. By @ErichDonGubler in [#4701](https://github.com/gfx-rs/wgpu/pull/4701).

- In WGSL output, always include the `i` suffix on `i32` literals. By @jimblandy in [#4863](https://github.com/gfx-rs/wgpu/pull/4863).

### Bug Fixes

#### General
Expand Down
2 changes: 1 addition & 1 deletion naga/src/back/wgsl/writer.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1093,7 +1093,7 @@ impl<W: Write> Writer<W> {
// decimal part even it's zero
crate::Literal::F32(value) => write!(self.out, "{:?}", value)?,
crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
crate::Literal::I32(value) => write!(self.out, "{}", value)?,
crate::Literal::I32(value) => write!(self.out, "{}i", value)?,
crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
crate::Literal::F64(value) => write!(self.out, "{:?}lf", value)?,
crate::Literal::I64(_) => {
Expand Down
2 changes: 1 addition & 1 deletion naga/tests/out/wgsl/900-implicit-conversions.frag.wgsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -55,7 +55,7 @@ fn implicit_dims_3(v_6: vec4<f32>) {
}

fn main_1() {
exact_1(1);
exact_1(1i);
implicit(1.0);
implicit_dims_2(vec3(1.0));
return;
Expand Down
2 changes: 1 addition & 1 deletion naga/tests/out/wgsl/931-constant-emitting.frag.wgsl
View file
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
const constant: i32 = 10;
const constant: i32 = 10i;

fn function() -> f32 {
return 0.0;
Expand Down
6 changes: 3 additions & 3 deletions naga/tests/out/wgsl/932-for-loop-if.frag.wgsl
View file
Original file line number Diff line number Diff line change
@@ -1,16 +1,16 @@
fn main_1() {
var i: i32 = 0;
var i: i32 = 0i;

loop {
let _e2 = i;
if !((_e2 < 1)) {
if !((_e2 < 1i)) {
break;
}
{
}
continuing {
let _e6 = i;
i = (_e6 + 1);
i = (_e6 + 1i);
}
}
return;
Expand Down
20 changes: 10 additions & 10 deletions naga/tests/out/wgsl/abstract-types-const.wgsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -18,9 +18,9 @@ const xmfpaiaiaiaf: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0
const imfpaiaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
const imfpafaiaiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
const imfpafafafaf: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
const ivispai: vec2<i32> = vec2(1);
const ivispai: vec2<i32> = vec2(1i);
const ivfspaf: vec2<f32> = vec2(1.0);
const ivis_ai: vec2<i32> = vec2(1);
const ivis_ai: vec2<i32> = vec2(1i);
const ivus_ai: vec2<u32> = vec2(1u);
const ivfs_ai: vec2<f32> = vec2(1.0);
const ivfs_af: vec2<f32> = vec2(1.0);
Expand All @@ -30,14 +30,14 @@ const iafpafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
const iafpaiaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
const iafpafai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
const xafpafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
const s_f_i_u: S = S(1.0, 1, 1u);
const s_f_iai: S = S(1.0, 1, 1u);
const s_fai_u: S = S(1.0, 1, 1u);
const s_faiai: S = S(1.0, 1, 1u);
const saf_i_u: S = S(1.0, 1, 1u);
const saf_iai: S = S(1.0, 1, 1u);
const safai_u: S = S(1.0, 1, 1u);
const safaiai: S = S(1.0, 1, 1u);
const s_f_i_u: S = S(1.0, 1i, 1u);
const s_f_iai: S = S(1.0, 1i, 1u);
const s_fai_u: S = S(1.0, 1i, 1u);
const s_faiai: S = S(1.0, 1i, 1u);
const saf_i_u: S = S(1.0, 1i, 1u);
const saf_iai: S = S(1.0, 1i, 1u);
const safai_u: S = S(1.0, 1i, 1u);
const safaiai: S = S(1.0, 1i, 1u);
const ivfr_f_f: vec3<f32> = vec3<f32>(vec2<f32>(1.0, 2.0), 3.0);
const ivfr_f_af: vec3<f32> = vec3<f32>(vec2<f32>(1.0, 2.0), 3.0);
const ivfraf_f: vec3<f32> = vec3<f32>(vec2<f32>(1.0, 2.0), 3.0);
Expand Down
28 changes: 14 additions & 14 deletions naga/tests/out/wgsl/abstract-types-var.wgsl
View file
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
var<private> xvipaiai_1: vec2<i32> = vec2<i32>(42, 43);
var<private> xvipaiai_1: vec2<i32> = vec2<i32>(42i, 43i);
var<private> xvupaiai_1: vec2<u32> = vec2<u32>(44u, 45u);
var<private> xvfpaiai_1: vec2<f32> = vec2<f32>(46.0, 47.0);
var<private> xvupuai_2: vec2<u32> = vec2<u32>(42u, 43u);
Expand All @@ -10,21 +10,21 @@ var<private> xmfpafaiaiai_1: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2
var<private> xmfpaiafaiai_1: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xmfpaiaiafai_1: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xmfpaiaiaiaf_1: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var<private> xvispai_1: vec2<i32> = vec2(1);
var<private> xvispai_1: vec2<i32> = vec2(1i);
var<private> xvfspaf_1: vec2<f32> = vec2(1.0);
var<private> xvis_ai_1: vec2<i32> = vec2(1);
var<private> xvis_ai_1: vec2<i32> = vec2(1i);
var<private> xvus_ai_1: vec2<u32> = vec2(1u);
var<private> xvfs_ai_1: vec2<f32> = vec2(1.0);
var<private> xvfs_af_1: vec2<f32> = vec2(1.0);
var<private> xafafaf_1: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafaiai_1: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafpaiai_1: array<i32, 2> = array<i32, 2>(1, 2);
var<private> xafpaiai_1: array<i32, 2> = array<i32, 2>(1i, 2i);
var<private> xafpaiaf_1: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafpafai_1: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var<private> xafpafaf_1: array<f32, 2> = array<f32, 2>(1.0, 2.0);

fn all_constant_arguments() {
var xvipaiai: vec2<i32> = vec2<i32>(42, 43);
var xvipaiai: vec2<i32> = vec2<i32>(42i, 43i);
var xvupaiai: vec2<u32> = vec2<u32>(44u, 45u);
var xvfpaiai: vec2<f32> = vec2<f32>(46.0, 47.0);
var xvupuai: vec2<u32> = vec2<u32>(42u, 43u);
Expand All @@ -40,19 +40,19 @@ fn all_constant_arguments() {
var xmfpai_faiai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xmfpaiai_fai: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xmfpaiaiai_f: mat2x2<f32> = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
var xvispai: vec2<i32> = vec2(1);
var xvispai: vec2<i32> = vec2(1i);
var xvfspaf: vec2<f32> = vec2(1.0);
var xvis_ai: vec2<i32> = vec2(1);
var xvis_ai: vec2<i32> = vec2(1i);
var xvus_ai: vec2<u32> = vec2(1u);
var xvfs_ai: vec2<f32> = vec2(1.0);
var xvfs_af: vec2<f32> = vec2(1.0);
var xafafaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xaf_faf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xafaf_f: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xafaiai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xai_iai: array<i32, 2> = array<i32, 2>(1, 2);
var xaiai_i: array<i32, 2> = array<i32, 2>(1, 2);
var xaipaiai: array<i32, 2> = array<i32, 2>(1, 2);
var xai_iai: array<i32, 2> = array<i32, 2>(1i, 2i);
var xaiai_i: array<i32, 2> = array<i32, 2>(1i, 2i);
var xaipaiai: array<i32, 2> = array<i32, 2>(1i, 2i);
var xafpaiai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xafpaiaf: array<f32, 2> = array<f32, 2>(1.0, 2.0);
var xafpafai: array<f32, 2> = array<f32, 2>(1.0, 2.0);
Expand Down Expand Up @@ -110,9 +110,9 @@ fn mixed_constant_and_runtime_arguments() {
let _e63 = f;
xafai_f = array<f32, 2>(1.0, _e63);
let _e67 = i;
xai_iai_1 = array<i32, 2>(_e67, 2);
xai_iai_1 = array<i32, 2>(_e67, 2i);
let _e71 = i;
xaiai_i_1 = array<i32, 2>(1, _e71);
xaiai_i_1 = array<i32, 2>(1i, _e71);
let _e75 = f;
xafp_faf = array<f32, 2>(_e75, 2.0);
let _e79 = f;
Expand All @@ -122,9 +122,9 @@ fn mixed_constant_and_runtime_arguments() {
let _e87 = f;
xafpai_f = array<f32, 2>(1.0, _e87);
let _e91 = i;
xaip_iai = array<i32, 2>(_e91, 2);
xaip_iai = array<i32, 2>(_e91, 2i);
let _e95 = i;
xaipai_i = array<i32, 2>(1, _e95);
xaipai_i = array<i32, 2>(1i, _e95);
return;
}

20 changes: 10 additions & 10 deletions naga/tests/out/wgsl/access.wgsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ struct MatCx2InArray {
am: array<mat4x2<f32>, 2>,
}

var<private> global_const: GlobalConst = GlobalConst(0u, vec3<u32>(0u, 0u, 0u), 0);
var<private> global_const: GlobalConst = GlobalConst(0u, vec3<u32>(0u, 0u, 0u), 0i);
@group(0) @binding(0)
var<storage, read_write> bar: Bar;
@group(0) @binding(1)
Expand All @@ -36,11 +36,11 @@ var<storage, read_write> qux: vec2<i32>;
var<uniform> nested_mat_cx2_: MatCx2InArray;

fn test_matrix_within_struct_accesses() {
var idx: i32 = 1;
var idx: i32 = 1i;
var t: Baz = Baz(mat3x2<f32>(vec2(1.0), vec2(2.0), vec2(3.0)));

let _e3 = idx;
idx = (_e3 - 1);
idx = (_e3 - 1i);
let l0_ = baz.m;
let l1_ = baz.m[0];
let _e14 = idx;
Expand All @@ -54,7 +54,7 @@ fn test_matrix_within_struct_accesses() {
let _e38 = idx;
let l6_ = baz.m[_e36][_e38];
let _e51 = idx;
idx = (_e51 + 1);
idx = (_e51 + 1i);
t.m = mat3x2<f32>(vec2(6.0), vec2(5.0), vec2(4.0));
t.m[0] = vec2(9.0);
let _e66 = idx;
Expand All @@ -71,11 +71,11 @@ fn test_matrix_within_struct_accesses() {
}

fn test_matrix_within_array_within_struct_accesses() {
var idx_1: i32 = 1;
var idx_1: i32 = 1i;
var t_1: MatCx2InArray = MatCx2InArray(array<mat4x2<f32>, 2>());

let _e3 = idx_1;
idx_1 = (_e3 - 1);
idx_1 = (_e3 - 1i);
let l0_1 = nested_mat_cx2_.am;
let l1_1 = nested_mat_cx2_.am[0];
let l2_1 = nested_mat_cx2_.am[0][0];
Expand All @@ -90,7 +90,7 @@ fn test_matrix_within_array_within_struct_accesses() {
let _e48 = idx_1;
let l7_ = nested_mat_cx2_.am[0][_e46][_e48];
let _e55 = idx_1;
idx_1 = (_e55 + 1);
idx_1 = (_e55 + 1i);
t_1.am = array<mat4x2<f32>, 2>();
t_1.am[0] = mat4x2<f32>(vec2(8.0), vec2(7.0), vec2(6.0), vec2(5.0));
t_1.am[0][0] = vec2(9.0);
Expand Down Expand Up @@ -142,8 +142,8 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
let c = qux;
let data_pointer = (&bar.data[0].value);
let _e33 = read_from_private((&foo));
c2_ = array<i32, 5>(a_1, i32(b), 3, 4, 5);
c2_[(vi + 1u)] = 42;
c2_ = array<i32, 5>(a_1, i32(b), 3i, 4i, 5i);
c2_[(vi + 1u)] = 42i;
let value = c2_[vi];
let _e47 = test_arr_as_arg(array<array<f32, 10>, 5>());
return vec4<f32>((_matrix * vec4<f32>(vec4(value))), 2.0);
Expand All @@ -154,7 +154,7 @@ fn foo_frag() -> @location(0) vec4<f32> {
bar._matrix[1][2] = 1.0;
bar._matrix = mat4x3<f32>(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0));
bar.arr = array<vec2<u32>, 2>(vec2(0u), vec2(1u));
bar.data[1].value = 1;
bar.data[1].value = 1i;
qux = vec2<i32>();
return vec4(0.0);
}
Expand Down
72 changes: 36 additions & 36 deletions naga/tests/out/wgsl/atomicOps.wgsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -16,13 +16,13 @@ var<workgroup> workgroup_struct: Struct;
@compute @workgroup_size(2, 1, 1)
fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
atomicStore((&storage_atomic_scalar), 1u);
atomicStore((&storage_atomic_arr[1]), 1);
atomicStore((&storage_atomic_arr[1]), 1i);
atomicStore((&storage_struct.atomic_scalar), 1u);
atomicStore((&storage_struct.atomic_arr[1]), 1);
atomicStore((&storage_struct.atomic_arr[1]), 1i);
atomicStore((&workgroup_atomic_scalar), 1u);
atomicStore((&workgroup_atomic_arr[1]), 1);
atomicStore((&workgroup_atomic_arr[1]), 1i);
atomicStore((&workgroup_struct.atomic_scalar), 1u);
atomicStore((&workgroup_struct.atomic_arr[1]), 1);
atomicStore((&workgroup_struct.atomic_arr[1]), 1i);
workgroupBarrier();
let l0_ = atomicLoad((&storage_atomic_scalar));
let l1_ = atomicLoad((&storage_atomic_arr[1]));
Expand All @@ -34,74 +34,74 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
let l7_ = atomicLoad((&workgroup_struct.atomic_arr[1]));
workgroupBarrier();
let _e51 = atomicAdd((&storage_atomic_scalar), 1u);
let _e55 = atomicAdd((&storage_atomic_arr[1]), 1);
let _e55 = atomicAdd((&storage_atomic_arr[1]), 1i);
let _e59 = atomicAdd((&storage_struct.atomic_scalar), 1u);
let _e64 = atomicAdd((&storage_struct.atomic_arr[1]), 1);
let _e64 = atomicAdd((&storage_struct.atomic_arr[1]), 1i);
let _e67 = atomicAdd((&workgroup_atomic_scalar), 1u);
let _e71 = atomicAdd((&workgroup_atomic_arr[1]), 1);
let _e71 = atomicAdd((&workgroup_atomic_arr[1]), 1i);
let _e75 = atomicAdd((&workgroup_struct.atomic_scalar), 1u);
let _e80 = atomicAdd((&workgroup_struct.atomic_arr[1]), 1);
let _e80 = atomicAdd((&workgroup_struct.atomic_arr[1]), 1i);
workgroupBarrier();
let _e83 = atomicSub((&storage_atomic_scalar), 1u);
let _e87 = atomicSub((&storage_atomic_arr[1]), 1);
let _e87 = atomicSub((&storage_atomic_arr[1]), 1i);
let _e91 = atomicSub((&storage_struct.atomic_scalar), 1u);
let _e96 = atomicSub((&storage_struct.atomic_arr[1]), 1);
let _e96 = atomicSub((&storage_struct.atomic_arr[1]), 1i);
let _e99 = atomicSub((&workgroup_atomic_scalar), 1u);
let _e103 = atomicSub((&workgroup_atomic_arr[1]), 1);
let _e103 = atomicSub((&workgroup_atomic_arr[1]), 1i);
let _e107 = atomicSub((&workgroup_struct.atomic_scalar), 1u);
let _e112 = atomicSub((&workgroup_struct.atomic_arr[1]), 1);
let _e112 = atomicSub((&workgroup_struct.atomic_arr[1]), 1i);
workgroupBarrier();
let _e115 = atomicMax((&storage_atomic_scalar), 1u);
let _e119 = atomicMax((&storage_atomic_arr[1]), 1);
let _e119 = atomicMax((&storage_atomic_arr[1]), 1i);
let _e123 = atomicMax((&storage_struct.atomic_scalar), 1u);
let _e128 = atomicMax((&storage_struct.atomic_arr[1]), 1);
let _e128 = atomicMax((&storage_struct.atomic_arr[1]), 1i);
let _e131 = atomicMax((&workgroup_atomic_scalar), 1u);
let _e135 = atomicMax((&workgroup_atomic_arr[1]), 1);
let _e135 = atomicMax((&workgroup_atomic_arr[1]), 1i);
let _e139 = atomicMax((&workgroup_struct.atomic_scalar), 1u);
let _e144 = atomicMax((&workgroup_struct.atomic_arr[1]), 1);
let _e144 = atomicMax((&workgroup_struct.atomic_arr[1]), 1i);
workgroupBarrier();
let _e147 = atomicMin((&storage_atomic_scalar), 1u);
let _e151 = atomicMin((&storage_atomic_arr[1]), 1);
let _e151 = atomicMin((&storage_atomic_arr[1]), 1i);
let _e155 = atomicMin((&storage_struct.atomic_scalar), 1u);
let _e160 = atomicMin((&storage_struct.atomic_arr[1]), 1);
let _e160 = atomicMin((&storage_struct.atomic_arr[1]), 1i);
let _e163 = atomicMin((&workgroup_atomic_scalar), 1u);
let _e167 = atomicMin((&workgroup_atomic_arr[1]), 1);
let _e167 = atomicMin((&workgroup_atomic_arr[1]), 1i);
let _e171 = atomicMin((&workgroup_struct.atomic_scalar), 1u);
let _e176 = atomicMin((&workgroup_struct.atomic_arr[1]), 1);
let _e176 = atomicMin((&workgroup_struct.atomic_arr[1]), 1i);
workgroupBarrier();
let _e179 = atomicAnd((&storage_atomic_scalar), 1u);
let _e183 = atomicAnd((&storage_atomic_arr[1]), 1);
let _e183 = atomicAnd((&storage_atomic_arr[1]), 1i);
let _e187 = atomicAnd((&storage_struct.atomic_scalar), 1u);
let _e192 = atomicAnd((&storage_struct.atomic_arr[1]), 1);
let _e192 = atomicAnd((&storage_struct.atomic_arr[1]), 1i);
let _e195 = atomicAnd((&workgroup_atomic_scalar), 1u);
let _e199 = atomicAnd((&workgroup_atomic_arr[1]), 1);
let _e199 = atomicAnd((&workgroup_atomic_arr[1]), 1i);
let _e203 = atomicAnd((&workgroup_struct.atomic_scalar), 1u);
let _e208 = atomicAnd((&workgroup_struct.atomic_arr[1]), 1);
let _e208 = atomicAnd((&workgroup_struct.atomic_arr[1]), 1i);
workgroupBarrier();
let _e211 = atomicOr((&storage_atomic_scalar), 1u);
let _e215 = atomicOr((&storage_atomic_arr[1]), 1);
let _e215 = atomicOr((&storage_atomic_arr[1]), 1i);
let _e219 = atomicOr((&storage_struct.atomic_scalar), 1u);
let _e224 = atomicOr((&storage_struct.atomic_arr[1]), 1);
let _e224 = atomicOr((&storage_struct.atomic_arr[1]), 1i);
let _e227 = atomicOr((&workgroup_atomic_scalar), 1u);
let _e231 = atomicOr((&workgroup_atomic_arr[1]), 1);
let _e231 = atomicOr((&workgroup_atomic_arr[1]), 1i);
let _e235 = atomicOr((&workgroup_struct.atomic_scalar), 1u);
let _e240 = atomicOr((&workgroup_struct.atomic_arr[1]), 1);
let _e240 = atomicOr((&workgroup_struct.atomic_arr[1]), 1i);
workgroupBarrier();
let _e243 = atomicXor((&storage_atomic_scalar), 1u);
let _e247 = atomicXor((&storage_atomic_arr[1]), 1);
let _e247 = atomicXor((&storage_atomic_arr[1]), 1i);
let _e251 = atomicXor((&storage_struct.atomic_scalar), 1u);
let _e256 = atomicXor((&storage_struct.atomic_arr[1]), 1);
let _e256 = atomicXor((&storage_struct.atomic_arr[1]), 1i);
let _e259 = atomicXor((&workgroup_atomic_scalar), 1u);
let _e263 = atomicXor((&workgroup_atomic_arr[1]), 1);
let _e263 = atomicXor((&workgroup_atomic_arr[1]), 1i);
let _e267 = atomicXor((&workgroup_struct.atomic_scalar), 1u);
let _e272 = atomicXor((&workgroup_struct.atomic_arr[1]), 1);
let _e272 = atomicXor((&workgroup_struct.atomic_arr[1]), 1i);
let _e275 = atomicExchange((&storage_atomic_scalar), 1u);
let _e279 = atomicExchange((&storage_atomic_arr[1]), 1);
let _e279 = atomicExchange((&storage_atomic_arr[1]), 1i);
let _e283 = atomicExchange((&storage_struct.atomic_scalar), 1u);
let _e288 = atomicExchange((&storage_struct.atomic_arr[1]), 1);
let _e288 = atomicExchange((&storage_struct.atomic_arr[1]), 1i);
let _e291 = atomicExchange((&workgroup_atomic_scalar), 1u);
let _e295 = atomicExchange((&workgroup_atomic_arr[1]), 1);
let _e295 = atomicExchange((&workgroup_atomic_arr[1]), 1i);
let _e299 = atomicExchange((&workgroup_struct.atomic_scalar), 1u);
let _e304 = atomicExchange((&workgroup_struct.atomic_arr[1]), 1);
let _e304 = atomicExchange((&workgroup_struct.atomic_arr[1]), 1i);
return;
}
Loading

0 comments on commit 6dc9cca

Please sign in to comment.