1 //
2 // Copyright (c) 2020 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16
17 #include <algorithm>
18 #include <numeric>
19 #include <string>
20 #include <vector>
21
22 #include "procs.h"
23 #include "harness/testHarness.h"
24
25 static std::string pragma_extension;
26
27 template <int N> struct TestInfo
28 {
29 };
30
31 template <> struct TestInfo<2>
32 {
33 static const size_t vector_size = 2;
34
35 static constexpr const char* kernel_source_xyzw = R"CLC(
36 __kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
37 int index = 0;
38
39 // lvalue swizzles
40 dst[index++].x = value.x;
41 dst[index++].y = value.x;
42 dst[index++].xy = value;
43 dst[index++].yx = value;
44
45 // rvalue swizzles
46 dst[index++] = value.x;
47 dst[index++] = value.y;
48 dst[index++] = value.xy;
49 dst[index++] = value.yx;
50 }
51 )CLC";
52
53 static constexpr const char* kernel_source_rgba = R"CLC(
54 __kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
55 int index = 0;
56
57 // lvalue swizzles
58 dst[index++].r = value.r;
59 dst[index++].g = value.r;
60 dst[index++].rg = value;
61 dst[index++].gr = value;
62
63 // rvalue swizzles
64 dst[index++] = value.r;
65 dst[index++] = value.g;
66 dst[index++] = value.rg;
67 dst[index++] = value.gr;
68 }
69 )CLC";
70
71 static constexpr const char* kernel_source_sN = R"CLC(
72 __kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
73 int index = 0;
74
75 // lvalue swizzles
76 dst[index++].s0 = value.s0;
77 dst[index++].s1 = value.s0;
78 dst[index++].s01 = value;
79 dst[index++].s10 = value;
80
81 // rvalue swizzles
82 dst[index++] = value.s0;
83 dst[index++] = value.s1;
84 dst[index++] = value.s01;
85 dst[index++] = value.s10;
86 }
87 )CLC";
88 };
89
90 template <> struct TestInfo<3>
91 {
92 static const size_t vector_size = 4; // sizeof(vec3) is four elements
93
94 static constexpr const char* kernel_source_xyzw = R"CLC(
95 __kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
96 int index = 0;
97
98 // lvalue swizzles
99 TYPE t;
100 t = dst[index]; t.x = value.x;
101 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
102 t = dst[index]; t.y = value.x;
103 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
104 t = dst[index]; t.z = value.x;
105 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
106 t = dst[index]; t.xyz = value;
107 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
108 t = dst[index]; t.zyx = value;
109 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
110
111 // rvalue swizzles
112 vstore3(value.x, 0, (__global BASETYPE*)(dst + index++));
113 vstore3(value.y, 0, (__global BASETYPE*)(dst + index++));
114 vstore3(value.z, 0, (__global BASETYPE*)(dst + index++));
115 vstore3(value.xyz, 0, (__global BASETYPE*)(dst + index++));
116 vstore3(value.zyx, 0, (__global BASETYPE*)(dst + index++));
117 }
118 )CLC";
119
120 static constexpr const char* kernel_source_rgba = R"CLC(
121 __kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
122 int index = 0;
123
124 // lvalue swizzles
125 TYPE t;
126 t = dst[index]; t.r = value.r;
127 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
128 t = dst[index]; t.g = value.r;
129 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
130 t = dst[index]; t.b = value.r;
131 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
132 t = dst[index]; t.rgb = value;
133 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
134 t = dst[index]; t.bgr = value;
135 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
136
137 // rvalue swizzles
138 vstore3(value.r, 0, (__global BASETYPE*)(dst + index++));
139 vstore3(value.g, 0, (__global BASETYPE*)(dst + index++));
140 vstore3(value.b, 0, (__global BASETYPE*)(dst + index++));
141 vstore3(value.rgb, 0, (__global BASETYPE*)(dst + index++));
142 vstore3(value.bgr, 0, (__global BASETYPE*)(dst + index++));
143 }
144 )CLC";
145
146 static constexpr const char* kernel_source_sN = R"CLC(
147 __kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
148 int index = 0;
149
150 // lvalue swizzles
151 TYPE t;
152 t = dst[index]; t.s0 = value.s0;
153 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
154 t = dst[index]; t.s1 = value.s0;
155 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
156 t = dst[index]; t.s2 = value.s0;
157 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
158 t = dst[index]; t.s012 = value;
159 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
160 t = dst[index]; t.s210 = value;
161 vstore3(t, 0, (__global BASETYPE*)(dst + index++));
162
163 // rvalue swizzles
164 vstore3(value.s0, 0, (__global BASETYPE*)(dst + index++));
165 vstore3(value.s1, 0, (__global BASETYPE*)(dst + index++));
166 vstore3(value.s2, 0, (__global BASETYPE*)(dst + index++));
167 vstore3(value.s012, 0, (__global BASETYPE*)(dst + index++));
168 vstore3(value.s210, 0, (__global BASETYPE*)(dst + index++));
169 }
170 )CLC";
171 };
172
173 template <> struct TestInfo<4>
174 {
175 static const size_t vector_size = 4;
176
177 static constexpr const char* kernel_source_xyzw = R"CLC(
178 __kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
179 int index = 0;
180
181 // lvalue swizzles
182 dst[index++].x = value.x;
183 dst[index++].y = value.x;
184 dst[index++].z = value.x;
185 dst[index++].w = value.x;
186 dst[index++].xyzw = value;
187 dst[index++].wzyx = value;
188
189 // rvalue swizzles
190 dst[index++] = value.x;
191 dst[index++] = value.y;
192 dst[index++] = value.z;
193 dst[index++] = value.w;
194 dst[index++] = value.xyzw;
195 dst[index++] = value.wzyx;
196 }
197 )CLC";
198
199 static constexpr const char* kernel_source_rgba = R"CLC(
200 __kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
201 int index = 0;
202
203 // lvalue swizzles
204 dst[index++].r = value.r;
205 dst[index++].g = value.r;
206 dst[index++].b = value.r;
207 dst[index++].a = value.r;
208 dst[index++].rgba = value;
209 dst[index++].abgr = value;
210
211 // rvalue swizzles
212 dst[index++] = value.r;
213 dst[index++] = value.g;
214 dst[index++] = value.b;
215 dst[index++] = value.a;
216 dst[index++] = value.rgba;
217 dst[index++] = value.abgr;
218 }
219 )CLC";
220
221 static constexpr const char* kernel_source_sN = R"CLC(
222 __kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
223 int index = 0;
224
225 // lvalue swizzles
226 dst[index++].s0 = value.s0;
227 dst[index++].s1 = value.s0;
228 dst[index++].s2 = value.s0;
229 dst[index++].s3 = value.s0;
230 dst[index++].s0123 = value;
231 dst[index++].s3210 = value;
232
233 // rvalue swizzles
234 dst[index++] = value.s0;
235 dst[index++] = value.s1;
236 dst[index++] = value.s2;
237 dst[index++] = value.s3;
238 dst[index++] = value.s0123;
239 dst[index++] = value.s3210;
240 }
241 )CLC";
242 };
243
244 template <> struct TestInfo<8>
245 {
246 static const size_t vector_size = 8;
247
248 static constexpr const char* kernel_source_xyzw = R"CLC(
249 __kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
250 int index = 0;
251
252 // xwzw only for first four components!
253
254 // lvalue swizzles
255 dst[index++].x = value.x;
256 dst[index++].y = value.x;
257 dst[index++].z = value.x;
258 dst[index++].w = value.x;
259 dst[index++].s4 = value.s0;
260 dst[index++].s5 = value.s0;
261 dst[index++].s6 = value.s0;
262 dst[index++].s7 = value.s0;
263 dst[index].xyzw = value.s0123;
264 dst[index++].s4567 = value.s4567;
265 dst[index].s7654 = value.s0123;
266 dst[index++].wzyx = value.s4567;
267
268 // rvalue swizzles
269 dst[index++] = value.x;
270 dst[index++] = value.y;
271 dst[index++] = value.z;
272 dst[index++] = value.w;
273 dst[index++] = value.s4;
274 dst[index++] = value.s5;
275 dst[index++] = value.s6;
276 dst[index++] = value.s7;
277 dst[index++] = (TYPE)(value.xyzw, value.s4567);
278 dst[index++] = (TYPE)(value.s7654, value.wzyx);
279 }
280 )CLC";
281 static constexpr const char* kernel_source_rgba = R"CLC(
282 __kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
283 int index = 0;
284
285 // rgba only for first four components!
286
287 // lvalue swizzles
288 dst[index++].r = value.r;
289 dst[index++].g = value.r;
290 dst[index++].b = value.r;
291 dst[index++].a = value.r;
292 dst[index++].s4 = value.s0;
293 dst[index++].s5 = value.s0;
294 dst[index++].s6 = value.s0;
295 dst[index++].s7 = value.s0;
296 dst[index].rgba = value.s0123;
297 dst[index++].s4567 = value.s4567;
298 dst[index].s7654 = value.s0123;
299 dst[index++].abgr = value.s4567;
300
301 // rvalue swizzles
302 dst[index++] = value.r;
303 dst[index++] = value.g;
304 dst[index++] = value.b;
305 dst[index++] = value.a;
306 dst[index++] = value.s4;
307 dst[index++] = value.s5;
308 dst[index++] = value.s6;
309 dst[index++] = value.s7;
310 dst[index++] = (TYPE)(value.rgba, value.s4567);
311 dst[index++] = (TYPE)(value.s7654, value.abgr);
312 }
313 )CLC";
314 static constexpr const char* kernel_source_sN = R"CLC(
315 __kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
316 int index = 0;
317
318 // lvalue swizzles
319 dst[index++].s0 = value.s0;
320 dst[index++].s1 = value.s0;
321 dst[index++].s2 = value.s0;
322 dst[index++].s3 = value.s0;
323 dst[index++].s4 = value.s0;
324 dst[index++].s5 = value.s0;
325 dst[index++].s6 = value.s0;
326 dst[index++].s7 = value.s0;
327 dst[index++].s01234567 = value;
328 dst[index++].s76543210 = value;
329
330 // rvalue swizzles
331 dst[index++] = value.s0;
332 dst[index++] = value.s1;
333 dst[index++] = value.s2;
334 dst[index++] = value.s3;
335 dst[index++] = value.s4;
336 dst[index++] = value.s5;
337 dst[index++] = value.s6;
338 dst[index++] = value.s7;
339 dst[index++] = value.s01234567;
340 dst[index++] = value.s76543210;
341 }
342 )CLC";
343 };
344
345 template <> struct TestInfo<16>
346 {
347 static const size_t vector_size = 16;
348
349 static constexpr const char* kernel_source_xyzw = R"CLC(
350 __kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
351 int index = 0;
352
353 // xwzw only for first four components!
354
355 // lvalue swizzles
356 dst[index++].x = value.x;
357 dst[index++].y = value.x;
358 dst[index++].z = value.x;
359 dst[index++].w = value.x;
360 dst[index++].s4 = value.s0;
361 dst[index++].s5 = value.s0;
362 dst[index++].s6 = value.s0;
363 dst[index++].s7 = value.s0;
364 dst[index++].s8 = value.s0;
365 dst[index++].s9 = value.s0;
366 dst[index++].sa = value.s0;
367 dst[index++].sb = value.s0;
368 dst[index++].sc = value.s0;
369 dst[index++].sd = value.s0;
370 dst[index++].se = value.s0;
371 dst[index++].sf = value.s0;
372 dst[index].xyzw = value.s0123;
373 dst[index].s4567 = value.s4567;
374 dst[index].s89ab = value.s89ab;
375 dst[index++].scdef = value.scdef;
376 dst[index].sfedc = value.s0123;
377 dst[index].sba98 = value.s4567;
378 dst[index].s7654 = value.s89ab;
379 dst[index++].wzyx = value.scdef;
380
381 // rvalue swizzles
382 dst[index++] = value.x;
383 dst[index++] = value.y;
384 dst[index++] = value.z;
385 dst[index++] = value.w;
386 dst[index++] = value.s4;
387 dst[index++] = value.s5;
388 dst[index++] = value.s6;
389 dst[index++] = value.s7;
390 dst[index++] = value.s8;
391 dst[index++] = value.s9;
392 dst[index++] = value.sa;
393 dst[index++] = value.sb;
394 dst[index++] = value.sc;
395 dst[index++] = value.sd;
396 dst[index++] = value.se;
397 dst[index++] = value.sf;
398 dst[index++] = (TYPE)(value.xyzw, value.s4567, value.s89abcdef);
399 dst[index++] = (TYPE)(value.sfedcba98, value.s7654, value.wzyx);
400 }
401 )CLC";
402 static constexpr const char* kernel_source_rgba = R"CLC(
403 __kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
404 int index = 0;
405
406 // rgba only for first four components!
407
408 // lvalue swizzles
409 dst[index++].r = value.r;
410 dst[index++].g = value.r;
411 dst[index++].b = value.r;
412 dst[index++].a = value.r;
413 dst[index++].s4 = value.s0;
414 dst[index++].s5 = value.s0;
415 dst[index++].s6 = value.s0;
416 dst[index++].s7 = value.s0;
417 dst[index++].s8 = value.s0;
418 dst[index++].s9 = value.s0;
419 dst[index++].sa = value.s0;
420 dst[index++].sb = value.s0;
421 dst[index++].sc = value.s0;
422 dst[index++].sd = value.s0;
423 dst[index++].se = value.s0;
424 dst[index++].sf = value.s0;
425 dst[index].rgba = value.s0123;
426 dst[index].s4567 = value.s4567;
427 dst[index].s89ab = value.s89ab;
428 dst[index++].scdef = value.scdef;
429 dst[index].sfedc = value.s0123;
430 dst[index].sba98 = value.s4567;
431 dst[index].s7654 = value.s89ab;
432 dst[index++].abgr = value.scdef;
433
434 // rvalue swizzles
435 dst[index++] = value.r;
436 dst[index++] = value.g;
437 dst[index++] = value.b;
438 dst[index++] = value.a;
439 dst[index++] = value.s4;
440 dst[index++] = value.s5;
441 dst[index++] = value.s6;
442 dst[index++] = value.s7;
443 dst[index++] = value.s8;
444 dst[index++] = value.s9;
445 dst[index++] = value.sa;
446 dst[index++] = value.sb;
447 dst[index++] = value.sc;
448 dst[index++] = value.sd;
449 dst[index++] = value.se;
450 dst[index++] = value.sf;
451 dst[index++] = (TYPE)(value.rgba, value.s4567, value.s89abcdef);
452 dst[index++] = (TYPE)(value.sfedcba98, value.s7654, value.abgr);
453 }
454 )CLC";
455 static constexpr const char* kernel_source_sN = R"CLC(
456 __kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
457 int index = 0;
458
459 // lvalue swizzles
460 dst[index++].s0 = value.s0;
461 dst[index++].s1 = value.s0;
462 dst[index++].s2 = value.s0;
463 dst[index++].s3 = value.s0;
464 dst[index++].s4 = value.s0;
465 dst[index++].s5 = value.s0;
466 dst[index++].s6 = value.s0;
467 dst[index++].s7 = value.s0;
468 dst[index++].s8 = value.s0;
469 dst[index++].s9 = value.s0;
470 dst[index++].sa = value.s0;
471 dst[index++].sb = value.s0;
472 dst[index++].sc = value.s0;
473 dst[index++].sd = value.s0;
474 dst[index++].se = value.s0;
475 dst[index++].sf = value.s0;
476 dst[index++].s0123456789abcdef = value; // lower-case
477 dst[index++].sFEDCBA9876543210 = value; // upper-case
478
479 // rvalue swizzles
480 dst[index++] = value.s0;
481 dst[index++] = value.s1;
482 dst[index++] = value.s2;
483 dst[index++] = value.s3;
484 dst[index++] = value.s4;
485 dst[index++] = value.s5;
486 dst[index++] = value.s6;
487 dst[index++] = value.s7;
488 dst[index++] = value.s8;
489 dst[index++] = value.s9;
490 dst[index++] = value.sa;
491 dst[index++] = value.sb;
492 dst[index++] = value.sc;
493 dst[index++] = value.sd;
494 dst[index++] = value.se;
495 dst[index++] = value.sf;
496 dst[index++] = value.s0123456789abcdef; // lower-case
497 dst[index++] = value.sFEDCBA9876543210; // upper-case
498 }
499 )CLC";
500 };
501
502 template <typename T, size_t N, size_t S>
makeReference(std::vector<T> & ref)503 static void makeReference(std::vector<T>& ref)
504 {
505 // N single channel lvalue tests
506 // 2 multi-value lvalue tests
507 // N single channel rvalue tests
508 // 2 multi-value rvalue tests
509 const size_t refSize = (N + 2 + N + 2) * S;
510
511 ref.resize(refSize);
512 std::fill(ref.begin(), ref.end(), 99);
513
514 size_t dstIndex = 0;
515
516 // single channel lvalue
517 for (size_t i = 0; i < N; i++)
518 {
519 ref[dstIndex * S + i] = 0;
520 ++dstIndex;
521 }
522
523 // normal lvalue
524 for (size_t c = 0; c < N; c++)
525 {
526 ref[dstIndex * S + c] = c;
527 }
528 ++dstIndex;
529
530 // reverse lvalue
531 for (size_t c = 0; c < N; c++)
532 {
533 ref[dstIndex * S + c] = N - c - 1;
534 }
535 ++dstIndex;
536
537 // single channel rvalue
538 for (size_t i = 0; i < N; i++)
539 {
540 for (size_t c = 0; c < N; c++)
541 {
542 ref[dstIndex * S + c] = i;
543 }
544 ++dstIndex;
545 }
546
547 // normal rvalue
548 for (size_t c = 0; c < N; c++)
549 {
550 ref[dstIndex * S + c] = c;
551 }
552 ++dstIndex;
553
554 // reverse rvalue
555 for (size_t c = 0; c < N; c++)
556 {
557 ref[dstIndex * S + c] = N - c - 1;
558 }
559 ++dstIndex;
560
561 assert(dstIndex * S == refSize);
562 }
563
564 template <typename T>
565 static int
test_vectype_case(const std::vector<T> & value,const std::vector<T> & reference,cl_context context,cl_kernel kernel,cl_command_queue queue)566 test_vectype_case(const std::vector<T>& value, const std::vector<T>& reference,
567 cl_context context, cl_kernel kernel, cl_command_queue queue)
568 {
569 cl_int error = CL_SUCCESS;
570
571 clMemWrapper mem;
572
573 std::vector<T> buffer(reference.size(), 99);
574 mem = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
575 buffer.size() * sizeof(T), buffer.data(), &error);
576 test_error(error, "Unable to create test buffer");
577
578 error = clSetKernelArg(kernel, 0, value.size() * sizeof(T), value.data());
579 test_error(error, "Unable to set value kernel arg");
580
581 error = clSetKernelArg(kernel, 1, sizeof(mem), &mem);
582 test_error(error, "Unable to set destination buffer kernel arg");
583
584 size_t global_work_size[] = { 1 };
585 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size,
586 NULL, 0, NULL, NULL);
587 test_error(error, "Unable to enqueue test kernel");
588
589 error = clFinish(queue);
590 test_error(error, "clFinish failed after test kernel");
591
592 error =
593 clEnqueueReadBuffer(queue, mem, CL_TRUE, 0, buffer.size() * sizeof(T),
594 buffer.data(), 0, NULL, NULL);
595 test_error(error, "Unable to read data after test kernel");
596
597 if (buffer != reference)
598 {
599 log_error("Result buffer did not match reference buffer!\n");
600 return TEST_FAIL;
601 }
602
603 return TEST_PASS;
604 }
605
606 template <typename T, size_t N>
test_vectype(const char * type_name,cl_device_id device,cl_context context,cl_command_queue queue)607 static int test_vectype(const char* type_name, cl_device_id device,
608 cl_context context, cl_command_queue queue)
609 {
610 log_info(" testing type %s%d\n", type_name, N);
611
612 cl_int error = CL_SUCCESS;
613 int result = TEST_PASS;
614
615 std::string buildOptions{ "-DTYPE=" };
616 buildOptions += type_name;
617 buildOptions += std::to_string(N);
618 buildOptions += " -DBASETYPE=";
619 buildOptions += type_name;
620
621 constexpr size_t S = TestInfo<N>::vector_size;
622
623 std::vector<T> value(S);
624 std::iota(value.begin(), value.end(), 0);
625
626 std::vector<T> reference;
627 makeReference<T, N, S>(reference);
628
629 // XYZW swizzles:
630 {
631 clProgramWrapper program;
632 clKernelWrapper kernel;
633
634 std::string program_src =
635 pragma_extension + std::string(TestInfo<N>::kernel_source_xyzw);
636 const char* xyzw_source = program_src.c_str();
637 error = create_single_kernel_helper(
638 context, &program, &kernel, 1, &xyzw_source,
639 "test_vector_swizzle_xyzw", buildOptions.c_str());
640 test_error(error, "Unable to create xyzw test kernel");
641
642 result |= test_vectype_case(value, reference, context, kernel, queue);
643 }
644
645 // sN swizzles:
646 {
647 clProgramWrapper program;
648 clKernelWrapper kernel;
649
650 std::string program_src =
651 pragma_extension + std::string(TestInfo<N>::kernel_source_sN);
652 const char* sN_source = program_src.c_str();
653 error = create_single_kernel_helper(
654 context, &program, &kernel, 1, &sN_source, "test_vector_swizzle_sN",
655 buildOptions.c_str());
656 test_error(error, "Unable to create sN test kernel");
657
658 result |= test_vectype_case(value, reference, context, kernel, queue);
659 }
660
661 // RGBA swizzles for OpenCL 3.0 and newer:
662 {
663 clProgramWrapper program;
664 clKernelWrapper kernel;
665
666 const Version device_version = get_device_cl_version(device);
667 if (device_version >= Version(3, 0))
668 {
669 std::string program_src =
670 pragma_extension + std::string(TestInfo<N>::kernel_source_rgba);
671 const char* rgba_source = program_src.c_str();
672 error = create_single_kernel_helper(
673 context, &program, &kernel, 1, &rgba_source,
674 "test_vector_swizzle_rgba", buildOptions.c_str());
675 test_error(error, "Unable to create rgba test kernel");
676
677 result |=
678 test_vectype_case(value, reference, context, kernel, queue);
679 }
680 }
681
682 return result;
683 }
684
685 template <typename T>
test_type(const char * type_name,cl_device_id device,cl_context context,cl_command_queue queue)686 static int test_type(const char* type_name, cl_device_id device,
687 cl_context context, cl_command_queue queue)
688 {
689 return test_vectype<T, 2>(type_name, device, context, queue)
690 | test_vectype<T, 3>(type_name, device, context, queue)
691 | test_vectype<T, 4>(type_name, device, context, queue)
692 | test_vectype<T, 8>(type_name, device, context, queue)
693 | test_vectype<T, 16>(type_name, device, context, queue);
694 }
695
test_vector_swizzle(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)696 int test_vector_swizzle(cl_device_id device, cl_context context,
697 cl_command_queue queue, int num_elements)
698 {
699 int hasDouble = is_extension_available(device, "cl_khr_fp64");
700 int hasHalf = is_extension_available(device, "cl_khr_fp16");
701
702 int result = TEST_PASS;
703 result |= test_type<cl_char>("char", device, context, queue);
704 result |= test_type<cl_uchar>("uchar", device, context, queue);
705 result |= test_type<cl_short>("short", device, context, queue);
706 result |= test_type<cl_ushort>("ushort", device, context, queue);
707 result |= test_type<cl_int>("int", device, context, queue);
708 result |= test_type<cl_uint>("uint", device, context, queue);
709 if (gHasLong)
710 {
711 result |= test_type<cl_long>("long", device, context, queue);
712 result |= test_type<cl_ulong>("ulong", device, context, queue);
713 }
714 result |= test_type<cl_float>("float", device, context, queue);
715 if (hasHalf)
716 {
717 pragma_extension = "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n";
718 result |= test_type<cl_half>("half", device, context, queue);
719 }
720 if (hasDouble)
721 {
722 pragma_extension = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
723 result |= test_type<cl_double>("double", device, context, queue);
724 }
725 return result;
726 }
727