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
|
// RUN: %metal-compile main
// RUN: %metal main 2>&1 | %check
struct T {
v2f: vec2<f32>,
v3f: vec3<f32>,
v4f: vec4<f32>,
v2u: vec2<u32>,
v3u: vec3<u32>,
v4u: vec4<u32>,
f: f32,
u: u32,
}
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, 1>;
@group(0) @binding(5) var<storage, read_write> at2: array<T, 1>;
fn testUnpacked() -> i32
{
_ = t.v3f * m3;
_ = m3 * t.v3f;
_ = t1.v2f * m2;
_ = m2 * t1.v2f;
return 0;
}
fn testAssignment() -> i32
{
// packed struct
// CHECK: local\d+ = __unpack\(parameter\d+\);
var t = t1;
// CHECK-NEXT: parameter\d+ = parameter\d+;
t1 = t2;
// CHECK-NEXT: parameter\d+ = __pack\(local\d+\);
t2 = t;
// array of packed structs
// CHECK-NEXT: local\d+ = __unpack_array\(parameter\d+\);
var at = at1;
// CHECK-NEXT: parameter\d+ = parameter\d+;
at1 = at2;
// CHECK-NEXT: parameter\d+ = __pack_array\(local\d+\);
at2 = at;
return 0;
}
fn testFieldAccess() -> i32
{
// CHECK: parameter\d+\.v2f\.x = parameter\d+\.v2f\.x;
// CHECK-NEXT: parameter\d+\.v3f\.x = parameter\d+\.v3f\.x;
// CHECK-NEXT: parameter\d+\.v4f\.x = parameter\d+\.v4f\.x;
// CHECK-NEXT: parameter\d+\.v2u\.x = parameter\d+\.v2u\.x;
// CHECK-NEXT: parameter\d+\.v3u\.x = parameter\d+\.v3u\.x;
// CHECK-NEXT: parameter\d+\.v4u\.x = parameter\d+\.v4u\.x;
// CHECK-NEXT: parameter\d+\.v2f = parameter\d+\.v2f;
// CHECK-NEXT: parameter\d+\.v3f = float3\(parameter\d+\.v3f\);
// CHECK-NEXT: parameter\d+\.v4f = parameter\d+\.v4f;
// CHECK-NEXT: parameter\d+\.v2u = parameter\d+\.v2u;
// CHECK-NEXT: parameter\d+\.v3u = uint3\(parameter\d+\.v3u\);
// CHECK-NEXT: parameter\d+\.v4u = parameter\d+\.v4u;
// CHECK-NEXT: parameter\d+\.f = parameter\d+\.f;
// CHECK-NEXT: parameter\d+\.u = parameter\d+\.u;
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: parameter\d+\.v2f\.x = parameter\d+\.v2f\.x;
// CHECK-NEXT: parameter\d+\.v3f\.x = parameter\d+\.v3f\.x;
// CHECK-NEXT: parameter\d+\.v4f\.x = parameter\d+\.v4f\.x;
// CHECK-NEXT: parameter\d+\.v2u\.x = parameter\d+\.v2u\.x;
// CHECK-NEXT: parameter\d+\.v3u\.x = parameter\d+\.v3u\.x;
// CHECK-NEXT: parameter\d+\.v4u\.x = parameter\d+\.v4u\.x;
// CHECK-NEXT: parameter\d+\.v2f = parameter\d+\.v2f;
// CHECK-NEXT: parameter\d+\.v3f = parameter\d+\.v3f;
// CHECK-NEXT: parameter\d+\.v4f = parameter\d+\.v4f;
// CHECK-NEXT: parameter\d+\.v2u = parameter\d+\.v2u;
// CHECK-NEXT: parameter\d+\.v3u = parameter\d+\.v3u;
// CHECK-NEXT: parameter\d+\.v4u = parameter\d+\.v4u;
// CHECK-NEXT: parameter\d+\.f = parameter\d+\.f;
// CHECK-NEXT: parameter\d+\.u = parameter\d+\.u;
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: parameter\d+\.v2f\.x = parameter\d+\.v2f\.x;
// CHECK-NEXT: parameter\d+\.v3f\.x = parameter\d+\.v3f\.x;
// CHECK-NEXT: parameter\d+\.v4f\.x = parameter\d+\.v4f\.x;
// CHECK-NEXT: parameter\d+\.v2u\.x = parameter\d+\.v2u\.x;
// CHECK-NEXT: parameter\d+\.v3u\.x = parameter\d+\.v3u\.x;
// CHECK-NEXT: parameter\d+\.v4u\.x = parameter\d+\.v4u\.x;
// CHECK-NEXT: parameter\d+\.v2f = parameter\d+\.v2f;
// CHECK-NEXT: parameter\d+\.v3f = packed_float3\(parameter\d+\.v3f\);
// CHECK-NEXT: parameter\d+\.v4f = parameter\d+\.v4f;
// CHECK-NEXT: parameter\d+\.v2u = parameter\d+\.v2u;
// CHECK-NEXT: parameter\d+\.v3u = packed_uint3\(parameter\d+\.v3u\);
// CHECK-NEXT: parameter\d+\.v4u = parameter\d+\.v4u;
// CHECK-NEXT: parameter\d+\.f = parameter\d+\.f;
// CHECK-NEXT: parameter\d+\.u = parameter\d+\.u;
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;
}
fn testIndexAccess() -> i32
{
// CHECK: local\d+ = __unpack_array\(parameter\d+\);
// CHECK-NEXT: local\d+\[0\] = __unpack\(parameter\d+\[0\]\);
// CHECK-NEXT: parameter\d+\[0\] = parameter\d+\[0\];
// CHECK-NEXT: parameter\d+\[0\] = __pack\(local\d+\[0\]\);
var at = at1;
at[0] = at1[0];
at1[0] = at2[0];
at2[0] = at[0];
return 0;
}
fn testBinaryOperations() -> i32
{
// CHECK: parameter\d+\.v2f\.x = \(2 \* parameter\d+\.v2f\.x\);
// CHECK-NEXT: parameter\d+\.v3f\.x = \(2 \* parameter\d+\.v3f\.x\);
// CHECK-NEXT: parameter\d+\.v4f\.x = \(2 \* parameter\d+\.v4f\.x\);
// CHECK-NEXT: parameter\d+\.v2u\.x = \(2u \* parameter\d+\.v2u\.x\);
// CHECK-NEXT: parameter\d+\.v3u\.x = \(2u \* parameter\d+\.v3u\.x\);
// CHECK-NEXT: parameter\d+\.v4u\.x = \(2u \* parameter\d+\.v4u\.x\);
// CHECK-NEXT: parameter\d+\.v2f = \(2 \* parameter\d+\.v2f\);
// CHECK-NEXT: parameter\d+\.v3f = \(2 \* float3\(parameter\d+\.v3f\)\);
// CHECK-NEXT: parameter\d+\.v4f = \(2 \* parameter\d+\.v4f\);
// CHECK-NEXT: parameter\d+\.v2u = \(2u \* parameter\d+\.v2u\);
// CHECK-NEXT: parameter\d+\.v3u = \(2u \* uint3\(parameter\d+\.v3u\)\);
// CHECK-NEXT: parameter\d+\.v4u = \(2u \* parameter\d+\.v4u\);
// CHECK-NEXT: parameter\d+\.f = \(2 \* parameter\d+\.f\);
// CHECK-NEXT: parameter\d+\.u = \(2u \* parameter\d+\.u\);
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: parameter\d+\.v2f\.x = \(2 \* parameter\d+\.v2f\.x\);
// CHECK-NEXT: parameter\d+\.v3f\.x = \(2 \* parameter\d+\.v3f\.x\);
// CHECK-NEXT: parameter\d+\.v4f\.x = \(2 \* parameter\d+\.v4f\.x\);
// CHECK-NEXT: parameter\d+\.v2u\.x = \(2u \* parameter\d+\.v2u\.x\);
// CHECK-NEXT: parameter\d+\.v3u\.x = \(2u \* parameter\d+\.v3u\.x\);
// CHECK-NEXT: parameter\d+\.v4u\.x = \(2u \* parameter\d+\.v4u\.x\);
// CHECK-NEXT: parameter\d+\.v2f = \(2 \* parameter\d+\.v2f\);
// CHECK-NEXT: parameter\d+\.v3f = packed_float3\(\(2 \* float3\(parameter\d+\.v3f\)\)\);
// CHECK-NEXT: parameter\d+\.v4f = \(2 \* parameter\d+\.v4f\);
// CHECK-NEXT: parameter\d+\.v2u = \(2u \* parameter\d+\.v2u\);
// CHECK-NEXT: parameter\d+\.v3u = packed_uint3\(\(2u \* uint3\(parameter\d+\.v3u\)\)\);
// CHECK-NEXT: parameter\d+\.v4u = \(2u \* parameter\d+\.v4u\);
// CHECK-NEXT: parameter\d+\.f = \(2 \* parameter\d+\.f\);
// CHECK-NEXT: parameter\d+\.u = \(2u \* parameter\d+\.u\);
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: parameter\d+\.v2f\.x = \(2 \* parameter\d+\.v2f\.x\);
// CHECK-NEXT: parameter\d+\.v3f\.x = \(2 \* parameter\d+\.v3f\.x\);
// CHECK-NEXT: parameter\d+\.v4f\.x = \(2 \* parameter\d+\.v4f\.x\);
// CHECK-NEXT: parameter\d+\.v2u\.x = \(2u \* parameter\d+\.v2u\.x\);
// CHECK-NEXT: parameter\d+\.v3u\.x = \(2u \* parameter\d+\.v3u\.x\);
// CHECK-NEXT: parameter\d+\.v4u\.x = \(2u \* parameter\d+\.v4u\.x\);
// CHECK-NEXT: parameter\d+\.v2f = \(2 \* parameter\d+\.v2f\);
// CHECK-NEXT: parameter\d+\.v3f = packed_float3\(\(2 \* parameter\d+\.v3f\)\);
// CHECK-NEXT: parameter\d+\.v4f = \(2 \* parameter\d+\.v4f\);
// CHECK-NEXT: parameter\d+\.v2u = \(2u \* parameter\d+\.v2u\);
// CHECK-NEXT: parameter\d+\.v3u = packed_uint3\(\(2u \* parameter\d+\.v3u\)\);
// CHECK-NEXT: parameter\d+\.v4u = \(2u \* parameter\d+\.v4u\);
// CHECK-NEXT: parameter\d+\.f = \(2 \* parameter\d+\.f\);
// CHECK-NEXT: parameter\d+\.u = \(2u \* parameter\d+\.u\);
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: parameter\d+\.v2f\.x = -parameter\d+\.v2f\.x;
// CHECK-NEXT: parameter\d+\.v3f\.x = -parameter\d+\.v3f\.x;
// CHECK-NEXT: parameter\d+\.v4f\.x = -parameter\d+\.v4f\.x;
// CHECK-NEXT: parameter\d+\.v2f = -parameter\d+\.v2f;
// CHECK-NEXT: parameter\d+\.v3f = -float3\(parameter\d+\.v3f\);
// CHECK-NEXT: parameter\d+\.v4f = -parameter\d+\.v4f;
// CHECK-NEXT: parameter\d+\.f = -parameter\d+\.f;
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: parameter\d+\.v2f\.x = -parameter\d+\.v2f\.x;
// CHECK-NEXT: parameter\d+\.v3f\.x = -parameter\d+\.v3f\.x;
// CHECK-NEXT: parameter\d+\.v4f\.x = -parameter\d+\.v4f\.x;
// CHECK-NEXT: parameter\d+\.v2f = -parameter\d+\.v2f;
// CHECK-NEXT: parameter\d+\.v3f = packed_float3\(-float3\(parameter\d+\.v3f\)\);
// CHECK-NEXT: parameter\d+\.v4f = -parameter\d+\.v4f;
// CHECK-NEXT: parameter\d+\.f = -parameter\d+\.f;
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: parameter\d+\.v2f\.x = -parameter\d+\.v2f\.x;
// CHECK-NEXT: parameter\d+\.v3f\.x = -parameter\d+\.v3f\.x;
// CHECK-NEXT: parameter\d+\.v4f\.x = -parameter\d+\.v4f\.x;
// CHECK-NEXT: parameter\d+\.v2f = -parameter\d+\.v2f;
// CHECK-NEXT: parameter\d+\.v3f = packed_float3\(-parameter\d+\.v3f\);
// CHECK-NEXT: parameter\d+\.v4f = -parameter\d+\.v4f;
// CHECK-NEXT: parameter\d+\.f = -parameter\d+\.f;
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: parameter\d+\.v2f\.x = abs\(parameter\d+\.v2f\.x\);
// CHECK-NEXT: parameter\d+\.v3f\.x = abs\(parameter\d+\.v3f\.x\);
// CHECK-NEXT: parameter\d+\.v4f\.x = abs\(parameter\d+\.v4f\.x\);
// CHECK-NEXT: parameter\d+\.v2u\.x = abs\(parameter\d+\.v2u\.x\);
// CHECK-NEXT: parameter\d+\.v3u\.x = abs\(parameter\d+\.v3u\.x\);
// CHECK-NEXT: parameter\d+\.v4u\.x = abs\(parameter\d+\.v4u\.x\);
// CHECK-NEXT: parameter\d+\.v2f = abs\(parameter\d+\.v2f\);
// CHECK-NEXT: parameter\d+\.v3f = abs\(float3\(parameter\d+\.v3f\)\);
// CHECK-NEXT: parameter\d+\.v4f = abs\(parameter\d+\.v4f\);
// CHECK-NEXT: parameter\d+\.v2u = abs\(parameter\d+\.v2u\);
// CHECK-NEXT: parameter\d+\.v3u = abs\(uint3\(parameter\d+\.v3u\)\);
// CHECK-NEXT: parameter\d+\.v4u = abs\(parameter\d+\.v4u\);
// CHECK-NEXT: parameter\d+\.f = abs\(parameter\d+\.f\);
// CHECK-NEXT: parameter\d+\.u = abs\(parameter\d+\.u\);
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: parameter\d+\.v2f\.x = abs\(parameter\d+\.v2f\.x\);
// CHECK-NEXT: parameter\d+\.v3f\.x = abs\(parameter\d+\.v3f\.x\);
// CHECK-NEXT: parameter\d+\.v4f\.x = abs\(parameter\d+\.v4f\.x\);
// CHECK-NEXT: parameter\d+\.v2u\.x = abs\(parameter\d+\.v2u\.x\);
// CHECK-NEXT: parameter\d+\.v3u\.x = abs\(parameter\d+\.v3u\.x\);
// CHECK-NEXT: parameter\d+\.v4u\.x = abs\(parameter\d+\.v4u\.x\);
// CHECK-NEXT: parameter\d+\.v2f = abs\(parameter\d+\.v2f\);
// CHECK-NEXT: parameter\d+\.v3f = packed_float3\(abs\(float3\(parameter\d+\.v3f\)\)\);
// CHECK-NEXT: parameter\d+\.v4f = abs\(parameter\d+\.v4f\);
// CHECK-NEXT: parameter\d+\.v2u = abs\(parameter\d+\.v2u\);
// CHECK-NEXT: parameter\d+\.v3u = packed_uint3\(abs\(uint3\(parameter\d+\.v3u\)\)\);
// CHECK-NEXT: parameter\d+\.v4u = abs\(parameter\d+\.v4u\);
// CHECK-NEXT: parameter\d+\.f = abs\(parameter\d+\.f\);
// CHECK-NEXT: parameter\d+\.u = abs\(parameter\d+\.u\);
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: parameter\d+\.v2f\.x = abs\(parameter\d+\.v2f\.x\);
// CHECK-NEXT: parameter\d+\.v3f\.x = abs\(parameter\d+\.v3f\.x\);
// CHECK-NEXT: parameter\d+\.v4f\.x = abs\(parameter\d+\.v4f\.x\);
// CHECK-NEXT: parameter\d+\.v2u\.x = abs\(parameter\d+\.v2u\.x\);
// CHECK-NEXT: parameter\d+\.v3u\.x = abs\(parameter\d+\.v3u\.x\);
// CHECK-NEXT: parameter\d+\.v4u\.x = abs\(parameter\d+\.v4u\.x\);
// CHECK-NEXT: parameter\d+\.v2f = abs\(parameter\d+\.v2f\);
// CHECK-NEXT: parameter\d+\.v3f = packed_float3\(abs\(parameter\d+\.v3f\)\);
// CHECK-NEXT: parameter\d+\.v4f = abs\(parameter\d+\.v4f\);
// CHECK-NEXT: parameter\d+\.v2u = abs\(parameter\d+\.v2u\);
// CHECK-NEXT: parameter\d+\.v3u = packed_uint3\(abs\(parameter\d+\.v3u\)\);
// CHECK-NEXT: parameter\d+\.v4u = abs\(parameter\d+\.v4u\);
// CHECK-NEXT: parameter\d+\.f = abs\(parameter\d+\.f\);
// CHECK-NEXT: parameter\d+\.u = abs\(parameter\d+\.u\);
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;
}
@compute @workgroup_size(1)
fn main()
{
_ = testUnpacked();
_ = testAssignment();
_ = testFieldAccess();
_ = testIndexAccess();
_ = testBinaryOperations();
_ = testUnaryOperations();
_ = testCall();
}
|
