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