Skip to content

Commit 4fcc744

Browse files
ayylolayylol
authored
[SYCL][E2E] Remove uses of OpenCL primitives in Basic/image e2e tests (#14084)
1 parent 25bfb0b commit 4fcc744

File tree

4 files changed

+180
-198
lines changed

4 files changed

+180
-198
lines changed

‎sycl/test-e2e/Basic/image/image.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -91,7 +91,7 @@ int main() {
9191

9292
constexpr int dims = 1;
9393

94-
using data_img = sycl::cl_float4;
94+
using data_img = sycl::float4;
9595
constexpr auto mode_img = sycl::access::mode::read;
9696
constexpr auto target_img = sycl::target::image;
9797
const auto range_img = sycl::range<dims>(3);

‎sycl/test-e2e/Basic/image/image_accessor_readsampler.cpp

Lines changed: 63 additions & 81 deletions
Original file line numberDiff line numberDiff line change
@@ -26,14 +26,13 @@ namespace s = sycl;
2626

2727
template <int unique_number> class kernel_class;
2828

29-
void validateReadData(s::cl_float4 ReadData, s::cl_float4 ExpectedColor,
30-
s::cl_int precision = 1) {
29+
void validateReadData(s::float4 ReadData, s::float4 ExpectedColor,
30+
int precision = 1) {
3131
// Maximum difference of 1.5 ULP is allowed when precision = 1.
32-
s::cl_int4 PixelDataInt = ReadData.template as<s::cl_int4>();
33-
s::cl_int4 ExpectedDataInt = ExpectedColor.template as<s::cl_int4>();
34-
s::cl_int4 Diff = ExpectedDataInt - PixelDataInt;
35-
s::cl_int DataIsCorrect =
36-
s::all((Diff <= precision) && (Diff >= (-precision)));
32+
s::int4 PixelDataInt = ReadData.template as<s::int4>();
33+
s::int4 ExpectedDataInt = ExpectedColor.template as<s::int4>();
34+
s::int4 Diff = ExpectedDataInt - PixelDataInt;
35+
int DataIsCorrect = s::all((Diff <= precision) && (Diff >= (-precision)));
3736
#if DEBUG_OUTPUT
3837
{
3938
if (DataIsCorrect) {
@@ -49,28 +48,30 @@ void validateReadData(s::cl_float4 ReadData, s::cl_float4 ExpectedColor,
4948
Diff.dump();
5049
}
5150
#else
52-
{ assert(DataIsCorrect); }
51+
{
52+
assert(DataIsCorrect);
53+
}
5354
#endif
5455
}
5556

5657
template <int i>
57-
void checkReadSampler(char *host_ptr, s::sampler Sampler, s::cl_float4 Coord,
58-
s::cl_float4 ExpectedColor, s::cl_int precision = 1) {
58+
void checkReadSampler(char *host_ptr, s::sampler Sampler, s::float4 Coord,
59+
s::float4 ExpectedColor, int precision = 1) {
5960

60-
s::cl_float4 ReadData;
61+
s::float4 ReadData;
6162
{
6263
// image with dim = 3
6364
s::image<3> Img(host_ptr, s::image_channel_order::rgba,
6465
s::image_channel_type::snorm_int8, s::range<3>{2, 3, 4});
6566
s::queue myQueue;
66-
s::buffer<s::cl_float4, 1> ReadDataBuf(&ReadData, s::range<1>(1));
67+
s::buffer<s::float4, 1> ReadDataBuf(&ReadData, s::range<1>(1));
6768
myQueue.submit([&](s::handler &cgh) {
68-
auto ReadAcc = Img.get_access<s::cl_float4, s::access::mode::read>(cgh);
69-
s::accessor<s::cl_float4, 1, s::access::mode::write> ReadDataBufAcc(
69+
auto ReadAcc = Img.get_access<s::float4, s::access::mode::read>(cgh);
70+
s::accessor<s::float4, 1, s::access::mode::write> ReadDataBufAcc(
7071
ReadDataBuf, cgh);
7172

7273
cgh.single_task<class kernel_class<i>>([=]() {
73-
s::cl_float4 RetColor = ReadAcc.read(Coord, Sampler);
74+
s::float4 RetColor = ReadAcc.read(Coord, Sampler);
7475
ReadDataBufAcc[0] = RetColor;
7576
});
7677
});
@@ -90,9 +91,8 @@ void checkSamplerNearest() {
9091
// addressing_mode::mirrored_repeat
9192
{
9293
// Out-of-range mirrored_repeat mode
93-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
94-
s::cl_float4 ExpectedValue =
95-
s::cl_float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
94+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
95+
s::float4 ExpectedValue = s::float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
9696
auto Sampler = s::sampler(s::coordinate_normalization_mode::normalized,
9797
s::addressing_mode::mirrored_repeat,
9898
s::filtering_mode::nearest);
@@ -102,9 +102,8 @@ void checkSamplerNearest() {
102102
// addressing_mode::repeat
103103
{
104104
// Out-of-range repeat mode
105-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
106-
s::cl_float4 ExpectedValue =
107-
s::cl_float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
105+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
106+
s::float4 ExpectedValue = s::float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
108107
auto Sampler =
109108
s::sampler(s::coordinate_normalization_mode::normalized,
110109
s::addressing_mode::repeat, s::filtering_mode::nearest);
@@ -114,9 +113,8 @@ void checkSamplerNearest() {
114113
// addressing_mode::clamp_to_edge
115114
{
116115
// Out-of-range Edge Color
117-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
118-
s::cl_float4 ExpectedValue =
119-
s::cl_float4(88.0f, 89.0f, 90.0f, 91.0f) / 127.0f;
116+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
117+
s::float4 ExpectedValue = s::float4(88.0f, 89.0f, 90.0f, 91.0f) / 127.0f;
120118
auto Sampler = s::sampler(s::coordinate_normalization_mode::normalized,
121119
s::addressing_mode::clamp_to_edge,
122120
s::filtering_mode::nearest);
@@ -126,8 +124,8 @@ void checkSamplerNearest() {
126124
// addressing_mode::clamp
127125
{
128126
// Out-of-range Border Color
129-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
130-
s::cl_float4 ExpectedValue = s::cl_float4(0.0f, 0.0f, 0.0f, 0.0f);
127+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
128+
s::float4 ExpectedValue = s::float4(0.0f, 0.0f, 0.0f, 0.0f);
131129
auto Sampler =
132130
s::sampler(s::coordinate_normalization_mode::normalized,
133131
s::addressing_mode::clamp, s::filtering_mode::nearest);
@@ -137,9 +135,8 @@ void checkSamplerNearest() {
137135
// addressing_mode::none
138136
{
139137
// In-range for consistent return value.
140-
s::cl_float4 Coord(0.0f, 0.5f, 0.75f, 0.0f);
141-
s::cl_float4 ExpectedValue =
142-
s::cl_float4(80.0f, 81.0f, 82.0f, 83.0f) / 127.0f;
138+
s::float4 Coord(0.0f, 0.5f, 0.75f, 0.0f);
139+
s::float4 ExpectedValue = s::float4(80.0f, 81.0f, 82.0f, 83.0f) / 127.0f;
143140
auto Sampler =
144141
s::sampler(s::coordinate_normalization_mode::normalized,
145142
s::addressing_mode::none, s::filtering_mode::nearest);
@@ -149,9 +146,8 @@ void checkSamplerNearest() {
149146
// B. coordinate_normalization_mode::unnormalized
150147
// addressing_mode::clamp_to_edge
151148
{
152-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
153-
s::cl_float4 ExpectedValue =
154-
s::cl_float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
149+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
150+
s::float4 ExpectedValue = s::float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
155151
auto Sampler = s::sampler(s::coordinate_normalization_mode::unnormalized,
156152
s::addressing_mode::clamp_to_edge,
157153
s::filtering_mode::nearest);
@@ -160,9 +156,8 @@ void checkSamplerNearest() {
160156

161157
// addressing_mode::clamp
162158
{
163-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
164-
s::cl_float4 ExpectedValue =
165-
s::cl_float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
159+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
160+
s::float4 ExpectedValue = s::float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
166161
auto Sampler =
167162
s::sampler(s::coordinate_normalization_mode::unnormalized,
168163
s::addressing_mode::clamp, s::filtering_mode::nearest);
@@ -172,9 +167,8 @@ void checkSamplerNearest() {
172167
// addressing_mode::none
173168
{
174169
// In-range for consistent return value.
175-
s::cl_float4 Coord(0.0f, 1.0f, 2.0f, 0.0f);
176-
s::cl_float4 ExpectedValue =
177-
s::cl_float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
170+
s::float4 Coord(0.0f, 1.0f, 2.0f, 0.0f);
171+
s::float4 ExpectedValue = s::float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
178172
auto Sampler =
179173
s::sampler(s::coordinate_normalization_mode::unnormalized,
180174
s::addressing_mode::none, s::filtering_mode::nearest);
@@ -190,7 +184,7 @@ void checkSamplerNearest() {
190184
// value of 15000 ULP is used.
191185
void checkSamplerLinear() {
192186

193-
const s::cl_int PrecisionInULP = 15000;
187+
const int PrecisionInULP = 15000;
194188
// create image:
195189
char host_ptr[100];
196190
for (int i = 0; i < 100; i++)
@@ -201,9 +195,8 @@ void checkSamplerLinear() {
201195
// addressing_mode::mirrored_repeat
202196
{
203197
// Out-of-range mirrored_repeat mode
204-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
205-
s::cl_float4 ExpectedValue =
206-
s::cl_float4(44.0f, 45.0f, 46.0f, 47.0f) / 127.0f;
198+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
199+
s::float4 ExpectedValue = s::float4(44.0f, 45.0f, 46.0f, 47.0f) / 127.0f;
207200
auto Sampler = s::sampler(s::coordinate_normalization_mode::normalized,
208201
s::addressing_mode::mirrored_repeat,
209202
s::filtering_mode::linear);
@@ -212,9 +205,8 @@ void checkSamplerLinear() {
212205
}
213206
{
214207
// In-range mirrored_repeat mode
215-
s::cl_float4 Coord(0.0f, 0.25f, 0.55f, 0.0f);
216-
s::cl_float4 ExpectedValue =
217-
s::cl_float4(42.8f, 43.8f, 44.8f, 45.8f) / 127.0f;
208+
s::float4 Coord(0.0f, 0.25f, 0.55f, 0.0f);
209+
s::float4 ExpectedValue = s::float4(42.8f, 43.8f, 44.8f, 45.8f) / 127.0f;
218210
auto Sampler = s::sampler(s::coordinate_normalization_mode::normalized,
219211
s::addressing_mode::mirrored_repeat,
220212
s::filtering_mode::linear);
@@ -225,9 +217,8 @@ void checkSamplerLinear() {
225217
// addressing_mode::repeat
226218
{
227219
// Out-of-range repeat mode
228-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
229-
s::cl_float4 ExpectedValue =
230-
s::cl_float4(46.0f, 47.0f, 48.0f, 49.0f) / 127.0f;
220+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
221+
s::float4 ExpectedValue = s::float4(46.0f, 47.0f, 48.0f, 49.0f) / 127.0f;
231222
auto Sampler =
232223
s::sampler(s::coordinate_normalization_mode::normalized,
233224
s::addressing_mode::repeat, s::filtering_mode::linear);
@@ -236,9 +227,8 @@ void checkSamplerLinear() {
236227
}
237228
{
238229
// In-range repeat mode
239-
s::cl_float4 Coord(0.0f, 0.25f, 0.55f, 0.0f);
240-
s::cl_float4 ExpectedValue =
241-
s::cl_float4(44.8f, 45.8f, 46.8f, 47.8f) / 127.0f;
230+
s::float4 Coord(0.0f, 0.25f, 0.55f, 0.0f);
231+
s::float4 ExpectedValue = s::float4(44.8f, 45.8f, 46.8f, 47.8f) / 127.0f;
242232
auto Sampler =
243233
s::sampler(s::coordinate_normalization_mode::normalized,
244234
s::addressing_mode::repeat, s::filtering_mode::linear);
@@ -249,19 +239,17 @@ void checkSamplerLinear() {
249239
// addressing_mode::clamp_to_edge
250240
{
251241
// Out-of-range Edge Color
252-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
253-
s::cl_float4 ExpectedValue =
254-
s::cl_float4(88.0f, 89.0f, 90.0f, 91.0f) / 127.0f;
242+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
243+
s::float4 ExpectedValue = s::float4(88.0f, 89.0f, 90.0f, 91.0f) / 127.0f;
255244
auto Sampler = s::sampler(s::coordinate_normalization_mode::normalized,
256245
s::addressing_mode::clamp_to_edge,
257246
s::filtering_mode::linear);
258247
checkReadSampler<3>(host_ptr, Sampler, Coord, ExpectedValue,
259248
PrecisionInULP);
260249
}
261250
{
262-
s::cl_float4 Coord(0.0f, 0.2f, 0.5f, 0.0f); // In-range
263-
s::cl_float4 ExpectedValue =
264-
s::cl_float4(36.8f, 37.8f, 38.8f, 39.8f) / 127.0f;
251+
s::float4 Coord(0.0f, 0.2f, 0.5f, 0.0f); // In-range
252+
s::float4 ExpectedValue = s::float4(36.8f, 37.8f, 38.8f, 39.8f) / 127.0f;
265253
auto Sampler = s::sampler(s::coordinate_normalization_mode::normalized,
266254
s::addressing_mode::clamp_to_edge,
267255
s::filtering_mode::linear);
@@ -272,8 +260,8 @@ void checkSamplerLinear() {
272260
// addressing_mode::clamp
273261
{
274262
// Out-of-range
275-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
276-
s::cl_float4 ExpectedValue = s::cl_float4(0.0f, 0.0f, 0.0f, 0.0f);
263+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
264+
s::float4 ExpectedValue = s::float4(0.0f, 0.0f, 0.0f, 0.0f);
277265
auto Sampler =
278266
s::sampler(s::coordinate_normalization_mode::normalized,
279267
s::addressing_mode::clamp, s::filtering_mode::linear);
@@ -282,9 +270,8 @@ void checkSamplerLinear() {
282270
}
283271
{
284272
// In-range
285-
s::cl_float4 Coord(0.0f, 0.2f, 0.5f, 0.0f);
286-
s::cl_float4 ExpectedValue =
287-
s::cl_float4(18.4f, 18.9f, 19.4f, 19.9f) / 127.0f;
273+
s::float4 Coord(0.0f, 0.2f, 0.5f, 0.0f);
274+
s::float4 ExpectedValue = s::float4(18.4f, 18.9f, 19.4f, 19.9f) / 127.0f;
288275
auto Sampler =
289276
s::sampler(s::coordinate_normalization_mode::normalized,
290277
s::addressing_mode::clamp, s::filtering_mode::linear);
@@ -295,9 +282,8 @@ void checkSamplerLinear() {
295282
// addressing_mode::none
296283
{
297284
// In-range for consistent return value.
298-
s::cl_float4 Coord(0.5f, 0.5f, 0.5f, 0.0f);
299-
s::cl_float4 ExpectedValue =
300-
s::cl_float4(46.0f, 47.0f, 48.0f, 49.0f) / 127.0f;
285+
s::float4 Coord(0.5f, 0.5f, 0.5f, 0.0f);
286+
s::float4 ExpectedValue = s::float4(46.0f, 47.0f, 48.0f, 49.0f) / 127.0f;
301287
auto Sampler =
302288
s::sampler(s::coordinate_normalization_mode::normalized,
303289
s::addressing_mode::none, s::filtering_mode::linear);
@@ -309,9 +295,8 @@ void checkSamplerLinear() {
309295
// addressing_mode::clamp_to_edge
310296
{
311297
// Out-of-range
312-
s::cl_float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
313-
s::cl_float4 ExpectedValue =
314-
s::cl_float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
298+
s::float4 Coord(0.0f, 1.5f, 2.5f, 0.0f);
299+
s::float4 ExpectedValue = s::float4(56.0f, 57.0f, 58.0f, 59.0f) / 127.0f;
315300
auto Sampler = s::sampler(s::coordinate_normalization_mode::unnormalized,
316301
s::addressing_mode::clamp_to_edge,
317302
s::filtering_mode::linear);
@@ -320,8 +305,8 @@ void checkSamplerLinear() {
320305
}
321306
{
322307
// In-range
323-
s::cl_float4 Coord(0.0f, 0.2f, 0.5f, 0.0f);
324-
s::cl_float4 ExpectedValue = s::cl_float4(0.0f, 1.0f, 2.0f, 3.0f) / 127.0f;
308+
s::float4 Coord(0.0f, 0.2f, 0.5f, 0.0f);
309+
s::float4 ExpectedValue = s::float4(0.0f, 1.0f, 2.0f, 3.0f) / 127.0f;
325310
auto Sampler = s::sampler(s::coordinate_normalization_mode::unnormalized,
326311
s::addressing_mode::clamp_to_edge,
327312
s::filtering_mode::linear);
@@ -332,9 +317,8 @@ void checkSamplerLinear() {
332317
// addressing_mode::clamp
333318
{
334319
// Out-of-range
335-
s::cl_float4 Coord(0.0f, 1.5f, 1.5f, 0.0f);
336-
s::cl_float4 ExpectedValue =
337-
s::cl_float4(16.0f, 16.5f, 17.0f, 17.5f) / 127.0f;
320+
s::float4 Coord(0.0f, 1.5f, 1.5f, 0.0f);
321+
s::float4 ExpectedValue = s::float4(16.0f, 16.5f, 17.0f, 17.5f) / 127.0f;
338322
auto Sampler =
339323
s::sampler(s::coordinate_normalization_mode::unnormalized,
340324
s::addressing_mode::clamp, s::filtering_mode::linear);
@@ -343,9 +327,8 @@ void checkSamplerLinear() {
343327
}
344328
{
345329
// In-range
346-
s::cl_float4 Coord(0.0f, 0.2f, 0.5f, 0.0f);
347-
s::cl_float4 ExpectedValue =
348-
s::cl_float4(0.0f, 0.35f, 0.7f, 1.05f) / 127.0f;
330+
s::float4 Coord(0.0f, 0.2f, 0.5f, 0.0f);
331+
s::float4 ExpectedValue = s::float4(0.0f, 0.35f, 0.7f, 1.05f) / 127.0f;
349332
auto Sampler =
350333
s::sampler(s::coordinate_normalization_mode::unnormalized,
351334
s::addressing_mode::clamp, s::filtering_mode::linear);
@@ -356,9 +339,8 @@ void checkSamplerLinear() {
356339
// addressing_mode::none
357340
{
358341
// In-range for consistent return value.
359-
s::cl_float4 Coord(1.0f, 2.0f, 3.0f, 0.0f);
360-
s::cl_float4 ExpectedValue =
361-
s::cl_float4(74.0f, 75.0f, 76.0f, 77.0f) / 127.0f;
342+
s::float4 Coord(1.0f, 2.0f, 3.0f, 0.0f);
343+
s::float4 ExpectedValue = s::float4(74.0f, 75.0f, 76.0f, 77.0f) / 127.0f;
362344
auto Sampler =
363345
s::sampler(s::coordinate_normalization_mode::unnormalized,
364346
s::addressing_mode::none, s::filtering_mode::linear);
@@ -369,8 +351,8 @@ void checkSamplerLinear() {
369351

370352
int main() {
371353

372-
// Note: Currently these functions only check for cl_float4 return datatype,
373-
// the test case can be extended to test all return datatypes.
354+
// Note: Currently these functions only check for vec<float, 4> return
355+
// datatype, the test case can be extended to test all return datatypes.
374356
checkSamplerNearest();
375357
checkSamplerLinear();
376358
}

0 commit comments

Comments
 (0)