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
|
commit 83ef5b48800a47cc30b3d4bcfdf31de9c3bd0dc8
Author: Patrick Lauer <patrick@gentoo.org>
Date: Sun Jul 28 07:43:54 2024 +0000
Revert "ockl: Don't use wave32 ballot builtin"
This reverts commit 066a0b2716b7ade96a2b3e79e5ddcd0c110e9f98.
diff --git a/ockl/src/dm.cl b/ockl/src/dm.cl
index 18efc54203b7..a3f06c448aee 100644
--- a/ockl/src/dm.cl
+++ b/ockl/src/dm.cl
@@ -287,14 +287,9 @@ first(__global void * v)
return __builtin_astype(w2, __global void *);
}
-// Read val from one active lane whose predicate is one.
-// If no lanes have the predicate set, return none
-// This is like first, except that first may not have its predicate set
+REQUIRES_WAVE64
static uint
-elect_uint(int pred, uint val, uint none)
-{
- // Pretend wave32 doesn't exist. The wave64 ballot works, and the high half
- // will fold out as 0.
+elect_uint_wave64(int pred, uint val, uint none) {
uint ret = none;
ulong mask = __builtin_amdgcn_ballot_w64(pred != 0);
@@ -306,14 +301,51 @@ elect_uint(int pred, uint val, uint none)
return ret;
}
-// Count the number of nonzero arguments across the wave
+REQUIRES_WAVE32
static uint
-votes(bool b)
+elect_uint_wave32(int pred, uint val, uint none) {
+ uint ret = none;
+ uint mask = __builtin_amdgcn_ballot_w32(pred != 0);
+ if (mask != 0U) {
+ uint l = __ockl_ctz_u32(mask);
+ ret = __builtin_amdgcn_ds_bpermute(l << 2, val);
+ }
+
+ return ret;
+}
+
+// Read val from one active lane whose predicate is one.
+// If no lanes have the predicate set, return none
+// This is like first, except that first may not have its predicate set
+static uint
+elect_uint(int pred, uint val, uint none)
+{
+ return __oclc_wavefrontsize64 ? elect_uint_wave64(pred, val, none) : elect_uint_wave32(pred, val, none);
+}
+
+REQUIRES_WAVE64
+static uint
+votes_wave64(bool b)
{
ulong mask = __builtin_amdgcn_ballot_w64(b);
return __builtin_popcountl(mask);
}
+REQUIRES_WAVE32
+static uint
+votes_wave32(bool b)
+{
+ uint mask = __builtin_amdgcn_ballot_w32(b);
+ return __builtin_popcount(mask);
+}
+
+// Count the number of nonzero arguments across the wave
+static uint
+votes(bool b)
+{
+ return __oclc_wavefrontsize64 ? votes_wave64(b) : votes_wave32(b);
+}
+
// The kind of the smallest block that can hold sz bytes
static uint
size_to_kind(uint sz)
diff --git a/ockl/src/wfaas.cl b/ockl/src/wfaas.cl
index 3861a5bb3eab..4dab97cea5f7 100644
--- a/ockl/src/wfaas.cl
+++ b/ockl/src/wfaas.cl
@@ -21,25 +21,60 @@ static int optimizationBarrierHack(int in_val)
return out_val;
}
+REQUIRES_WAVE32
+static bool wfany_impl_w32(int e) {
+ return __builtin_amdgcn_ballot_w32(e) != 0;
+}
+
+REQUIRES_WAVE64
+static bool wfany_impl_w64(int e) {
+ return __builtin_amdgcn_ballot_w64(e) != 0;
+}
+
ATTR bool
OCKL_MANGLE_I32(wfany)(int e)
{
e = optimizationBarrierHack(e);
- return __builtin_amdgcn_ballot_w64(e) != 0;
+ return __oclc_wavefrontsize64 ?
+ wfany_impl_w64(e) : wfany_impl_w32(e);
+}
+
+REQUIRES_WAVE32
+static bool wfall_impl_w32(int e) {
+ return __builtin_amdgcn_ballot_w32(e) == __builtin_amdgcn_read_exec_lo();
+}
+
+REQUIRES_WAVE64
+static bool wfall_impl_w64(int e) {
+ return __builtin_amdgcn_ballot_w64(e) == __builtin_amdgcn_read_exec();
}
ATTR bool
OCKL_MANGLE_I32(wfall)(int e)
{
e = optimizationBarrierHack(e);
- return __builtin_amdgcn_ballot_w64(e) == __builtin_amdgcn_read_exec();
+ return __oclc_wavefrontsize64 ?
+ wfall_impl_w64(e) : wfall_impl_w32(e);
+}
+
+
+REQUIRES_WAVE32
+static bool wfsame_impl_w32(int e) {
+ uint u = __builtin_amdgcn_ballot_w32(e);
+ return (u == 0) | (u == __builtin_amdgcn_read_exec_lo());
+}
+
+REQUIRES_WAVE64
+static bool wfsame_impl_w64(int e) {
+ ulong u = __builtin_amdgcn_ballot_w64(e);
+ return (u == 0UL) | (u == __builtin_amdgcn_read_exec());
}
ATTR bool
OCKL_MANGLE_I32(wfsame)(int e)
{
e = optimizationBarrierHack(e);
- ulong u = __builtin_amdgcn_ballot_w64(e);
- return (u == 0UL) | (u == __builtin_amdgcn_read_exec());
+ return __oclc_wavefrontsize64 ?
+ wfsame_impl_w64(e) : wfsame_impl_w32(e);
}
|
