1 //
2 // Copyright (c) 2022 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 <cstdint>
18 #include <functional>
19 #include <iostream>
20 #include <map>
21 #include <memory>
22 #include <stdexcept>
23 #include <vector>
24
25 #include "harness/stringHelpers.h"
26
27 #include <CL/cl_half.h>
28
29 #include "test_comparisons_fp.h"
30
31 #define TEST_SIZE 512
32
33 static char ftype[32] = { 0 };
34 static char ftype_vec[32] = { 0 };
35 static char itype[32] = { 0 };
36 static char itype_vec[32] = { 0 };
37 static char extension[128] = { 0 };
38
39 // clang-format off
40 // for readability sake keep this section unformatted
41 const char* equivTestKernPat[] = {
42 extension,
43 "__kernel void sample_test(__global ", ftype_vec, " *sourceA, __global ", ftype_vec,
44 " *sourceB, __global ", itype_vec, " *destValues, __global ", itype_vec, " *destValuesB)\n"
45 "{\n"
46 " int tid = get_global_id(0);\n"
47 " destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n"
48 " destValuesB[tid] = sourceA[tid] %s sourceB[tid];\n"
49 "}\n"};
50
51 const char* equivTestKernPatLessGreater[] = {
52 extension,
53 "__kernel void sample_test(__global ", ftype_vec, " *sourceA, __global ", ftype_vec,
54 " *sourceB, __global ", itype_vec, " *destValues, __global ", itype_vec, " *destValuesB)\n"
55 "{\n"
56 " int tid = get_global_id(0);\n"
57 " destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n"
58 " destValuesB[tid] = (sourceA[tid] < sourceB[tid]) | (sourceA[tid] > sourceB[tid]);\n"
59 "}\n"};
60
61 const char* equivTestKerPat_3[] = {
62 extension,
63 "__kernel void sample_test(__global ", ftype_vec, " *sourceA, __global ", ftype_vec,
64 " *sourceB, __global ", itype_vec, " *destValues, __global ", itype_vec, " *destValuesB)\n"
65 "{\n"
66 " int tid = get_global_id(0);\n"
67 " ",ftype_vec," sampA = vload3(tid, (__global ",ftype," *)sourceA);\n"
68 " ",ftype_vec," sampB = vload3(tid, (__global ",ftype," *)sourceB);\n"
69 " vstore3(%s( sampA, sampB ), tid, (__global ",itype," *)destValues);\n"
70 " vstore3(( sampA %s sampB ), tid, (__global ",itype," *)destValuesB);\n"
71 "}\n"};
72
73 const char* equivTestKerPatLessGreater_3[] = {
74 extension,
75 "__kernel void sample_test(__global ", ftype_vec, " *sourceA, __global ", ftype_vec,
76 " *sourceB, __global ", itype_vec, " *destValues, __global ", itype_vec, " *destValuesB)\n"
77 "{\n"
78 " int tid = get_global_id(0);\n"
79 " ", ftype_vec, " sampA = vload3(tid, (__global ", ftype, " *)sourceA);\n"
80 " ", ftype_vec, " sampB = vload3(tid, (__global ", ftype, " *)sourceB);\n"
81 " vstore3(%s( sampA, sampB ), tid, (__global ", itype, " *)destValues);\n"
82 " vstore3(( sampA < sampB ) | (sampA > sampB), tid, (__global ", itype, " *)destValuesB);\n"
83 "}\n"
84 };
85 // clang-format on
86
87
verify(const T & A,const T & B)88 template <typename T, typename F> bool verify(const T& A, const T& B)
89 {
90 return F()(A, B);
91 }
92
RelationalsFPTest(cl_context context,cl_device_id device,cl_command_queue queue,const char * fn,const char * op)93 RelationalsFPTest::RelationalsFPTest(cl_context context, cl_device_id device,
94 cl_command_queue queue, const char* fn,
95 const char* op)
96 : context(context), device(device), queue(queue), fnName(fn), opName(op),
97 halfFlushDenormsToZero(0)
98 {
99 // hardcoded for now, to be changed into typeid().name solution in future
100 // for now C++ spec doesn't guarantee human readable type name
101
102 eqTypeNames = { { kHalf, "short" },
103 { kFloat, "int" },
104 { kDouble, "long" } };
105 }
106
107 template <typename T>
generate_equiv_test_data(T * outData,unsigned int vecSize,bool alpha,const RelTestParams<T> & param,const MTdata & d)108 void RelationalsFPTest::generate_equiv_test_data(T* outData,
109 unsigned int vecSize,
110 bool alpha,
111 const RelTestParams<T>& param,
112 const MTdata& d)
113 {
114 unsigned int i;
115
116 generate_random_data(param.dataType, vecSize * TEST_SIZE, d, outData);
117
118 // Fill the first few vectors with NAN in each vector element (or the second
119 // set if we're alpha, so we can test either case)
120 if (alpha) outData += vecSize * vecSize;
121 for (i = 0; i < vecSize; i++)
122 {
123 outData[0] = param.nan;
124 outData += vecSize + 1;
125 }
126 // Make sure the third set is filled regardless, to test the case where both
127 // have NANs
128 if (!alpha) outData += vecSize * vecSize;
129 for (i = 0; i < vecSize; i++)
130 {
131 outData[0] = param.nan;
132 outData += vecSize + 1;
133 }
134 }
135
136 template <typename T, typename U>
verify_equiv_values(unsigned int vecSize,const T * const inDataA,const T * const inDataB,U * const outData,const VerifyFunc<T> & verifyFn)137 void RelationalsFPTest::verify_equiv_values(unsigned int vecSize,
138 const T* const inDataA,
139 const T* const inDataB,
140 U* const outData,
141 const VerifyFunc<T>& verifyFn)
142 {
143 unsigned int i;
144 int trueResult;
145 bool result;
146
147 trueResult = (vecSize == 1) ? 1 : -1;
148 for (i = 0; i < vecSize; i++)
149 {
150 result = verifyFn(inDataA[i], inDataB[i]);
151 outData[i] = result ? trueResult : 0;
152 }
153 }
154
155 template <typename T>
test_equiv_kernel(unsigned int vecSize,const RelTestParams<T> & param,const MTdata & d)156 int RelationalsFPTest::test_equiv_kernel(unsigned int vecSize,
157 const RelTestParams<T>& param,
158 const MTdata& d)
159 {
160 clProgramWrapper program;
161 clKernelWrapper kernel;
162 clMemWrapper streams[4];
163 T inDataA[TEST_SIZE * 16], inDataB[TEST_SIZE * 16];
164
165 // support half, float, double equivalents - otherwise assert
166 typedef typename std::conditional<
167 (sizeof(T) == sizeof(std::int16_t)), std::int16_t,
168 typename std::conditional<(sizeof(T) == sizeof(std::int32_t)),
169 std::int32_t, std::int64_t>::type>::type U;
170
171 U outData[TEST_SIZE * 16], expected[16];
172 int error, i, j;
173 size_t threads[1], localThreads[1];
174 std::string kernelSource;
175 char sizeName[4];
176
177 /* Create the source */
178 if (vecSize == 1)
179 sizeName[0] = 0;
180 else
181 sprintf(sizeName, "%d", vecSize);
182
183 if (eqTypeNames.find(param.dataType) == eqTypeNames.end())
184 log_error(
185 "RelationalsFPTest::test_equiv_kernel: unsupported fp data type");
186
187 sprintf(ftype, "%s", get_explicit_type_name(param.dataType));
188 sprintf(ftype_vec, "%s%s", get_explicit_type_name(param.dataType),
189 sizeName);
190
191 sprintf(itype, "%s", eqTypeNames[param.dataType].c_str());
192 sprintf(itype_vec, "%s%s", eqTypeNames[param.dataType].c_str(), sizeName);
193
194 if (std::is_same<T, double>::value)
195 strcpy(extension, "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n");
196 else if (std::is_same<T, cl_half>::value)
197 strcpy(extension, "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n");
198 else
199 extension[0] = '\0';
200
201 if (DENSE_PACK_VECS && vecSize == 3)
202 {
203 if (strcmp(fnName.c_str(), "islessgreater"))
204 {
205 auto str =
206 concat_kernel(equivTestKerPat_3,
207 sizeof(equivTestKerPat_3) / sizeof(const char*));
208 kernelSource = str_sprintf(str, fnName.c_str(), opName.c_str());
209 }
210 else
211 {
212 auto str = concat_kernel(equivTestKerPatLessGreater_3,
213 sizeof(equivTestKerPatLessGreater_3)
214 / sizeof(const char*));
215 kernelSource = str_sprintf(str, fnName.c_str());
216 }
217 }
218 else
219 {
220 if (strcmp(fnName.c_str(), "islessgreater"))
221 {
222 auto str =
223 concat_kernel(equivTestKernPat,
224 sizeof(equivTestKernPat) / sizeof(const char*));
225 kernelSource = str_sprintf(str, fnName.c_str(), opName.c_str());
226 }
227 else
228 {
229 auto str = concat_kernel(equivTestKernPatLessGreater,
230 sizeof(equivTestKernPatLessGreater)
231 / sizeof(const char*));
232 kernelSource = str_sprintf(str, fnName.c_str());
233 }
234 }
235
236 /* Create kernels */
237 const char* programPtr = kernelSource.c_str();
238 if (create_single_kernel_helper(context, &program, &kernel, 1,
239 (const char**)&programPtr, "sample_test"))
240 {
241 return -1;
242 }
243
244 /* Generate some streams */
245 generate_equiv_test_data<T>(inDataA, vecSize, true, param, d);
246 generate_equiv_test_data<T>(inDataB, vecSize, false, param, d);
247
248 streams[0] =
249 clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
250 sizeof(T) * vecSize * TEST_SIZE, &inDataA, &error);
251 if (streams[0] == NULL)
252 {
253 print_error(error, "Creating input array A failed!\n");
254 return -1;
255 }
256 streams[1] =
257 clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
258 sizeof(T) * vecSize * TEST_SIZE, &inDataB, &error);
259 if (streams[1] == NULL)
260 {
261 print_error(error, "Creating input array A failed!\n");
262 return -1;
263 }
264 streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
265 sizeof(U) * vecSize * TEST_SIZE, NULL, &error);
266 if (streams[2] == NULL)
267 {
268 print_error(error, "Creating output array failed!\n");
269 return -1;
270 }
271 streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE,
272 sizeof(U) * vecSize * TEST_SIZE, NULL, &error);
273 if (streams[3] == NULL)
274 {
275 print_error(error, "Creating output array failed!\n");
276 return -1;
277 }
278
279 /* Assign streams and execute */
280 error = clSetKernelArg(kernel, 0, sizeof(streams[0]), &streams[0]);
281 test_error(error, "Unable to set indexed kernel arguments");
282 error = clSetKernelArg(kernel, 1, sizeof(streams[1]), &streams[1]);
283 test_error(error, "Unable to set indexed kernel arguments");
284 error = clSetKernelArg(kernel, 2, sizeof(streams[2]), &streams[2]);
285 test_error(error, "Unable to set indexed kernel arguments");
286 error = clSetKernelArg(kernel, 3, sizeof(streams[3]), &streams[3]);
287 test_error(error, "Unable to set indexed kernel arguments");
288
289 /* Run the kernel */
290 threads[0] = TEST_SIZE;
291
292 error = get_max_common_work_group_size(context, kernel, threads[0],
293 &localThreads[0]);
294 test_error(error, "Unable to get work group size to use");
295
296 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads,
297 localThreads, 0, NULL, NULL);
298 test_error(error, "Unable to execute test kernel");
299
300 /* Now get the results */
301 error = clEnqueueReadBuffer(queue, streams[2], true, 0,
302 sizeof(U) * TEST_SIZE * vecSize, outData, 0,
303 NULL, NULL);
304 test_error(error, "Unable to read output array!");
305
306 auto verror_msg = [](const int& i, const int& j, const unsigned& vs,
307 const U& e, const U& o, const T& iA, const T& iB) {
308 std::stringstream sstr;
309 sstr << "ERROR: Data sample " << i << ":" << j << " at size " << vs
310 << " does not validate! Expected " << e << ", got " << o
311 << ", source " << iA << ":" << iB << std::endl;
312 log_error(sstr.str().c_str());
313 };
314
315 /* And verify! */
316 for (i = 0; i < TEST_SIZE; i++)
317 {
318 verify_equiv_values<T, U>(vecSize, &inDataA[i * vecSize],
319 &inDataB[i * vecSize], expected,
320 param.verifyFn);
321
322 for (j = 0; j < (int)vecSize; j++)
323 {
324 if (expected[j] != outData[i * vecSize + j])
325 {
326 bool acceptFail = true;
327 if (std::is_same<T, cl_half>::value)
328 {
329 bool in_denorm = IsHalfSubnormal(inDataA[i * vecSize + j])
330 || IsHalfSubnormal(inDataB[i * vecSize + j]);
331
332 if (halfFlushDenormsToZero && in_denorm)
333 {
334 acceptFail = false;
335 }
336 }
337
338 if (acceptFail)
339 {
340 verror_msg(
341 i, j, vecSize, expected[j], outData[i * vecSize + j],
342 inDataA[i * vecSize + j], inDataB[i * vecSize + j]);
343 return -1;
344 }
345 }
346 }
347 }
348
349 /* Now get the results */
350 error = clEnqueueReadBuffer(queue, streams[3], true, 0,
351 sizeof(U) * TEST_SIZE * vecSize, outData, 0,
352 NULL, NULL);
353 test_error(error, "Unable to read output array!");
354
355 /* And verify! */
356 int fail = 0;
357 for (i = 0; i < TEST_SIZE; i++)
358 {
359 verify_equiv_values<T, U>(vecSize, &inDataA[i * vecSize],
360 &inDataB[i * vecSize], expected,
361 param.verifyFn);
362
363 for (j = 0; j < (int)vecSize; j++)
364 {
365 if (expected[j] != outData[i * vecSize + j])
366 {
367 if (std::is_same<T, float>::value)
368 {
369 if (gInfNanSupport == 0)
370 {
371 if (isnan(inDataA[i * vecSize + j])
372 || isnan(inDataB[i * vecSize + j]))
373 fail = 0;
374 else
375 fail = 1;
376 }
377 if (fail)
378 {
379 verror_msg(i, j, vecSize, expected[j],
380 outData[i * vecSize + j],
381 inDataA[i * vecSize + j],
382 inDataB[i * vecSize + j]);
383 return -1;
384 }
385 }
386 else if (std::is_same<T, cl_half>::value)
387 {
388 bool in_denorm = IsHalfSubnormal(inDataA[i * vecSize + j])
389 || IsHalfSubnormal(inDataB[i * vecSize + j]);
390
391 if (!(halfFlushDenormsToZero && in_denorm))
392 {
393 verror_msg(i, j, vecSize, expected[j],
394 outData[i * vecSize + j],
395 inDataA[i * vecSize + j],
396 inDataB[i * vecSize + j]);
397 return -1;
398 }
399 }
400 else
401 {
402 verror_msg(
403 i, j, vecSize, expected[j], outData[i * vecSize + j],
404 inDataA[i * vecSize + j], inDataB[i * vecSize + j]);
405 return -1;
406 }
407 }
408 }
409 }
410 return 0;
411 }
412
413 template <typename T>
test_relational(int numElements,const RelTestParams<T> & param)414 int RelationalsFPTest::test_relational(int numElements,
415 const RelTestParams<T>& param)
416 {
417 RandomSeed seed(gRandomSeed);
418 unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
419 unsigned int index;
420 int retVal = 0;
421
422 for (index = 0; vecSizes[index] != 0; index++)
423 {
424 // Test!
425 if (test_equiv_kernel<T>(vecSizes[index], param, seed) != 0)
426 {
427 log_error(" Vector %s%d FAILED\n", ftype, vecSizes[index]);
428 retVal = -1;
429 }
430 }
431 return retVal;
432 }
433
SetUp(int elements)434 cl_int RelationalsFPTest::SetUp(int elements)
435 {
436 if (is_extension_available(device, "cl_khr_fp16"))
437 {
438 cl_device_fp_config config = 0;
439 cl_int error = clGetDeviceInfo(device, CL_DEVICE_HALF_FP_CONFIG,
440 sizeof(config), &config, NULL);
441 test_error(error, "Unable to get device CL_DEVICE_HALF_FP_CONFIG");
442
443 halfFlushDenormsToZero = (0 == (config & CL_FP_DENORM));
444 log_info("Supports half precision denormals: %s\n",
445 halfFlushDenormsToZero ? "NO" : "YES");
446 }
447
448 return CL_SUCCESS;
449 }
450
Run()451 cl_int RelationalsFPTest::Run()
452 {
453 cl_int error = CL_SUCCESS;
454 for (auto&& param : params)
455 {
456 switch (param->dataType)
457 {
458 case kHalf:
459 error = test_relational<cl_half>(
460 num_elements, *((RelTestParams<cl_half>*)param.get()));
461 break;
462 case kFloat:
463 error = test_relational<float>(
464 num_elements, *((RelTestParams<float>*)param.get()));
465 break;
466 case kDouble:
467 error = test_relational<double>(
468 num_elements, *((RelTestParams<double>*)param.get()));
469 break;
470 default:
471 test_error(-1, "RelationalsFPTest::Run: incorrect fp type");
472 break;
473 }
474 test_error(error, "RelationalsFPTest::Run: test_relational failed");
475 }
476 return CL_SUCCESS;
477 }
478
SetUp(int elements)479 cl_int IsEqualFPTest::SetUp(int elements)
480 {
481 num_elements = elements;
482 if (is_extension_available(device, "cl_khr_fp16"))
483 params.emplace_back(new RelTestParams<cl_half>(
484 &verify<cl_half, half_equals_to>, kHalf, HALF_NAN));
485
486 params.emplace_back(new RelTestParams<float>(
487 &verify<float, std::equal_to<float>>, kFloat, NAN));
488
489 if (is_extension_available(device, "cl_khr_fp64"))
490 params.emplace_back(new RelTestParams<double>(
491 &verify<double, std::equal_to<double>>, kDouble, NAN));
492
493 return RelationalsFPTest::SetUp(elements);
494 }
495
SetUp(int elements)496 cl_int IsNotEqualFPTest::SetUp(int elements)
497 {
498 num_elements = elements;
499 if (is_extension_available(device, "cl_khr_fp16"))
500 params.emplace_back(new RelTestParams<cl_half>(
501 &verify<cl_half, half_not_equals_to>, kHalf, HALF_NAN));
502
503 params.emplace_back(new RelTestParams<float>(
504 &verify<float, std::not_equal_to<float>>, kFloat, NAN));
505
506 if (is_extension_available(device, "cl_khr_fp64"))
507 params.emplace_back(new RelTestParams<double>(
508 &verify<double, std::not_equal_to<double>>, kDouble, NAN));
509
510 return RelationalsFPTest::SetUp(elements);
511 }
512
SetUp(int elements)513 cl_int IsGreaterFPTest::SetUp(int elements)
514 {
515 num_elements = elements;
516 if (is_extension_available(device, "cl_khr_fp16"))
517 params.emplace_back(new RelTestParams<cl_half>(
518 &verify<cl_half, half_greater>, kHalf, HALF_NAN));
519
520 params.emplace_back(new RelTestParams<float>(
521 &verify<float, std::greater<float>>, kFloat, NAN));
522
523 if (is_extension_available(device, "cl_khr_fp64"))
524 params.emplace_back(new RelTestParams<double>(
525 &verify<double, std::greater<double>>, kDouble, NAN));
526
527 return RelationalsFPTest::SetUp(elements);
528 }
529
SetUp(int elements)530 cl_int IsGreaterEqualFPTest::SetUp(int elements)
531 {
532 num_elements = elements;
533 if (is_extension_available(device, "cl_khr_fp16"))
534 params.emplace_back(new RelTestParams<cl_half>(
535 &verify<cl_half, half_greater_equal>, kHalf, HALF_NAN));
536
537 params.emplace_back(new RelTestParams<float>(
538 &verify<float, std::greater_equal<float>>, kFloat, NAN));
539
540 if (is_extension_available(device, "cl_khr_fp64"))
541 params.emplace_back(new RelTestParams<double>(
542 &verify<double, std::greater_equal<double>>, kDouble, NAN));
543
544 return RelationalsFPTest::SetUp(elements);
545 }
546
SetUp(int elements)547 cl_int IsLessFPTest::SetUp(int elements)
548 {
549 num_elements = elements;
550 if (is_extension_available(device, "cl_khr_fp16"))
551 params.emplace_back(new RelTestParams<cl_half>(
552 &verify<cl_half, half_less>, kHalf, HALF_NAN));
553
554 params.emplace_back(new RelTestParams<float>(
555 &verify<float, std::less<float>>, kFloat, NAN));
556
557 if (is_extension_available(device, "cl_khr_fp64"))
558 params.emplace_back(new RelTestParams<double>(
559 &verify<double, std::less<double>>, kDouble, NAN));
560
561 return RelationalsFPTest::SetUp(elements);
562 }
563
SetUp(int elements)564 cl_int IsLessEqualFPTest::SetUp(int elements)
565 {
566 num_elements = elements;
567 if (is_extension_available(device, "cl_khr_fp16"))
568 params.emplace_back(new RelTestParams<cl_half>(
569 &verify<cl_half, half_less_equal>, kHalf, HALF_NAN));
570
571 params.emplace_back(new RelTestParams<float>(
572 &verify<float, std::less_equal<float>>, kFloat, NAN));
573
574 if (is_extension_available(device, "cl_khr_fp64"))
575 params.emplace_back(new RelTestParams<double>(
576 &verify<double, std::less_equal<double>>, kDouble, NAN));
577
578 return RelationalsFPTest::SetUp(elements);
579 }
580
SetUp(int elements)581 cl_int IsLessGreaterFPTest::SetUp(int elements)
582 {
583 num_elements = elements;
584 if (is_extension_available(device, "cl_khr_fp16"))
585 params.emplace_back(new RelTestParams<cl_half>(
586 &verify<cl_half, half_less_greater>, kHalf, HALF_NAN));
587
588 params.emplace_back(new RelTestParams<float>(
589 &verify<float, less_greater<float>>, kFloat, NAN));
590
591 if (is_extension_available(device, "cl_khr_fp64"))
592 params.emplace_back(new RelTestParams<double>(
593 &verify<double, less_greater<double>>, kDouble, NAN));
594
595 return RelationalsFPTest::SetUp(elements);
596 }
597
test_relational_isequal(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)598 int test_relational_isequal(cl_device_id device, cl_context context,
599 cl_command_queue queue, int numElements)
600 {
601 return MakeAndRunTest<IsEqualFPTest>(device, context, queue, numElements);
602 }
603
test_relational_isnotequal(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)604 int test_relational_isnotequal(cl_device_id device, cl_context context,
605 cl_command_queue queue, int numElements)
606 {
607 return MakeAndRunTest<IsNotEqualFPTest>(device, context, queue,
608 numElements);
609 }
610
test_relational_isgreater(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)611 int test_relational_isgreater(cl_device_id device, cl_context context,
612 cl_command_queue queue, int numElements)
613 {
614 return MakeAndRunTest<IsGreaterFPTest>(device, context, queue, numElements);
615 }
616
test_relational_isgreaterequal(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)617 int test_relational_isgreaterequal(cl_device_id device, cl_context context,
618 cl_command_queue queue, int numElements)
619 {
620 return MakeAndRunTest<IsGreaterEqualFPTest>(device, context, queue,
621 numElements);
622 }
623
test_relational_isless(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)624 int test_relational_isless(cl_device_id device, cl_context context,
625 cl_command_queue queue, int numElements)
626 {
627 return MakeAndRunTest<IsLessFPTest>(device, context, queue, numElements);
628 }
629
test_relational_islessequal(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)630 int test_relational_islessequal(cl_device_id device, cl_context context,
631 cl_command_queue queue, int numElements)
632 {
633 return MakeAndRunTest<IsLessEqualFPTest>(device, context, queue,
634 numElements);
635 }
636
test_relational_islessgreater(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)637 int test_relational_islessgreater(cl_device_id device, cl_context context,
638 cl_command_queue queue, int numElements)
639 {
640 return MakeAndRunTest<IsLessGreaterFPTest>(device, context, queue,
641 numElements);
642 }
643