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
|
// RUN: %metal-compile main
fn f1(x: ptr<function, i32>) -> i32
{
return *x;
}
fn f2(x: ptr<private, i32>) -> i32
{
return *x;
}
fn f3(x: vec3<f32>) { }
fn f4(p: ptr<function, vec2<i32>>) -> i32
{
return (*p).x;
}
fn f5(p: ptr<function, vec2<i32>>)
{
*p = vec2(0);
(*p).x = 42;
}
var<private> global: i32;
struct S { x: vec3<f32> }
@group(0) @binding(0) var<storage, read_write> global2 : S;
fn testSimplePointerRewriting()
{
let x = &global2;
f3((*x).x);
}
fn testIndirectPointerRewriting()
{
let x = &global2;
let y = x;
let z = y;
f3((*z).x);
}
fn testPhonyPointerElimination()
{
_ = &global2;
}
fn testShadowedGlobalRewriting()
{
let x = &global2;
let global2 = 0;
f3((*x).x);
}
fn testShadowedLocalRewriting()
{
var x: S;
let y = &x;
{
let x = 0;
f3((*y).x);
}
}
fn testVectorAccessPrecedence()
{
var v = vec2(0);
let p = &v;
let x = f4(p);
}
fn testAssignment()
{
var v = vec2(0);
let p = &v;
f5(p);
*&v = vec2(13);
}
@compute @workgroup_size(1)
fn main()
{
var local: i32;
_ = &global2;
let x = &global2;
let y = x;
f1(&local);
f2(&global);
testSimplePointerRewriting();
testIndirectPointerRewriting();
testPhonyPointerElimination();
testShadowedGlobalRewriting();
testShadowedLocalRewriting();
testVectorAccessPrecedence();
testAssignment();
}
|
