1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
|
// RUN: %metal-compile main
// RUN: %metal main 2>&1 | %check
struct Unpacked {
x: i32,
}
struct T {
v2f: vec2<f32>,
v3f: vec3<f32>,
v4f: vec4<f32>,
v2u: vec2<u32>,
v3u: vec3<u32>,
v4u: vec4<u32>,
f: f32,
u: u32,
unpacked: Unpacked,
}
struct U {
// CHECK: array<type\d::PackedType, 1> field\d
ts: array<T>,
}
var<private> t: T;
var<private> m2: mat2x2<f32>;
var<private> m3: mat3x3<f32>;
@group(0) @binding(0) var<storage, read_write> t1: T;
@group(0) @binding(1) var<storage, read_write> t2: T;
@group(0) @binding(2) var<storage, read_write> v1: vec3<f32>;
@group(0) @binding(3) var<storage, read_write> v2: vec3<f32>;
@group(0) @binding(4) var<storage, read_write> at1: array<T, 2>;
@group(0) @binding(5) var<storage, read_write> at2: array<T, 2>;
@group(0) @binding(6) var<storage, read_write> av1: array<vec3f, 2>;
@group(0) @binding(7) var<storage, read_write> av2: array<vec3f, 2>;
@group(0) @binding(8) var<storage, read_write> u1: U;
fn testUnpacked() -> i32
{
{ let x = t.v3f * m3; }
{ let x = m3 * t.v3f; }
{ let x = t1.v2f * m2; }
{ let x = m2 * t1.v2f; }
{ let x = u1.ts[0].v3f * m3; }
{ let x = m3 * u1.ts[0].v3f; }
{ let x = av1[0] * m3; }
{ let x = m3 * av1[0]; }
return 0;
}
fn testAssignment() -> i32
{
// packed struct
// CHECK: local\d+ = __unpack\(global\d+\);
var t = t1;
// CHECK: local\d+ = __unpack\(global\d+\);
t = t1;
// CHECK-NEXT: global\d+ = global\d+;
t1 = t2;
// CHECK-NEXT: global\d+ = __pack\(local\d+\);
t2 = t;
// array of packed structs
// CHECK-NEXT: local\d+ = __unpack\(global\d+\);
var at = at1;
// CHECK-NEXT: local\d+ = __unpack\(global\d+\);
at = at1;
// CHECK-NEXT: global\d+ = global\d+;
at1 = at2;
// CHECK-NEXT: global\d+ = __pack\(local\d+\);
at2 = at;
// array of vec3
// CHECK-NEXT: local\d+ = __unpack\(global\d+\);
var av = av1;
// CHECK-NEXT: local\d+ = __unpack\(global\d+\);
av = av1;
// CHECK-NEXT: global\d+ = global\d+;
av1 = av2;
// CHECK-NEXT: global\d+ = __pack\(local\d+\);
av2 = av;
return 0;
}
fn testFieldAccess() -> i32
{
// CHECK: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = __unpack\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = __unpack\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
t.v2f.x = t1.v2f.x;
t.v3f.x = t1.v3f.x;
t.v4f.x = t1.v4f.x;
t.v2u.x = t1.v2u.x;
t.v3u.x = t1.v3u.x;
t.v4u.x = t1.v4u.x;
t.v2f = t1.v2f;
t.v3f = t1.v3f;
t.v4f = t1.v4f;
t.v2u = t1.v2u;
t.v3u = t1.v3u;
t.v4u = t1.v4u;
t.f = t1.f;
t.u = t1.u;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
t1.v2f.x = t2.v2f.x;
t1.v3f.x = t2.v3f.x;
t1.v4f.x = t2.v4f.x;
t1.v2u.x = t2.v2u.x;
t1.v3u.x = t2.v3u.x;
t1.v4u.x = t2.v4u.x;
t1.v2f = t2.v2f;
t1.v3f = t2.v3f;
t1.v4f = t2.v4f;
t1.v2u = t2.v2u;
t1.v3u = t2.v3u;
t1.v4u = t2.v4u;
t1.f = t2.f;
t1.u = t2.u;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d\.x = global\d+\.field\d\.x;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d;
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d
// CHECK-NEXT: global\d+\.field\d = global\d+\.field\d
t2.v2f.x = t.v2f.x;
t2.v3f.x = t.v3f.x;
t2.v4f.x = t.v4f.x;
t2.v2u.x = t.v2u.x;
t2.v3u.x = t.v3u.x;
t2.v4u.x = t.v4u.x;
t2.v2f = t.v2f;
t2.v3f = t.v3f;
t2.v4f = t.v4f;
t2.v2u = t.v2u;
t2.v3u = t.v3u;
t2.v4u = t.v4u;
t2.f = t.f;
t2.u = t.u;
return 0;
}
@group(0) @binding(9) var<storage, read_write> index: u32;
fn testIndexAccess() -> i32
{
// CHECK: local\d+ = __unpack\(global\d+\);
// CHECK-NEXT: local\d+\[__wgslMin\(unsigned\(0\), \(2u - 1u\)\)\] = __unpack\(global\d+\[__wgslMin\(unsigned\(0\), \(2u - 1u\)\)\]\);
// CHECK-NEXT: global\d+\[__wgslMin\(unsigned\(0\), \(2u - 1u\)\)\] = global\d+\[__wgslMin\(unsigned\(0\), \(2u - 1u\)\)\];
// CHECK-NEXT: global\d+\[__wgslMin\(unsigned\(0\), \(2u - 1u\)\)\] = __pack\(local\d+\[__wgslMin\(unsigned\(0\), \(2u - 1u\)\)\]\);
// CHECK-NEXT: global\d+\[__wgslMin\(unsigned\(global\d+\), \(2u - 1u\)\)\] = __pack\(local\d+\[__wgslMin\(unsigned\(global\d+\), \(2u - 1u\)\)\]\);
var at = at1;
at[0] = at1[0];
at1[0] = at2[0];
at2[0] = at[0];
at2[index] = at[index];
return 0;
}
fn testBinaryOperations() -> i32
{
// CHECK: global\d+\.field\d\.x = \(2. \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2. \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2. \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2u \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2u \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2u \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* __unpack\(global\d+\.field\d\)\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* __unpack\(global\d+\.field\d\)\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* global\d+\.field\d\);
t.v2f.x = 2 * t1.v2f.x;
t.v3f.x = 2 * t1.v3f.x;
t.v4f.x = 2 * t1.v4f.x;
t.v2u.x = 2 * t1.v2u.x;
t.v3u.x = 2 * t1.v3u.x;
t.v4u.x = 2 * t1.v4u.x;
t.v2f = 2 * t1.v2f;
t.v3f = 2 * t1.v3f;
t.v4f = 2 * t1.v4f;
t.v2u = 2 * t1.v2u;
t.v3u = 2 * t1.v3u;
t.v4u = 2 * t1.v4u;
t.f = 2 * t1.f;
t.u = 2 * t1.u;
// CHECK-NEXT: global\d+\.field\d\.x = \(2. \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2. \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2. \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2u \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2u \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2u \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* __unpack\(global\d+\.field\d\)\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* __unpack\(global\d+\.field\d\)\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* global\d+\.field\d\);
t1.v2f.x = 2 * t2.v2f.x;
t1.v3f.x = 2 * t2.v3f.x;
t1.v4f.x = 2 * t2.v4f.x;
t1.v2u.x = 2 * t2.v2u.x;
t1.v3u.x = 2 * t2.v3u.x;
t1.v4u.x = 2 * t2.v4u.x;
t1.v2f = 2 * t2.v2f;
t1.v3f = 2 * t2.v3f;
t1.v4f = 2 * t2.v4f;
t1.v2u = 2 * t2.v2u;
t1.v3u = 2 * t2.v3u;
t1.v4u = 2 * t2.v4u;
t1.f = 2 * t2.f;
t1.u = 2 * t2.u;
// CHECK-NEXT: global\d+\.field\d\.x = \(2. \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2. \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2. \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2u \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2u \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(2u \* global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2. \* global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(2u \* global\d+\.field\d\);
t2.v2f.x = 2 * t.v2f.x;
t2.v3f.x = 2 * t.v3f.x;
t2.v4f.x = 2 * t.v4f.x;
t2.v2u.x = 2 * t.v2u.x;
t2.v3u.x = 2 * t.v3u.x;
t2.v4u.x = 2 * t.v4u.x;
t2.v2f = 2 * t.v2f;
t2.v3f = 2 * t.v3f;
t2.v4f = 2 * t.v4f;
t2.v2u = 2 * t.v2u;
t2.v3u = 2 * t.v3u;
t2.v4u = 2 * t.v4u;
t2.f = 2 * t.f;
t2.u = 2 * t.u;
return 0;
}
fn testUnaryOperations() -> i32
{
// CHECK: global\d+\.field\d\.x = \(-global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(-global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(-global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d = \(-global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(-__unpack\(global\d+\.field\d\)\);
// CHECK-NEXT: global\d+\.field\d = \(-global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(-global\d+\.field\d\);
t.v2f.x = -t1.v2f.x;
t.v3f.x = -t1.v3f.x;
t.v4f.x = -t1.v4f.x;
t.v2f = -t1.v2f;
t.v3f = -t1.v3f;
t.v4f = -t1.v4f;
t.f = -t1.f;
// CHECK-NEXT: global\d+\.field\d\.x = \(-global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(-global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(-global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d = \(-global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(-__unpack\(global\d+\.field\d\)\);
// CHECK-NEXT: global\d+\.field\d = \(-global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(-global\d+\.field\d\);
t1.v2f.x = -t2.v2f.x;
t1.v3f.x = -t2.v3f.x;
t1.v4f.x = -t2.v4f.x;
t1.v2f = -t2.v2f;
t1.v3f = -t2.v3f;
t1.v4f = -t2.v4f;
t1.f = -t2.f;
// CHECK-NEXT: global\d+\.field\d\.x = \(-global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(-global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = \(-global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d = \(-global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(-global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(-global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = \(-global\d+\.field\d\);
t2.v2f.x = -t.v2f.x;
t2.v3f.x = -t.v3f.x;
t2.v4f.x = -t.v4f.x;
t2.v2f = -t.v2f;
t2.v3f = -t.v3f;
t2.v4f = -t.v4f;
t2.f = -t.f;
return 0;
}
fn testCall() -> i32
{
// CHECK: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(__unpack\(global\d+\.field\d\)\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(__unpack\(global\d+\.field\d\)\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
t.v2f.x = abs(t1.v2f.x);
t.v3f.x = abs(t1.v3f.x);
t.v4f.x = abs(t1.v4f.x);
t.v2u.x = abs(t1.v2u.x);
t.v3u.x = abs(t1.v3u.x);
t.v4u.x = abs(t1.v4u.x);
t.v2f = abs(t1.v2f);
t.v3f = abs(t1.v3f);
t.v4f = abs(t1.v4f);
t.v2u = abs(t1.v2u);
t.v3u = abs(t1.v3u);
t.v4u = abs(t1.v4u);
t.f = abs(t1.f);
t.u = abs(t1.u);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(__unpack\(global\d+\.field\d\)\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(__unpack\(global\d+\.field\d\)\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
t1.v2f.x = abs(t2.v2f.x);
t1.v3f.x = abs(t2.v3f.x);
t1.v4f.x = abs(t2.v4f.x);
t1.v2u.x = abs(t2.v2u.x);
t1.v3u.x = abs(t2.v3u.x);
t1.v4u.x = abs(t2.v4u.x);
t1.v2f = abs(t2.v2f);
t1.v3f = abs(t2.v3f);
t1.v4f = abs(t2.v4f);
t1.v2u = abs(t2.v2u);
t1.v3u = abs(t2.v3u);
t1.v4u = abs(t2.v4u);
t1.f = abs(t2.f);
t1.u = abs(t2.u);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d\.x = abs\(global\d+\.field\d\.x\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
// CHECK-NEXT: global\d+\.field\d = abs\(global\d+\.field\d\);
t2.v2f.x = abs(t.v2f.x);
t2.v3f.x = abs(t.v3f.x);
t2.v4f.x = abs(t.v4f.x);
t2.v2u.x = abs(t.v2u.x);
t2.v3u.x = abs(t.v3u.x);
t2.v4u.x = abs(t.v4u.x);
t2.v2f = abs(t.v2f);
t2.v3f = abs(t.v3f);
t2.v4f = abs(t.v4f);
t2.v2u = abs(t.v2u);
t2.v3u = abs(t.v3u);
t2.v4u = abs(t.v4u);
t2.f = abs(t.f);
t2.u = abs(t.u);
return 0;
}
struct S {
@size(16) x: i32,
@align(32) y: array<i32>,
}
@group(0) @binding(10) var<storage, read_write> s: S;
fn testRuntimeArray()
{
s.x = 0;
s.y[0] = 0;
}
@group(0) @binding(11) var<storage,read_write> D: array<vec3f>;
fn testPackedVec3CompoundAssignment()
{
D[0] += vec3f(1);
}
struct InconstructibleS {
x: vec3f,
y: atomic<u32>
}
struct InconstructibleT {
x: InconstructibleS,
}
@group(0) @binding(12) var<storage, read_write> global: InconstructibleT;
fn testInconstructible() {
var x = global.x.x;
}
@compute @workgroup_size(1)
fn main()
{
_ = testUnpacked();
_ = testAssignment();
_ = testFieldAccess();
_ = testIndexAccess();
_ = testBinaryOperations();
_ = testUnaryOperations();
_ = testCall();
testRuntimeArray();
testPackedVec3CompoundAssignment();
testInconstructible();
}
|
