gfx-rs · teoxoy · Dec 12, 2023 · Dec 12, 2023 · Dec 12, 2023
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -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

@@ -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(_) => {

@@ -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;

@@ -1,4 +1,4 @@
-const constant: i32 = 10;
+const constant: i32 = 10i;
 
 fn function() -> f32 {
     return 0.0;

@@ -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;

@@ -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);
@@ -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);

@@ -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);
@@ -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);
@@ -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);
@@ -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;
@@ -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;
 }
 
@@ -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) 
@@ -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;
@@ -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;
@@ -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];
@@ -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);
@@ -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);
@@ -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);
 }

@@ -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]));
@@ -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;
 }