Skip to content

Commit 17f0d2b

Browse files
vladimirlazvladimirlaz
authored
[NFC][SYCL] Update sub_group LIT tests (#2219)
- Remove sub_group/common_ocl.cpp because it duplicates sub_group/common.cpp and directly use OpenCL API that causes instability on some configurations. - Fix sub_group/shuffle*.cpp tests to align with shuffle_xor restrictions mentioned in spec: "If the result of the XOR is greater than max_sub_group_size then it is considered out-of-range"
1 parent 4a1eb53 commit 17f0d2b

File tree

3 files changed

+69
-133
lines changed

3 files changed

+69
-133
lines changed

‎sycl/test/sub_group/common_ocl.cpp‎

Lines changed: 0 additions & 106 deletions
This file was deleted.

‎sycl/test/sub_group/generic-shuffle.cpp‎

Lines changed: 39 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,7 @@
1818
#include "helper.hpp"
1919
#include <CL/sycl.hpp>
2020
#include <complex>
21-
template <typename T>
22-
class pointer_kernel;
21+
template <typename T> class pointer_kernel;
2322

2423
using namespace cl::sycl;
2524

@@ -59,8 +58,9 @@ void check_pointer(queue &Queue, size_t G = 240, size_t L = 60) {
5958
/* Save GID+SGID */
6059
acc_down[NdItem.get_global_id()] = SG.shuffle_down(ptr, sgid);
6160

62-
/* Save GID XOR SGID */
63-
acc_xor[NdItem.get_global_id()] = SG.shuffle_xor(ptr, sgid);
61+
/* Save GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
62+
acc_xor[NdItem.get_global_id()] =
63+
SG.shuffle_xor(ptr, sgid % SG.get_max_local_range()[0]);
6464
});
6565
});
6666
auto acc = buf.template get_access<access::mode::read_write>();
@@ -71,30 +71,44 @@ void check_pointer(queue &Queue, size_t G = 240, size_t L = 60) {
7171

7272
size_t sg_size = sgsizeacc[0];
7373
int SGid = 0;
74+
int SGLid = 0;
75+
int SGBeginGid = 0;
7476
for (int j = 0; j < G; j++) {
7577
if (j % L % sg_size == 0) {
7678
SGid++;
79+
SGLid = 0;
80+
SGBeginGid = j;
7781
}
7882
if (j % L == 0) {
7983
SGid = 0;
84+
SGLid = 0;
85+
SGBeginGid = j;
8086
}
8187

8288
/*GID of middle element in every subgroup*/
83-
exit_if_not_equal(acc[j], static_cast<T *>(0x0) + (j / L * L + SGid * sg_size + sg_size / 2),
89+
exit_if_not_equal(acc[j],
90+
static_cast<T *>(0x0) +
91+
(j / L * L + SGid * sg_size + sg_size / 2),
8492
"shuffle");
8593

8694
/* Value GID+SGID for all element except last SGID in SG*/
8795
if (j % L % sg_size + SGid < sg_size && j % L + SGid < L) {
88-
exit_if_not_equal(acc_down[j], static_cast<T *>(0x0) + (j + SGid), "shuffle_down");
96+
exit_if_not_equal(acc_down[j], static_cast<T *>(0x0) + (j + SGid),
97+
"shuffle_down");
8998
}
9099

91100
/* Value GID-SGID for all element except first SGID in SG*/
92101
if (j % L % sg_size >= SGid) {
93-
exit_if_not_equal(acc_up[j], static_cast<T *>(0x0) + (j - SGid), "shuffle_up");
102+
exit_if_not_equal(acc_up[j], static_cast<T *>(0x0) + (j - SGid),
103+
"shuffle_up");
94104
}
95105

96-
/* GID XOR SGID */
97-
exit_if_not_equal(acc_xor[j], static_cast<T *>(0x0) + (j ^ SGid), "shuffle_xor");
106+
/* Value GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
107+
exit_if_not_equal(acc_xor[j],
108+
static_cast<T *>(0x0) +
109+
(SGBeginGid + (SGLid ^ (SGid % sg_size))),
110+
"shuffle_xor");
111+
SGLid++;
98112
}
99113
} catch (exception e) {
100114
std::cout << "SYCL exception caught: " << e.what();
@@ -145,8 +159,9 @@ void check_struct(queue &Queue, Generator &Gen, size_t G = 240, size_t L = 60) {
145159
/* Save GID+SGID */
146160
acc_down[NdItem.get_global_id()] = SG.shuffle_down(val, sgid);
147161

148-
/* Save GID XOR SGID */
149-
acc_xor[NdItem.get_global_id()] = SG.shuffle_xor(val, sgid);
162+
/* Save GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
163+
acc_xor[NdItem.get_global_id()] =
164+
SG.shuffle_xor(val, sgid % SG.get_max_local_range()[0]);
150165
});
151166
});
152167
auto acc = buf.template get_access<access::mode::read_write>();
@@ -157,17 +172,23 @@ void check_struct(queue &Queue, Generator &Gen, size_t G = 240, size_t L = 60) {
157172

158173
size_t sg_size = sgsizeacc[0];
159174
int SGid = 0;
175+
int SGLid = 0;
176+
int SGBeginGid = 0;
160177
for (int j = 0; j < G; j++) {
161178
if (j % L % sg_size == 0) {
162179
SGid++;
180+
SGLid = 0;
181+
SGBeginGid = j;
163182
}
164183
if (j % L == 0) {
165184
SGid = 0;
185+
SGLid = 0;
186+
SGBeginGid = j;
166187
}
167188

168189
/*GID of middle element in every subgroup*/
169-
exit_if_not_equal(acc[j], values[j / L * L + SGid * sg_size + sg_size / 2],
170-
"shuffle");
190+
exit_if_not_equal(
191+
acc[j], values[j / L * L + SGid * sg_size + sg_size / 2], "shuffle");
171192

172193
/* Value GID+SGID for all element except last SGID in SG*/
173194
if (j % L % sg_size + SGid < sg_size && j % L + SGid < L) {
@@ -179,8 +200,11 @@ void check_struct(queue &Queue, Generator &Gen, size_t G = 240, size_t L = 60) {
179200
exit_if_not_equal(acc_up[j], values[j - SGid], "shuffle_up");
180201
}
181202

182-
/* GID XOR SGID */
183-
exit_if_not_equal(acc_xor[j], values[j ^ SGid], "shuffle_xor");
203+
/* Value GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
204+
exit_if_not_equal(acc_xor[j],
205+
values[SGBeginGid + (SGLid ^ (SGid % sg_size))],
206+
"shuffle_xor");
207+
SGLid++;
184208
}
185209
} catch (exception e) {
186210
std::cout << "SYCL exception caught: " << e.what();

‎sycl/test/sub_group/shuffle.hpp‎

Lines changed: 30 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -8,8 +8,7 @@
88

99
#include "helper.hpp"
1010
#include <CL/sycl.hpp>
11-
template <typename T, int N>
12-
class sycl_subgr;
11+
template <typename T, int N> class sycl_subgr;
1312

1413
using namespace cl::sycl;
1514

@@ -66,8 +65,9 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
6665
acc_up[NdItem.get_global_id()] = SG.shuffle_up(vwggid, sgid);
6766
/* Save GID+SGID */
6867
acc_down[NdItem.get_global_id()] = SG.shuffle_down(vwggid, sgid);
69-
/* Save GID XOR SGID */
70-
acc_xor[NdItem.get_global_id()] = SG.shuffle_xor(vwggid, sgid);
68+
/* Save GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
69+
acc_xor[NdItem.get_global_id()] =
70+
SG.shuffle_xor(vwggid, sgid % SG.get_max_local_range()[0]);
7171
});
7272
});
7373
auto acc = buf.template get_access<access::mode::read_write>();
@@ -81,12 +81,18 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
8181

8282
size_t sg_size = sgsizeacc[0];
8383
int SGid = 0;
84+
int SGLid = 0;
85+
int SGBeginGid = 0;
8486
for (int j = 0; j < G; j++) {
8587
if (j % L % sg_size == 0) {
8688
SGid++;
89+
SGLid = 0;
90+
SGBeginGid = j;
8791
}
8892
if (j % L == 0) {
8993
SGid = 0;
94+
SGLid = 0;
95+
SGBeginGid = j;
9096
}
9197
/*GID of middle element in every subgroup*/
9298
exit_if_not_equal_vec<T, N>(
@@ -115,17 +121,19 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
115121
exit_if_not_equal_vec(acc2_up[j], vec<T, N>(j - SGid + sg_size),
116122
"shuffle2_up");
117123
}
118-
/* GID XOR SGID */
119-
exit_if_not_equal_vec(acc_xor[j], vec<T, N>(j ^ SGid), "shuffle_xor");
124+
/* Value GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
125+
exit_if_not_equal_vec(acc_xor[j],
126+
vec<T, N>(SGBeginGid + (SGLid ^ (SGid % sg_size))),
127+
"shuffle_xor");
128+
SGLid++;
120129
}
121130
} catch (exception e) {
122131
std::cout << "SYCL exception caught: " << e.what();
123132
exit(1);
124133
}
125134
}
126135

127-
template <typename T>
128-
void check(queue &Queue, size_t G = 240, size_t L = 60) {
136+
template <typename T> void check(queue &Queue, size_t G = 240, size_t L = 60) {
129137
try {
130138
nd_range<1> NdRange(G, L);
131139
buffer<T> buf2(G);
@@ -171,8 +179,9 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
171179
acc_up[NdItem.get_global_id()] = SG.shuffle_up<T>(wggid, sgid);
172180
/* Save GID+SGID */
173181
acc_down[NdItem.get_global_id()] = SG.shuffle_down<T>(wggid, sgid);
174-
/* Save GID XOR SGID */
175-
acc_xor[NdItem.get_global_id()] = SG.shuffle_xor<T>(wggid, sgid);
182+
/* Save GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
183+
acc_xor[NdItem.get_global_id()] =
184+
SG.shuffle_xor<T>(wggid, sgid % SG.get_max_local_range()[0]);
176185
});
177186
});
178187
auto acc = buf.template get_access<access::mode::read_write>();
@@ -186,13 +195,20 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
186195

187196
size_t sg_size = sgsizeacc[0];
188197
int SGid = 0;
198+
int SGLid = 0;
199+
int SGBeginGid = 0;
189200
for (int j = 0; j < G; j++) {
190201
if (j % L % sg_size == 0) {
191202
SGid++;
203+
SGLid = 0;
204+
SGBeginGid = j;
192205
}
193206
if (j % L == 0) {
194207
SGid = 0;
208+
SGLid = 0;
209+
SGBeginGid = j;
195210
}
211+
196212
/*GID of middle element in every subgroup*/
197213
exit_if_not_equal<T>(acc[j], j / L * L + SGid * sg_size + sg_size / 2,
198214
"shuffle");
@@ -215,8 +231,10 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
215231
if (j % L - SGid + sg_size < L) /* Do not go out LG*/
216232
exit_if_not_equal<T>(acc2_up[j], j - SGid + sg_size, "shuffle2_up");
217233
}
218-
/* GID XOR SGID */
219-
exit_if_not_equal<T>(acc_xor[j], j ^ SGid, "shuffle_xor");
234+
/* Value GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
235+
exit_if_not_equal<T>(acc_xor[j], SGBeginGid + (SGLid ^ (SGid % sg_size)),
236+
"shuffle_xor");
237+
SGLid++;
220238
}
221239
} catch (exception e) {
222240
std::cout << "SYCL exception caught: " << e.what();

0 commit comments

Comments
 (0)