1 // Copyright 2021 The Tint Authors.
2 //
3 // Licensed under the Apache License, Version 2.0 (the "License");
4 // you may not use this file except in compliance with the License.
5 // You may obtain a copy of the License at
6 //
7 // http://www.apache.org/licenses/LICENSE-2.0
8 //
9 // Unless required by applicable law or agreed to in writing, software
10 // distributed under the License is distributed on an "AS IS" BASIS,
11 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 // See the License for the specific language governing permissions and
13 // limitations under the License.
14
15 #include "src/transform/decompose_memory_access.h"
16
17 #include "src/transform/test_helper.h"
18
19 namespace tint {
20 namespace transform {
21 namespace {
22
23 using DecomposeMemoryAccessTest = TransformTest;
24
TEST_F(DecomposeMemoryAccessTest,SB_BasicLoad)25 TEST_F(DecomposeMemoryAccessTest, SB_BasicLoad) {
26 auto* src = R"(
27 [[block]]
28 struct SB {
29 a : i32;
30 b : u32;
31 c : f32;
32 d : vec2<i32>;
33 e : vec2<u32>;
34 f : vec2<f32>;
35 g : vec3<i32>;
36 h : vec3<u32>;
37 i : vec3<f32>;
38 j : vec4<i32>;
39 k : vec4<u32>;
40 l : vec4<f32>;
41 m : mat2x2<f32>;
42 n : mat2x3<f32>;
43 o : mat2x4<f32>;
44 p : mat3x2<f32>;
45 q : mat3x3<f32>;
46 r : mat3x4<f32>;
47 s : mat4x2<f32>;
48 t : mat4x3<f32>;
49 u : mat4x4<f32>;
50 v : array<vec3<f32>, 2>;
51 };
52
53 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
54
55 [[stage(compute), workgroup_size(1)]]
56 fn main() {
57 var a : i32 = sb.a;
58 var b : u32 = sb.b;
59 var c : f32 = sb.c;
60 var d : vec2<i32> = sb.d;
61 var e : vec2<u32> = sb.e;
62 var f : vec2<f32> = sb.f;
63 var g : vec3<i32> = sb.g;
64 var h : vec3<u32> = sb.h;
65 var i : vec3<f32> = sb.i;
66 var j : vec4<i32> = sb.j;
67 var k : vec4<u32> = sb.k;
68 var l : vec4<f32> = sb.l;
69 var m : mat2x2<f32> = sb.m;
70 var n : mat2x3<f32> = sb.n;
71 var o : mat2x4<f32> = sb.o;
72 var p : mat3x2<f32> = sb.p;
73 var q : mat3x3<f32> = sb.q;
74 var r : mat3x4<f32> = sb.r;
75 var s : mat4x2<f32> = sb.s;
76 var t : mat4x3<f32> = sb.t;
77 var u : mat4x4<f32> = sb.u;
78 var v : array<vec3<f32>, 2> = sb.v;
79 }
80 )";
81
82 auto* expect = R"(
83 [[block]]
84 struct SB {
85 a : i32;
86 b : u32;
87 c : f32;
88 d : vec2<i32>;
89 e : vec2<u32>;
90 f : vec2<f32>;
91 g : vec3<i32>;
92 h : vec3<u32>;
93 i : vec3<f32>;
94 j : vec4<i32>;
95 k : vec4<u32>;
96 l : vec4<f32>;
97 m : mat2x2<f32>;
98 n : mat2x3<f32>;
99 o : mat2x4<f32>;
100 p : mat3x2<f32>;
101 q : mat3x3<f32>;
102 r : mat3x4<f32>;
103 s : mat4x2<f32>;
104 t : mat4x3<f32>;
105 u : mat4x4<f32>;
106 v : array<vec3<f32>, 2>;
107 };
108
109 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
110
111 [[internal(intrinsic_load_storage_i32), internal(disable_validation__function_has_no_body)]]
112 fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32
113
114 [[internal(intrinsic_load_storage_u32), internal(disable_validation__function_has_no_body)]]
115 fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
116
117 [[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]]
118 fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
119
120 [[internal(intrinsic_load_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
121 fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<i32>
122
123 [[internal(intrinsic_load_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
124 fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<u32>
125
126 [[internal(intrinsic_load_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
127 fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<f32>
128
129 [[internal(intrinsic_load_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
130 fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<i32>
131
132 [[internal(intrinsic_load_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
133 fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<u32>
134
135 [[internal(intrinsic_load_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
136 fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<f32>
137
138 [[internal(intrinsic_load_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
139 fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<i32>
140
141 [[internal(intrinsic_load_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
142 fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<u32>
143
144 [[internal(intrinsic_load_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
145 fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<f32>
146
147 fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x2<f32> {
148 return mat2x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)));
149 }
150
151 fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x3<f32> {
152 return mat2x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
153 }
154
155 fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x4<f32> {
156 return mat2x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)));
157 }
158
159 fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x2<f32> {
160 return mat3x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)));
161 }
162
163 fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x3<f32> {
164 return mat3x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)));
165 }
166
167 fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x4<f32> {
168 return mat3x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)));
169 }
170
171 fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x2<f32> {
172 return mat4x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)));
173 }
174
175 fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x3<f32> {
176 return mat4x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)), tint_symbol_8(buffer, (offset + 48u)));
177 }
178
179 fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x4<f32> {
180 return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
181 }
182
183 fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2u> {
184 var arr : array<vec3<f32>, 2u>;
185 for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
186 arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u)));
187 }
188 return arr;
189 }
190
191 [[stage(compute), workgroup_size(1)]]
192 fn main() {
193 var a : i32 = tint_symbol(sb, 0u);
194 var b : u32 = tint_symbol_1(sb, 4u);
195 var c : f32 = tint_symbol_2(sb, 8u);
196 var d : vec2<i32> = tint_symbol_3(sb, 16u);
197 var e : vec2<u32> = tint_symbol_4(sb, 24u);
198 var f : vec2<f32> = tint_symbol_5(sb, 32u);
199 var g : vec3<i32> = tint_symbol_6(sb, 48u);
200 var h : vec3<u32> = tint_symbol_7(sb, 64u);
201 var i : vec3<f32> = tint_symbol_8(sb, 80u);
202 var j : vec4<i32> = tint_symbol_9(sb, 96u);
203 var k : vec4<u32> = tint_symbol_10(sb, 112u);
204 var l : vec4<f32> = tint_symbol_11(sb, 128u);
205 var m : mat2x2<f32> = tint_symbol_12(sb, 144u);
206 var n : mat2x3<f32> = tint_symbol_13(sb, 160u);
207 var o : mat2x4<f32> = tint_symbol_14(sb, 192u);
208 var p : mat3x2<f32> = tint_symbol_15(sb, 224u);
209 var q : mat3x3<f32> = tint_symbol_16(sb, 256u);
210 var r : mat3x4<f32> = tint_symbol_17(sb, 304u);
211 var s : mat4x2<f32> = tint_symbol_18(sb, 352u);
212 var t : mat4x3<f32> = tint_symbol_19(sb, 384u);
213 var u : mat4x4<f32> = tint_symbol_20(sb, 448u);
214 var v : array<vec3<f32>, 2> = tint_symbol_21(sb, 512u);
215 }
216 )";
217
218 auto got = Run<DecomposeMemoryAccess>(src);
219
220 EXPECT_EQ(expect, str(got));
221 }
222
TEST_F(DecomposeMemoryAccessTest,UB_BasicLoad)223 TEST_F(DecomposeMemoryAccessTest, UB_BasicLoad) {
224 auto* src = R"(
225 [[block]]
226 struct UB {
227 a : i32;
228 b : u32;
229 c : f32;
230 d : vec2<i32>;
231 e : vec2<u32>;
232 f : vec2<f32>;
233 g : vec3<i32>;
234 h : vec3<u32>;
235 i : vec3<f32>;
236 j : vec4<i32>;
237 k : vec4<u32>;
238 l : vec4<f32>;
239 m : mat2x2<f32>;
240 n : mat2x3<f32>;
241 o : mat2x4<f32>;
242 p : mat3x2<f32>;
243 q : mat3x3<f32>;
244 r : mat3x4<f32>;
245 s : mat4x2<f32>;
246 t : mat4x3<f32>;
247 u : mat4x4<f32>;
248 v : array<vec3<f32>, 2>;
249 };
250
251 [[group(0), binding(0)]] var<uniform> ub : UB;
252
253 [[stage(compute), workgroup_size(1)]]
254 fn main() {
255 var a : i32 = ub.a;
256 var b : u32 = ub.b;
257 var c : f32 = ub.c;
258 var d : vec2<i32> = ub.d;
259 var e : vec2<u32> = ub.e;
260 var f : vec2<f32> = ub.f;
261 var g : vec3<i32> = ub.g;
262 var h : vec3<u32> = ub.h;
263 var i : vec3<f32> = ub.i;
264 var j : vec4<i32> = ub.j;
265 var k : vec4<u32> = ub.k;
266 var l : vec4<f32> = ub.l;
267 var m : mat2x2<f32> = ub.m;
268 var n : mat2x3<f32> = ub.n;
269 var o : mat2x4<f32> = ub.o;
270 var p : mat3x2<f32> = ub.p;
271 var q : mat3x3<f32> = ub.q;
272 var r : mat3x4<f32> = ub.r;
273 var s : mat4x2<f32> = ub.s;
274 var t : mat4x3<f32> = ub.t;
275 var u : mat4x4<f32> = ub.u;
276 var v : array<vec3<f32>, 2> = ub.v;
277 }
278 )";
279
280 auto* expect = R"(
281 [[block]]
282 struct UB {
283 a : i32;
284 b : u32;
285 c : f32;
286 d : vec2<i32>;
287 e : vec2<u32>;
288 f : vec2<f32>;
289 g : vec3<i32>;
290 h : vec3<u32>;
291 i : vec3<f32>;
292 j : vec4<i32>;
293 k : vec4<u32>;
294 l : vec4<f32>;
295 m : mat2x2<f32>;
296 n : mat2x3<f32>;
297 o : mat2x4<f32>;
298 p : mat3x2<f32>;
299 q : mat3x3<f32>;
300 r : mat3x4<f32>;
301 s : mat4x2<f32>;
302 t : mat4x3<f32>;
303 u : mat4x4<f32>;
304 v : array<vec3<f32>, 2>;
305 };
306
307 [[group(0), binding(0)]] var<uniform> ub : UB;
308
309 [[internal(intrinsic_load_uniform_i32), internal(disable_validation__function_has_no_body)]]
310 fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> i32
311
312 [[internal(intrinsic_load_uniform_u32), internal(disable_validation__function_has_no_body)]]
313 fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> u32
314
315 [[internal(intrinsic_load_uniform_f32), internal(disable_validation__function_has_no_body)]]
316 fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> f32
317
318 [[internal(intrinsic_load_uniform_vec2_i32), internal(disable_validation__function_has_no_body)]]
319 fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec2<i32>
320
321 [[internal(intrinsic_load_uniform_vec2_u32), internal(disable_validation__function_has_no_body)]]
322 fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec2<u32>
323
324 [[internal(intrinsic_load_uniform_vec2_f32), internal(disable_validation__function_has_no_body)]]
325 fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec2<f32>
326
327 [[internal(intrinsic_load_uniform_vec3_i32), internal(disable_validation__function_has_no_body)]]
328 fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec3<i32>
329
330 [[internal(intrinsic_load_uniform_vec3_u32), internal(disable_validation__function_has_no_body)]]
331 fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec3<u32>
332
333 [[internal(intrinsic_load_uniform_vec3_f32), internal(disable_validation__function_has_no_body)]]
334 fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec3<f32>
335
336 [[internal(intrinsic_load_uniform_vec4_i32), internal(disable_validation__function_has_no_body)]]
337 fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec4<i32>
338
339 [[internal(intrinsic_load_uniform_vec4_u32), internal(disable_validation__function_has_no_body)]]
340 fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec4<u32>
341
342 [[internal(intrinsic_load_uniform_vec4_f32), internal(disable_validation__function_has_no_body)]]
343 fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec4<f32>
344
345 fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat2x2<f32> {
346 return mat2x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)));
347 }
348
349 fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat2x3<f32> {
350 return mat2x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
351 }
352
353 fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat2x4<f32> {
354 return mat2x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)));
355 }
356
357 fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat3x2<f32> {
358 return mat3x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)));
359 }
360
361 fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat3x3<f32> {
362 return mat3x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)));
363 }
364
365 fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat3x4<f32> {
366 return mat3x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)));
367 }
368
369 fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat4x2<f32> {
370 return mat4x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)));
371 }
372
373 fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat4x3<f32> {
374 return mat4x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)), tint_symbol_8(buffer, (offset + 48u)));
375 }
376
377 fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat4x4<f32> {
378 return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
379 }
380
381 fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> array<vec3<f32>, 2u> {
382 var arr : array<vec3<f32>, 2u>;
383 for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
384 arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u)));
385 }
386 return arr;
387 }
388
389 [[stage(compute), workgroup_size(1)]]
390 fn main() {
391 var a : i32 = tint_symbol(ub, 0u);
392 var b : u32 = tint_symbol_1(ub, 4u);
393 var c : f32 = tint_symbol_2(ub, 8u);
394 var d : vec2<i32> = tint_symbol_3(ub, 16u);
395 var e : vec2<u32> = tint_symbol_4(ub, 24u);
396 var f : vec2<f32> = tint_symbol_5(ub, 32u);
397 var g : vec3<i32> = tint_symbol_6(ub, 48u);
398 var h : vec3<u32> = tint_symbol_7(ub, 64u);
399 var i : vec3<f32> = tint_symbol_8(ub, 80u);
400 var j : vec4<i32> = tint_symbol_9(ub, 96u);
401 var k : vec4<u32> = tint_symbol_10(ub, 112u);
402 var l : vec4<f32> = tint_symbol_11(ub, 128u);
403 var m : mat2x2<f32> = tint_symbol_12(ub, 144u);
404 var n : mat2x3<f32> = tint_symbol_13(ub, 160u);
405 var o : mat2x4<f32> = tint_symbol_14(ub, 192u);
406 var p : mat3x2<f32> = tint_symbol_15(ub, 224u);
407 var q : mat3x3<f32> = tint_symbol_16(ub, 256u);
408 var r : mat3x4<f32> = tint_symbol_17(ub, 304u);
409 var s : mat4x2<f32> = tint_symbol_18(ub, 352u);
410 var t : mat4x3<f32> = tint_symbol_19(ub, 384u);
411 var u : mat4x4<f32> = tint_symbol_20(ub, 448u);
412 var v : array<vec3<f32>, 2> = tint_symbol_21(ub, 512u);
413 }
414 )";
415
416 auto got = Run<DecomposeMemoryAccess>(src);
417
418 EXPECT_EQ(expect, str(got));
419 }
420
TEST_F(DecomposeMemoryAccessTest,SB_BasicStore)421 TEST_F(DecomposeMemoryAccessTest, SB_BasicStore) {
422 auto* src = R"(
423 [[block]]
424 struct SB {
425 a : i32;
426 b : u32;
427 c : f32;
428 d : vec2<i32>;
429 e : vec2<u32>;
430 f : vec2<f32>;
431 g : vec3<i32>;
432 h : vec3<u32>;
433 i : vec3<f32>;
434 j : vec4<i32>;
435 k : vec4<u32>;
436 l : vec4<f32>;
437 m : mat2x2<f32>;
438 n : mat2x3<f32>;
439 o : mat2x4<f32>;
440 p : mat3x2<f32>;
441 q : mat3x3<f32>;
442 r : mat3x4<f32>;
443 s : mat4x2<f32>;
444 t : mat4x3<f32>;
445 u : mat4x4<f32>;
446 v : array<vec3<f32>, 2>;
447 };
448
449 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
450
451 [[stage(compute), workgroup_size(1)]]
452 fn main() {
453 sb.a = i32();
454 sb.b = u32();
455 sb.c = f32();
456 sb.d = vec2<i32>();
457 sb.e = vec2<u32>();
458 sb.f = vec2<f32>();
459 sb.g = vec3<i32>();
460 sb.h = vec3<u32>();
461 sb.i = vec3<f32>();
462 sb.j = vec4<i32>();
463 sb.k = vec4<u32>();
464 sb.l = vec4<f32>();
465 sb.m = mat2x2<f32>();
466 sb.n = mat2x3<f32>();
467 sb.o = mat2x4<f32>();
468 sb.p = mat3x2<f32>();
469 sb.q = mat3x3<f32>();
470 sb.r = mat3x4<f32>();
471 sb.s = mat4x2<f32>();
472 sb.t = mat4x3<f32>();
473 sb.u = mat4x4<f32>();
474 sb.v = array<vec3<f32>, 2>();
475 }
476 )";
477
478 auto* expect = R"(
479 [[block]]
480 struct SB {
481 a : i32;
482 b : u32;
483 c : f32;
484 d : vec2<i32>;
485 e : vec2<u32>;
486 f : vec2<f32>;
487 g : vec3<i32>;
488 h : vec3<u32>;
489 i : vec3<f32>;
490 j : vec4<i32>;
491 k : vec4<u32>;
492 l : vec4<f32>;
493 m : mat2x2<f32>;
494 n : mat2x3<f32>;
495 o : mat2x4<f32>;
496 p : mat3x2<f32>;
497 q : mat3x3<f32>;
498 r : mat3x4<f32>;
499 s : mat4x2<f32>;
500 t : mat4x3<f32>;
501 u : mat4x4<f32>;
502 v : array<vec3<f32>, 2>;
503 };
504
505 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
506
507 [[internal(intrinsic_store_storage_i32), internal(disable_validation__function_has_no_body)]]
508 fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : i32)
509
510 [[internal(intrinsic_store_storage_u32), internal(disable_validation__function_has_no_body)]]
511 fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : u32)
512
513 [[internal(intrinsic_store_storage_f32), internal(disable_validation__function_has_no_body)]]
514 fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : f32)
515
516 [[internal(intrinsic_store_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
517 fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<i32>)
518
519 [[internal(intrinsic_store_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
520 fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<u32>)
521
522 [[internal(intrinsic_store_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
523 fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<f32>)
524
525 [[internal(intrinsic_store_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
526 fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<i32>)
527
528 [[internal(intrinsic_store_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
529 fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<u32>)
530
531 [[internal(intrinsic_store_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
532 fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<f32>)
533
534 [[internal(intrinsic_store_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
535 fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<i32>)
536
537 [[internal(intrinsic_store_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
538 fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<u32>)
539
540 [[internal(intrinsic_store_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
541 fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<f32>)
542
543 fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x2<f32>) {
544 tint_symbol_5(buffer, (offset + 0u), value[0u]);
545 tint_symbol_5(buffer, (offset + 8u), value[1u]);
546 }
547
548 fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x3<f32>) {
549 tint_symbol_8(buffer, (offset + 0u), value[0u]);
550 tint_symbol_8(buffer, (offset + 16u), value[1u]);
551 }
552
553 fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x4<f32>) {
554 tint_symbol_11(buffer, (offset + 0u), value[0u]);
555 tint_symbol_11(buffer, (offset + 16u), value[1u]);
556 }
557
558 fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x2<f32>) {
559 tint_symbol_5(buffer, (offset + 0u), value[0u]);
560 tint_symbol_5(buffer, (offset + 8u), value[1u]);
561 tint_symbol_5(buffer, (offset + 16u), value[2u]);
562 }
563
564 fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x3<f32>) {
565 tint_symbol_8(buffer, (offset + 0u), value[0u]);
566 tint_symbol_8(buffer, (offset + 16u), value[1u]);
567 tint_symbol_8(buffer, (offset + 32u), value[2u]);
568 }
569
570 fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x4<f32>) {
571 tint_symbol_11(buffer, (offset + 0u), value[0u]);
572 tint_symbol_11(buffer, (offset + 16u), value[1u]);
573 tint_symbol_11(buffer, (offset + 32u), value[2u]);
574 }
575
576 fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x2<f32>) {
577 tint_symbol_5(buffer, (offset + 0u), value[0u]);
578 tint_symbol_5(buffer, (offset + 8u), value[1u]);
579 tint_symbol_5(buffer, (offset + 16u), value[2u]);
580 tint_symbol_5(buffer, (offset + 24u), value[3u]);
581 }
582
583 fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x3<f32>) {
584 tint_symbol_8(buffer, (offset + 0u), value[0u]);
585 tint_symbol_8(buffer, (offset + 16u), value[1u]);
586 tint_symbol_8(buffer, (offset + 32u), value[2u]);
587 tint_symbol_8(buffer, (offset + 48u), value[3u]);
588 }
589
590 fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x4<f32>) {
591 tint_symbol_11(buffer, (offset + 0u), value[0u]);
592 tint_symbol_11(buffer, (offset + 16u), value[1u]);
593 tint_symbol_11(buffer, (offset + 32u), value[2u]);
594 tint_symbol_11(buffer, (offset + 48u), value[3u]);
595 }
596
597 fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2u>) {
598 var array = value;
599 for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
600 tint_symbol_8(buffer, (offset + (i_1 * 16u)), array[i_1]);
601 }
602 }
603
604 [[stage(compute), workgroup_size(1)]]
605 fn main() {
606 tint_symbol(sb, 0u, i32());
607 tint_symbol_1(sb, 4u, u32());
608 tint_symbol_2(sb, 8u, f32());
609 tint_symbol_3(sb, 16u, vec2<i32>());
610 tint_symbol_4(sb, 24u, vec2<u32>());
611 tint_symbol_5(sb, 32u, vec2<f32>());
612 tint_symbol_6(sb, 48u, vec3<i32>());
613 tint_symbol_7(sb, 64u, vec3<u32>());
614 tint_symbol_8(sb, 80u, vec3<f32>());
615 tint_symbol_9(sb, 96u, vec4<i32>());
616 tint_symbol_10(sb, 112u, vec4<u32>());
617 tint_symbol_11(sb, 128u, vec4<f32>());
618 tint_symbol_12(sb, 144u, mat2x2<f32>());
619 tint_symbol_13(sb, 160u, mat2x3<f32>());
620 tint_symbol_14(sb, 192u, mat2x4<f32>());
621 tint_symbol_15(sb, 224u, mat3x2<f32>());
622 tint_symbol_16(sb, 256u, mat3x3<f32>());
623 tint_symbol_17(sb, 304u, mat3x4<f32>());
624 tint_symbol_18(sb, 352u, mat4x2<f32>());
625 tint_symbol_19(sb, 384u, mat4x3<f32>());
626 tint_symbol_20(sb, 448u, mat4x4<f32>());
627 tint_symbol_21(sb, 512u, array<vec3<f32>, 2>());
628 }
629 )";
630
631 auto got = Run<DecomposeMemoryAccess>(src);
632
633 EXPECT_EQ(expect, str(got));
634 }
635
TEST_F(DecomposeMemoryAccessTest,LoadStructure)636 TEST_F(DecomposeMemoryAccessTest, LoadStructure) {
637 auto* src = R"(
638 [[block]]
639 struct SB {
640 a : i32;
641 b : u32;
642 c : f32;
643 d : vec2<i32>;
644 e : vec2<u32>;
645 f : vec2<f32>;
646 g : vec3<i32>;
647 h : vec3<u32>;
648 i : vec3<f32>;
649 j : vec4<i32>;
650 k : vec4<u32>;
651 l : vec4<f32>;
652 m : mat2x2<f32>;
653 n : mat2x3<f32>;
654 o : mat2x4<f32>;
655 p : mat3x2<f32>;
656 q : mat3x3<f32>;
657 r : mat3x4<f32>;
658 s : mat4x2<f32>;
659 t : mat4x3<f32>;
660 u : mat4x4<f32>;
661 v : array<vec3<f32>, 2>;
662 };
663
664 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
665
666 [[stage(compute), workgroup_size(1)]]
667 fn main() {
668 var x : SB = sb;
669 }
670 )";
671
672 auto* expect = R"(
673 [[block]]
674 struct SB {
675 a : i32;
676 b : u32;
677 c : f32;
678 d : vec2<i32>;
679 e : vec2<u32>;
680 f : vec2<f32>;
681 g : vec3<i32>;
682 h : vec3<u32>;
683 i : vec3<f32>;
684 j : vec4<i32>;
685 k : vec4<u32>;
686 l : vec4<f32>;
687 m : mat2x2<f32>;
688 n : mat2x3<f32>;
689 o : mat2x4<f32>;
690 p : mat3x2<f32>;
691 q : mat3x3<f32>;
692 r : mat3x4<f32>;
693 s : mat4x2<f32>;
694 t : mat4x3<f32>;
695 u : mat4x4<f32>;
696 v : array<vec3<f32>, 2>;
697 };
698
699 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
700
701 [[internal(intrinsic_load_storage_i32), internal(disable_validation__function_has_no_body)]]
702 fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32
703
704 [[internal(intrinsic_load_storage_u32), internal(disable_validation__function_has_no_body)]]
705 fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
706
707 [[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]]
708 fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
709
710 [[internal(intrinsic_load_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
711 fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<i32>
712
713 [[internal(intrinsic_load_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
714 fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<u32>
715
716 [[internal(intrinsic_load_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
717 fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<f32>
718
719 [[internal(intrinsic_load_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
720 fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<i32>
721
722 [[internal(intrinsic_load_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
723 fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<u32>
724
725 [[internal(intrinsic_load_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
726 fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<f32>
727
728 [[internal(intrinsic_load_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
729 fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<i32>
730
731 [[internal(intrinsic_load_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
732 fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<u32>
733
734 [[internal(intrinsic_load_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
735 fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<f32>
736
737 fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x2<f32> {
738 return mat2x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)));
739 }
740
741 fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x3<f32> {
742 return mat2x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)));
743 }
744
745 fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x4<f32> {
746 return mat2x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)));
747 }
748
749 fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x2<f32> {
750 return mat3x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)));
751 }
752
753 fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x3<f32> {
754 return mat3x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)));
755 }
756
757 fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x4<f32> {
758 return mat3x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)));
759 }
760
761 fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x2<f32> {
762 return mat4x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)), tint_symbol_6(buffer, (offset + 24u)));
763 }
764
765 fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x3<f32> {
766 return mat4x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)), tint_symbol_9(buffer, (offset + 48u)));
767 }
768
769 fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x4<f32> {
770 return mat4x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)), tint_symbol_12(buffer, (offset + 48u)));
771 }
772
773 fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2u> {
774 var arr : array<vec3<f32>, 2u>;
775 for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
776 arr[i_1] = tint_symbol_9(buffer, (offset + (i_1 * 16u)));
777 }
778 return arr;
779 }
780
781 fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> SB {
782 return SB(tint_symbol_1(buffer, (offset + 0u)), tint_symbol_2(buffer, (offset + 4u)), tint_symbol_3(buffer, (offset + 8u)), tint_symbol_4(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)), tint_symbol_6(buffer, (offset + 32u)), tint_symbol_7(buffer, (offset + 48u)), tint_symbol_8(buffer, (offset + 64u)), tint_symbol_9(buffer, (offset + 80u)), tint_symbol_10(buffer, (offset + 96u)), tint_symbol_11(buffer, (offset + 112u)), tint_symbol_12(buffer, (offset + 128u)), tint_symbol_13(buffer, (offset + 144u)), tint_symbol_14(buffer, (offset + 160u)), tint_symbol_15(buffer, (offset + 192u)), tint_symbol_16(buffer, (offset + 224u)), tint_symbol_17(buffer, (offset + 256u)), tint_symbol_18(buffer, (offset + 304u)), tint_symbol_19(buffer, (offset + 352u)), tint_symbol_20(buffer, (offset + 384u)), tint_symbol_21(buffer, (offset + 448u)), tint_symbol_22(buffer, (offset + 512u)));
783 }
784
785 [[stage(compute), workgroup_size(1)]]
786 fn main() {
787 var x : SB = tint_symbol(sb, 0u);
788 }
789 )";
790
791 auto got = Run<DecomposeMemoryAccess>(src);
792
793 EXPECT_EQ(expect, str(got));
794 }
795
TEST_F(DecomposeMemoryAccessTest,StoreStructure)796 TEST_F(DecomposeMemoryAccessTest, StoreStructure) {
797 auto* src = R"(
798 [[block]]
799 struct SB {
800 a : i32;
801 b : u32;
802 c : f32;
803 d : vec2<i32>;
804 e : vec2<u32>;
805 f : vec2<f32>;
806 g : vec3<i32>;
807 h : vec3<u32>;
808 i : vec3<f32>;
809 j : vec4<i32>;
810 k : vec4<u32>;
811 l : vec4<f32>;
812 m : mat2x2<f32>;
813 n : mat2x3<f32>;
814 o : mat2x4<f32>;
815 p : mat3x2<f32>;
816 q : mat3x3<f32>;
817 r : mat3x4<f32>;
818 s : mat4x2<f32>;
819 t : mat4x3<f32>;
820 u : mat4x4<f32>;
821 v : array<vec3<f32>, 2>;
822 };
823
824 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
825
826 [[stage(compute), workgroup_size(1)]]
827 fn main() {
828 sb = SB();
829 }
830 )";
831
832 auto* expect = R"(
833 [[block]]
834 struct SB {
835 a : i32;
836 b : u32;
837 c : f32;
838 d : vec2<i32>;
839 e : vec2<u32>;
840 f : vec2<f32>;
841 g : vec3<i32>;
842 h : vec3<u32>;
843 i : vec3<f32>;
844 j : vec4<i32>;
845 k : vec4<u32>;
846 l : vec4<f32>;
847 m : mat2x2<f32>;
848 n : mat2x3<f32>;
849 o : mat2x4<f32>;
850 p : mat3x2<f32>;
851 q : mat3x3<f32>;
852 r : mat3x4<f32>;
853 s : mat4x2<f32>;
854 t : mat4x3<f32>;
855 u : mat4x4<f32>;
856 v : array<vec3<f32>, 2>;
857 };
858
859 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
860
861 [[internal(intrinsic_store_storage_i32), internal(disable_validation__function_has_no_body)]]
862 fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : i32)
863
864 [[internal(intrinsic_store_storage_u32), internal(disable_validation__function_has_no_body)]]
865 fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : u32)
866
867 [[internal(intrinsic_store_storage_f32), internal(disable_validation__function_has_no_body)]]
868 fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : f32)
869
870 [[internal(intrinsic_store_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
871 fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<i32>)
872
873 [[internal(intrinsic_store_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
874 fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<u32>)
875
876 [[internal(intrinsic_store_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
877 fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<f32>)
878
879 [[internal(intrinsic_store_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
880 fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<i32>)
881
882 [[internal(intrinsic_store_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
883 fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<u32>)
884
885 [[internal(intrinsic_store_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
886 fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<f32>)
887
888 [[internal(intrinsic_store_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
889 fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<i32>)
890
891 [[internal(intrinsic_store_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
892 fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<u32>)
893
894 [[internal(intrinsic_store_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
895 fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<f32>)
896
897 fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x2<f32>) {
898 tint_symbol_6(buffer, (offset + 0u), value[0u]);
899 tint_symbol_6(buffer, (offset + 8u), value[1u]);
900 }
901
902 fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x3<f32>) {
903 tint_symbol_9(buffer, (offset + 0u), value[0u]);
904 tint_symbol_9(buffer, (offset + 16u), value[1u]);
905 }
906
907 fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x4<f32>) {
908 tint_symbol_12(buffer, (offset + 0u), value[0u]);
909 tint_symbol_12(buffer, (offset + 16u), value[1u]);
910 }
911
912 fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x2<f32>) {
913 tint_symbol_6(buffer, (offset + 0u), value[0u]);
914 tint_symbol_6(buffer, (offset + 8u), value[1u]);
915 tint_symbol_6(buffer, (offset + 16u), value[2u]);
916 }
917
918 fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x3<f32>) {
919 tint_symbol_9(buffer, (offset + 0u), value[0u]);
920 tint_symbol_9(buffer, (offset + 16u), value[1u]);
921 tint_symbol_9(buffer, (offset + 32u), value[2u]);
922 }
923
924 fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x4<f32>) {
925 tint_symbol_12(buffer, (offset + 0u), value[0u]);
926 tint_symbol_12(buffer, (offset + 16u), value[1u]);
927 tint_symbol_12(buffer, (offset + 32u), value[2u]);
928 }
929
930 fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x2<f32>) {
931 tint_symbol_6(buffer, (offset + 0u), value[0u]);
932 tint_symbol_6(buffer, (offset + 8u), value[1u]);
933 tint_symbol_6(buffer, (offset + 16u), value[2u]);
934 tint_symbol_6(buffer, (offset + 24u), value[3u]);
935 }
936
937 fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x3<f32>) {
938 tint_symbol_9(buffer, (offset + 0u), value[0u]);
939 tint_symbol_9(buffer, (offset + 16u), value[1u]);
940 tint_symbol_9(buffer, (offset + 32u), value[2u]);
941 tint_symbol_9(buffer, (offset + 48u), value[3u]);
942 }
943
944 fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x4<f32>) {
945 tint_symbol_12(buffer, (offset + 0u), value[0u]);
946 tint_symbol_12(buffer, (offset + 16u), value[1u]);
947 tint_symbol_12(buffer, (offset + 32u), value[2u]);
948 tint_symbol_12(buffer, (offset + 48u), value[3u]);
949 }
950
951 fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2u>) {
952 var array = value;
953 for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
954 tint_symbol_9(buffer, (offset + (i_1 * 16u)), array[i_1]);
955 }
956 }
957
958 fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : SB) {
959 tint_symbol_1(buffer, (offset + 0u), value.a);
960 tint_symbol_2(buffer, (offset + 4u), value.b);
961 tint_symbol_3(buffer, (offset + 8u), value.c);
962 tint_symbol_4(buffer, (offset + 16u), value.d);
963 tint_symbol_5(buffer, (offset + 24u), value.e);
964 tint_symbol_6(buffer, (offset + 32u), value.f);
965 tint_symbol_7(buffer, (offset + 48u), value.g);
966 tint_symbol_8(buffer, (offset + 64u), value.h);
967 tint_symbol_9(buffer, (offset + 80u), value.i);
968 tint_symbol_10(buffer, (offset + 96u), value.j);
969 tint_symbol_11(buffer, (offset + 112u), value.k);
970 tint_symbol_12(buffer, (offset + 128u), value.l);
971 tint_symbol_13(buffer, (offset + 144u), value.m);
972 tint_symbol_14(buffer, (offset + 160u), value.n);
973 tint_symbol_15(buffer, (offset + 192u), value.o);
974 tint_symbol_16(buffer, (offset + 224u), value.p);
975 tint_symbol_17(buffer, (offset + 256u), value.q);
976 tint_symbol_18(buffer, (offset + 304u), value.r);
977 tint_symbol_19(buffer, (offset + 352u), value.s);
978 tint_symbol_20(buffer, (offset + 384u), value.t);
979 tint_symbol_21(buffer, (offset + 448u), value.u);
980 tint_symbol_22(buffer, (offset + 512u), value.v);
981 }
982
983 [[stage(compute), workgroup_size(1)]]
984 fn main() {
985 tint_symbol(sb, 0u, SB());
986 }
987 )";
988
989 auto got = Run<DecomposeMemoryAccess>(src);
990
991 EXPECT_EQ(expect, str(got));
992 }
993
TEST_F(DecomposeMemoryAccessTest,ComplexStaticAccessChain)994 TEST_F(DecomposeMemoryAccessTest, ComplexStaticAccessChain) {
995 auto* src = R"(
996 struct S1 {
997 a : i32;
998 b : vec3<f32>;
999 c : i32;
1000 };
1001
1002 struct S2 {
1003 a : i32;
1004 b : [[stride(32)]] array<S1, 3>;
1005 c : i32;
1006 };
1007
1008 [[block]]
1009 struct SB {
1010 [[size(128)]]
1011 a : i32;
1012 b : [[stride(256)]] array<S2>;
1013 };
1014
1015 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
1016
1017 [[stage(compute), workgroup_size(1)]]
1018 fn main() {
1019 var x : f32 = sb.b[4].b[1].b.z;
1020 }
1021 )";
1022
1023 // sb.b[4].b[1].b.z
1024 // ^ ^ ^ ^ ^ ^
1025 // | | | | | |
1026 // 128 | |1200| 1224
1027 // | | |
1028 // 1152 1168 1216
1029
1030 auto* expect = R"(
1031 struct S1 {
1032 a : i32;
1033 b : vec3<f32>;
1034 c : i32;
1035 };
1036
1037 struct S2 {
1038 a : i32;
1039 b : [[stride(32)]] array<S1, 3>;
1040 c : i32;
1041 };
1042
1043 [[block]]
1044 struct SB {
1045 [[size(128)]]
1046 a : i32;
1047 b : [[stride(256)]] array<S2>;
1048 };
1049
1050 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
1051
1052 [[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]]
1053 fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
1054
1055 [[stage(compute), workgroup_size(1)]]
1056 fn main() {
1057 var x : f32 = tint_symbol(sb, 1224u);
1058 }
1059 )";
1060
1061 auto got = Run<DecomposeMemoryAccess>(src);
1062
1063 EXPECT_EQ(expect, str(got));
1064 }
1065
TEST_F(DecomposeMemoryAccessTest,ComplexDynamicAccessChain)1066 TEST_F(DecomposeMemoryAccessTest, ComplexDynamicAccessChain) {
1067 auto* src = R"(
1068 struct S1 {
1069 a : i32;
1070 b : vec3<f32>;
1071 c : i32;
1072 };
1073
1074 struct S2 {
1075 a : i32;
1076 b : [[stride(32)]] array<S1, 3>;
1077 c : i32;
1078 };
1079
1080 [[block]]
1081 struct SB {
1082 [[size(128)]]
1083 a : i32;
1084 b : [[stride(256)]] array<S2>;
1085 };
1086
1087 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
1088
1089 [[stage(compute), workgroup_size(1)]]
1090 fn main() {
1091 var i : i32 = 4;
1092 var j : u32 = 1u;
1093 var k : i32 = 2;
1094 var x : f32 = sb.b[i].b[j].b[k];
1095 }
1096 )";
1097
1098 auto* expect = R"(
1099 struct S1 {
1100 a : i32;
1101 b : vec3<f32>;
1102 c : i32;
1103 };
1104
1105 struct S2 {
1106 a : i32;
1107 b : [[stride(32)]] array<S1, 3>;
1108 c : i32;
1109 };
1110
1111 [[block]]
1112 struct SB {
1113 [[size(128)]]
1114 a : i32;
1115 b : [[stride(256)]] array<S2>;
1116 };
1117
1118 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
1119
1120 [[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]]
1121 fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
1122
1123 [[stage(compute), workgroup_size(1)]]
1124 fn main() {
1125 var i : i32 = 4;
1126 var j : u32 = 1u;
1127 var k : i32 = 2;
1128 var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
1129 }
1130 )";
1131
1132 auto got = Run<DecomposeMemoryAccess>(src);
1133
1134 EXPECT_EQ(expect, str(got));
1135 }
1136
TEST_F(DecomposeMemoryAccessTest,ComplexDynamicAccessChainWithAliases)1137 TEST_F(DecomposeMemoryAccessTest, ComplexDynamicAccessChainWithAliases) {
1138 auto* src = R"(
1139 struct S1 {
1140 a : i32;
1141 b : vec3<f32>;
1142 c : i32;
1143 };
1144
1145 type A1 = S1;
1146
1147 type A1_Array = [[stride(32)]] array<S1, 3>;
1148
1149 struct S2 {
1150 a : i32;
1151 b : A1_Array;
1152 c : i32;
1153 };
1154
1155 type A2 = S2;
1156
1157 type A2_Array = [[stride(256)]] array<S2>;
1158
1159 [[block]]
1160 struct SB {
1161 [[size(128)]]
1162 a : i32;
1163 b : A2_Array;
1164 };
1165
1166 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
1167
1168 [[stage(compute), workgroup_size(1)]]
1169 fn main() {
1170 var i : i32 = 4;
1171 var j : u32 = 1u;
1172 var k : i32 = 2;
1173 var x : f32 = sb.b[i].b[j].b[k];
1174 }
1175 )";
1176
1177 auto* expect = R"(
1178 struct S1 {
1179 a : i32;
1180 b : vec3<f32>;
1181 c : i32;
1182 };
1183
1184 type A1 = S1;
1185
1186 type A1_Array = [[stride(32)]] array<S1, 3>;
1187
1188 struct S2 {
1189 a : i32;
1190 b : A1_Array;
1191 c : i32;
1192 };
1193
1194 type A2 = S2;
1195
1196 type A2_Array = [[stride(256)]] array<S2>;
1197
1198 [[block]]
1199 struct SB {
1200 [[size(128)]]
1201 a : i32;
1202 b : A2_Array;
1203 };
1204
1205 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
1206
1207 [[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]]
1208 fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
1209
1210 [[stage(compute), workgroup_size(1)]]
1211 fn main() {
1212 var i : i32 = 4;
1213 var j : u32 = 1u;
1214 var k : i32 = 2;
1215 var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
1216 }
1217 )";
1218
1219 auto got = Run<DecomposeMemoryAccess>(src);
1220
1221 EXPECT_EQ(expect, str(got));
1222 }
1223
TEST_F(DecomposeMemoryAccessTest,StorageBufferAtomics)1224 TEST_F(DecomposeMemoryAccessTest, StorageBufferAtomics) {
1225 auto* src = R"(
1226 [[block]]
1227 struct SB {
1228 padding : vec4<f32>;
1229 a : atomic<i32>;
1230 b : atomic<u32>;
1231 };
1232
1233 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
1234
1235 [[stage(compute), workgroup_size(1)]]
1236 fn main() {
1237 atomicStore(&sb.a, 123);
1238 atomicLoad(&sb.a);
1239 atomicAdd(&sb.a, 123);
1240 atomicSub(&sb.a, 123);
1241 atomicMax(&sb.a, 123);
1242 atomicMin(&sb.a, 123);
1243 atomicAnd(&sb.a, 123);
1244 atomicOr(&sb.a, 123);
1245 atomicXor(&sb.a, 123);
1246 atomicExchange(&sb.a, 123);
1247 atomicCompareExchangeWeak(&sb.a, 123, 345);
1248
1249 atomicStore(&sb.b, 123u);
1250 atomicLoad(&sb.b);
1251 atomicAdd(&sb.b, 123u);
1252 atomicSub(&sb.b, 123u);
1253 atomicMax(&sb.b, 123u);
1254 atomicMin(&sb.b, 123u);
1255 atomicAnd(&sb.b, 123u);
1256 atomicOr(&sb.b, 123u);
1257 atomicXor(&sb.b, 123u);
1258 atomicExchange(&sb.b, 123u);
1259 atomicCompareExchangeWeak(&sb.b, 123u, 345u);
1260 }
1261 )";
1262
1263 auto* expect = R"(
1264 [[block]]
1265 struct SB {
1266 padding : vec4<f32>;
1267 a : atomic<i32>;
1268 b : atomic<u32>;
1269 };
1270
1271 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
1272
1273 [[internal(intrinsic_atomic_store_storage_i32), internal(disable_validation__function_has_no_body)]]
1274 fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32)
1275
1276 [[internal(intrinsic_atomic_load_storage_i32), internal(disable_validation__function_has_no_body)]]
1277 fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32
1278
1279 [[internal(intrinsic_atomic_add_storage_i32), internal(disable_validation__function_has_no_body)]]
1280 fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
1281
1282 [[internal(intrinsic_atomic_sub_storage_i32), internal(disable_validation__function_has_no_body)]]
1283 fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
1284
1285 [[internal(intrinsic_atomic_max_storage_i32), internal(disable_validation__function_has_no_body)]]
1286 fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
1287
1288 [[internal(intrinsic_atomic_min_storage_i32), internal(disable_validation__function_has_no_body)]]
1289 fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
1290
1291 [[internal(intrinsic_atomic_and_storage_i32), internal(disable_validation__function_has_no_body)]]
1292 fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
1293
1294 [[internal(intrinsic_atomic_or_storage_i32), internal(disable_validation__function_has_no_body)]]
1295 fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
1296
1297 [[internal(intrinsic_atomic_xor_storage_i32), internal(disable_validation__function_has_no_body)]]
1298 fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
1299
1300 [[internal(intrinsic_atomic_exchange_storage_i32), internal(disable_validation__function_has_no_body)]]
1301 fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
1302
1303 [[internal(intrinsic_atomic_compare_exchange_weak_storage_i32), internal(disable_validation__function_has_no_body)]]
1304 fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> vec2<i32>
1305
1306 [[internal(intrinsic_atomic_store_storage_u32), internal(disable_validation__function_has_no_body)]]
1307 fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32)
1308
1309 [[internal(intrinsic_atomic_load_storage_u32), internal(disable_validation__function_has_no_body)]]
1310 fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
1311
1312 [[internal(intrinsic_atomic_add_storage_u32), internal(disable_validation__function_has_no_body)]]
1313 fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
1314
1315 [[internal(intrinsic_atomic_sub_storage_u32), internal(disable_validation__function_has_no_body)]]
1316 fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
1317
1318 [[internal(intrinsic_atomic_max_storage_u32), internal(disable_validation__function_has_no_body)]]
1319 fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
1320
1321 [[internal(intrinsic_atomic_min_storage_u32), internal(disable_validation__function_has_no_body)]]
1322 fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
1323
1324 [[internal(intrinsic_atomic_and_storage_u32), internal(disable_validation__function_has_no_body)]]
1325 fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
1326
1327 [[internal(intrinsic_atomic_or_storage_u32), internal(disable_validation__function_has_no_body)]]
1328 fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
1329
1330 [[internal(intrinsic_atomic_xor_storage_u32), internal(disable_validation__function_has_no_body)]]
1331 fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
1332
1333 [[internal(intrinsic_atomic_exchange_storage_u32), internal(disable_validation__function_has_no_body)]]
1334 fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
1335
1336 [[internal(intrinsic_atomic_compare_exchange_weak_storage_u32), internal(disable_validation__function_has_no_body)]]
1337 fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> vec2<u32>
1338
1339 [[stage(compute), workgroup_size(1)]]
1340 fn main() {
1341 tint_symbol(sb, 16u, 123);
1342 tint_symbol_1(sb, 16u);
1343 tint_symbol_2(sb, 16u, 123);
1344 tint_symbol_3(sb, 16u, 123);
1345 tint_symbol_4(sb, 16u, 123);
1346 tint_symbol_5(sb, 16u, 123);
1347 tint_symbol_6(sb, 16u, 123);
1348 tint_symbol_7(sb, 16u, 123);
1349 tint_symbol_8(sb, 16u, 123);
1350 tint_symbol_9(sb, 16u, 123);
1351 tint_symbol_10(sb, 16u, 123, 345);
1352 tint_symbol_11(sb, 20u, 123u);
1353 tint_symbol_12(sb, 20u);
1354 tint_symbol_13(sb, 20u, 123u);
1355 tint_symbol_14(sb, 20u, 123u);
1356 tint_symbol_15(sb, 20u, 123u);
1357 tint_symbol_16(sb, 20u, 123u);
1358 tint_symbol_17(sb, 20u, 123u);
1359 tint_symbol_18(sb, 20u, 123u);
1360 tint_symbol_19(sb, 20u, 123u);
1361 tint_symbol_20(sb, 20u, 123u);
1362 tint_symbol_21(sb, 20u, 123u, 345u);
1363 }
1364 )";
1365
1366 auto got = Run<DecomposeMemoryAccess>(src);
1367
1368 EXPECT_EQ(expect, str(got));
1369 }
1370
TEST_F(DecomposeMemoryAccessTest,WorkgroupBufferAtomics)1371 TEST_F(DecomposeMemoryAccessTest, WorkgroupBufferAtomics) {
1372 auto* src = R"(
1373 struct S {
1374 padding : vec4<f32>;
1375 a : atomic<i32>;
1376 b : atomic<u32>;
1377 };
1378
1379 var<workgroup> w : S;
1380
1381 [[stage(compute), workgroup_size(1)]]
1382 fn main() {
1383 atomicStore(&(w.a), 123);
1384 atomicLoad(&(w.a));
1385 atomicAdd(&(w.a), 123);
1386 atomicSub(&(w.a), 123);
1387 atomicMax(&(w.a), 123);
1388 atomicMin(&(w.a), 123);
1389 atomicAnd(&(w.a), 123);
1390 atomicOr(&(w.a), 123);
1391 atomicXor(&(w.a), 123);
1392 atomicExchange(&(w.a), 123);
1393 atomicCompareExchangeWeak(&(w.a), 123, 345);
1394 atomicStore(&(w.b), 123u);
1395 atomicLoad(&(w.b));
1396 atomicAdd(&(w.b), 123u);
1397 atomicSub(&(w.b), 123u);
1398 atomicMax(&(w.b), 123u);
1399 atomicMin(&(w.b), 123u);
1400 atomicAnd(&(w.b), 123u);
1401 atomicOr(&(w.b), 123u);
1402 atomicXor(&(w.b), 123u);
1403 atomicExchange(&(w.b), 123u);
1404 atomicCompareExchangeWeak(&(w.b), 123u, 345u);
1405 }
1406 )";
1407
1408 auto* expect = src;
1409
1410 auto got = Run<DecomposeMemoryAccess>(src);
1411
1412 EXPECT_EQ(expect, str(got));
1413 }
1414
1415 } // namespace
1416 } // namespace transform
1417 } // namespace tint
1418