From f31b2f029c9e33b018460666ebf47950fa9d6224 Mon Sep 17 00:00:00 2001 From: Marcin Hajder Date: Tue, 16 May 2023 17:43:47 +0200 Subject: [PATCH] Added cl_half support for test_relationals (#1623) * Added cl_khr_fp16 support for test_relationals (issue #142, relationals) * Added cl_khr_fp16 support for any and bitselect test cases (issue #142, relationals) * correction related to automated travis build for macOS (issue #142, relationals) * more corrections related to automated travis build for macOS (issue #142, relationals) * Added few cosmetic corrections (issue #142, test_relationals) * Added missing clang format * Added corrections related to order of initialization * Added corrections due to code review (issue #142, relationals) * Correction for prev commit * Added subnormals related condition for test verification (issue #142, relationals) * Added indexing correction due to code review * Replaced hardcoded iteration limit (issue #142, relationals) --- test_conformance/printf/test_printf.cpp | 13 +- test_conformance/relationals/CMakeLists.txt | 3 +- .../relationals/test_comparisons_double.cpp | 363 ---------- .../relationals/test_comparisons_float.cpp | 362 ---------- .../relationals/test_comparisons_fp.cpp | 661 ++++++++++++++++++ .../relationals/test_comparisons_fp.h | 227 ++++++ .../relationals/test_relationals.cpp | 224 +++--- 7 files changed, 1008 insertions(+), 845 deletions(-) delete mode 100644 test_conformance/relationals/test_comparisons_double.cpp delete mode 100644 test_conformance/relationals/test_comparisons_float.cpp create mode 100644 test_conformance/relationals/test_comparisons_fp.cpp create mode 100644 test_conformance/relationals/test_comparisons_fp.h diff --git a/test_conformance/printf/test_printf.cpp b/test_conformance/printf/test_printf.cpp index e789e0ca..e43e302f 100644 --- a/test_conformance/printf/test_printf.cpp +++ b/test_conformance/printf/test_printf.cpp @@ -268,7 +268,7 @@ static cl_program makePrintfProgram(cl_kernel *kernel_ptr, const cl_context cont }; //Update testname - sprintf(testname,"%s%d","test",testId); + std::snprintf(testname, sizeof(testname), "%s%d", "test", testId); if (allTestCase[testId]->_type == TYPE_HALF || allTestCase[testId]->_type == TYPE_HALF_LIMITS) @@ -278,13 +278,18 @@ static cl_program makePrintfProgram(cl_kernel *kernel_ptr, const cl_context cont //Update addrSpaceArgument and addrSpacePAddArgument types, based on FULL_PROFILE/EMBEDDED_PROFILE if(allTestCase[testId]->_type == TYPE_ADDRESS_SPACE) { - sprintf(addrSpaceArgument, "%s",allTestCase[testId]->_genParameters[testNum].addrSpaceArgumentTypeQualifier); + std::snprintf(addrSpaceArgument, sizeof(addrSpaceArgument), "%s", + allTestCase[testId] + ->_genParameters[testNum] + .addrSpaceArgumentTypeQualifier); - sprintf(addrSpacePAddArgument, "%s", allTestCase[testId]->_genParameters[testNum].addrSpacePAdd); + std::snprintf( + addrSpacePAddArgument, sizeof(addrSpacePAddArgument), "%s", + allTestCase[testId]->_genParameters[testNum].addrSpacePAdd); } if (strlen(addrSpaceArgument) == 0) - sprintf(addrSpaceArgument,"void"); + std::snprintf(addrSpaceArgument, sizeof(addrSpaceArgument), "void"); // create program based on its type diff --git a/test_conformance/relationals/CMakeLists.txt b/test_conformance/relationals/CMakeLists.txt index ecaa056c..aa5dd6a1 100644 --- a/test_conformance/relationals/CMakeLists.txt +++ b/test_conformance/relationals/CMakeLists.txt @@ -3,8 +3,7 @@ set(MODULE_NAME RELATIONALS) set(${MODULE_NAME}_SOURCES main.cpp test_relationals.cpp - test_comparisons_float.cpp - test_comparisons_double.cpp + test_comparisons_fp.cpp test_shuffles.cpp ) diff --git a/test_conformance/relationals/test_comparisons_double.cpp b/test_conformance/relationals/test_comparisons_double.cpp deleted file mode 100644 index 3fe1124c..00000000 --- a/test_conformance/relationals/test_comparisons_double.cpp +++ /dev/null @@ -1,363 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "testBase.h" -#include "harness/conversions.h" -#include "harness/typeWrappers.h" - -#define TEST_SIZE 512 - -const char *equivTestKernelPattern_double = -"#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" -"__kernel void sample_test(__global double%s *sourceA, __global double%s *sourceB, __global long%s *destValues, __global long%s *destValuesB)\n" -"{\n" -" int tid = get_global_id(0);\n" -" destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n" -" destValuesB[tid] = sourceA[tid] %s sourceB[tid];\n" -"\n" -"}\n"; - -const char *equivTestKernelPatternLessGreater_double = -"#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" -"__kernel void sample_test(__global double%s *sourceA, __global double%s *sourceB, __global long%s *destValues, __global long%s *destValuesB)\n" -"{\n" -" int tid = get_global_id(0);\n" -" destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n" -" destValuesB[tid] = (sourceA[tid] < sourceB[tid]) | (sourceA[tid] > sourceB[tid]);\n" -"\n" -"}\n"; - - -const char *equivTestKernelPattern_double3 = -"#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" -"__kernel void sample_test(__global double%s *sourceA, __global double%s *sourceB, __global long%s *destValues, __global long%s *destValuesB)\n" -"{\n" -" int tid = get_global_id(0);\n" -" double3 sampA = vload3(tid, (__global double *)sourceA);\n" -" double3 sampB = vload3(tid, (__global double *)sourceB);\n" -" vstore3(%s( sampA, sampB ), tid, (__global long *)destValues);\n" -" vstore3(( sampA %s sampB ), tid, (__global long *)destValuesB);\n" -"\n" -"}\n"; - -const char *equivTestKernelPatternLessGreater_double3 = -"#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" -"__kernel void sample_test(__global double%s *sourceA, __global double%s *sourceB, __global long%s *destValues, __global long%s *destValuesB)\n" -"{\n" -" int tid = get_global_id(0);\n" -" double3 sampA = vload3(tid, (__global double *)sourceA);\n" -" double3 sampB = vload3(tid, (__global double *)sourceB);\n" -" vstore3(%s( sampA, sampB ), tid, (__global long *)destValues);\n" -" vstore3(( sampA < sampB ) | (sampA > sampB), tid, (__global long *)destValuesB);\n" -"\n" -"}\n"; - - -typedef bool (*equivVerifyFn)( double inDataA, double inDataB ); - -void verify_equiv_values_double( unsigned int vecSize, double *inDataA, double *inDataB, cl_long *outData, equivVerifyFn verifyFn ) -{ - unsigned int i; - cl_long trueResult; - bool result; - - trueResult = ( vecSize == 1 ) ? 1 : -1; - for( i = 0; i < vecSize; i++ ) - { - result = verifyFn( inDataA[ i ], inDataB[ i ] ); - outData[ i ] = result ? trueResult : 0; - } -} - -void generate_equiv_test_data_double( double *outData, unsigned int vecSize, bool alpha, MTdata d ) -{ - unsigned int i; - - generate_random_data( kDouble, vecSize * TEST_SIZE, d, outData ); - - // Fill the first few vectors with NAN in each vector element (or the second set if we're alpha, so we can test either case) - if( alpha ) - outData += vecSize * vecSize; - for( i = 0; i < vecSize; i++ ) - { - outData[ 0 ] = NAN; - outData += vecSize + 1; - } - // Make sure the third set is filled regardless, to test the case where both have NANs - if( !alpha ) - outData += vecSize * vecSize; - for( i = 0; i < vecSize; i++ ) - { - outData[ 0 ] = NAN; - outData += vecSize + 1; - } -} - -int test_equiv_kernel_double(cl_context context, cl_command_queue queue, const char *fnName, const char *opName, - unsigned int vecSize, equivVerifyFn verifyFn, MTdata d ) -{ - clProgramWrapper program; - clKernelWrapper kernel; - clMemWrapper streams[4]; - double inDataA[TEST_SIZE * 16], inDataB[ TEST_SIZE * 16 ]; - cl_long outData[TEST_SIZE * 16], expected[16]; - int error, i, j; - size_t threads[1], localThreads[1]; - char kernelSource[10240]; - char *programPtr; - char sizeName[4]; - - - /* Create the source */ - if( vecSize == 1 ) - sizeName[ 0 ] = 0; - else - sprintf( sizeName, "%d", vecSize ); - - if(DENSE_PACK_VECS && vecSize == 3) { - if (strcmp(fnName, "islessgreater")) { - sprintf( kernelSource, equivTestKernelPattern_double3, sizeName, sizeName, sizeName, sizeName, fnName, opName ); - } else { - sprintf( kernelSource, equivTestKernelPatternLessGreater_double3, sizeName, sizeName, sizeName, sizeName, fnName ); - } - } else { - if (strcmp(fnName, "islessgreater")) { - sprintf( kernelSource, equivTestKernelPattern_double, sizeName, sizeName, sizeName, sizeName, fnName, opName ); - } else { - sprintf( kernelSource, equivTestKernelPatternLessGreater_double, sizeName, sizeName, sizeName, sizeName, fnName ); - } - } - - /* Create kernels */ - programPtr = kernelSource; - if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) ) - { - return -1; - } - - /* Generate some streams */ - generate_equiv_test_data_double( inDataA, vecSize, true, d ); - generate_equiv_test_data_double( inDataB, vecSize, false, d ); - - streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, - sizeof(cl_double) * vecSize * TEST_SIZE, - &inDataA, &error); - if( streams[0] == NULL ) - { - print_error( error, "Creating input array A failed!\n"); - return -1; - } - streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, - sizeof(cl_double) * vecSize * TEST_SIZE, - &inDataB, &error); - if( streams[1] == NULL ) - { - print_error( error, "Creating input array A failed!\n"); - return -1; - } - streams[2] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizeof( cl_long ) * vecSize * TEST_SIZE, NULL, &error); - if( streams[2] == NULL ) - { - print_error( error, "Creating output array failed!\n"); - return -1; - } - streams[3] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizeof( cl_long ) * vecSize * TEST_SIZE, NULL, &error); - if( streams[3] == NULL ) - { - print_error( error, "Creating output array failed!\n"); - return -1; - } - - - /* Assign streams and execute */ - error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] ); - test_error( error, "Unable to set indexed kernel arguments" ); - error = clSetKernelArg( kernel, 1, sizeof( streams[1] ), &streams[1] ); - test_error( error, "Unable to set indexed kernel arguments" ); - error = clSetKernelArg( kernel, 2, sizeof( streams[2] ), &streams[2] ); - test_error( error, "Unable to set indexed kernel arguments" ); - error = clSetKernelArg( kernel, 3, sizeof( streams[3] ), &streams[3] ); - test_error( error, "Unable to set indexed kernel arguments" ); - - - /* Run the kernel */ - threads[0] = TEST_SIZE; - - error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] ); - test_error( error, "Unable to get work group size to use" ); - - error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL ); - test_error( error, "Unable to execute test kernel" ); - - /* Now get the results */ - error = clEnqueueReadBuffer( queue, streams[2], true, 0, sizeof( cl_long ) * TEST_SIZE * vecSize, outData, 0, NULL, NULL ); - test_error( error, "Unable to read output array!" ); - - /* And verify! */ - for( i = 0; i < TEST_SIZE; i++ ) - { - verify_equiv_values_double( vecSize, &inDataA[ i * vecSize ], &inDataB[ i * vecSize ], expected, verifyFn); - - for( j = 0; j < (int)vecSize; j++ ) - { - if( expected[ j ] != outData[ i * vecSize + j ] ) - { - log_error( "ERROR: Data sample %d:%d at size %d does not validate! Expected %lld, got %lld, source %f,%f\n", - i, j, vecSize, expected[ j ], outData[ i * vecSize + j ], inDataA[i*vecSize + j], inDataB[i*vecSize + j] ); - return -1; - } - } - } - - /* Now get the results */ - error = clEnqueueReadBuffer( queue, streams[3], true, 0, sizeof( cl_long ) * TEST_SIZE * vecSize, outData, 0, NULL, NULL ); - test_error( error, "Unable to read output array!" ); - - /* And verify! */ - for( i = 0; i < TEST_SIZE; i++ ) - { - verify_equiv_values_double( vecSize, &inDataA[ i * vecSize ], &inDataB[ i * vecSize ], expected, verifyFn); - - for( j = 0; j < (int)vecSize; j++ ) - { - if( expected[ j ] != outData[ i * vecSize + j ] ) - { - log_error( "ERROR: Data sample %d:%d at size %d does not validate! Expected %lld, got %lld, source %f,%f\n", - i, j, vecSize, expected[ j ], outData[ i * vecSize + j ], inDataA[i*vecSize + j], inDataB[i*vecSize + j] ); - return -1; - } - } - } - - return 0; -} - -int test_equiv_kernel_set_double(cl_device_id device, cl_context context, cl_command_queue queue, const char *fnName, const char *opName, equivVerifyFn verifyFn, MTdata d ) -{ - unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 }; - unsigned int index; - int retVal = 0; - - if (!is_extension_available(device, "cl_khr_fp64")) { - log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n"); - return 0; - } - log_info("Testing doubles.\n"); - - for( index = 0; vecSizes[ index ] != 0; index++ ) - { - // Test! - if( test_equiv_kernel_double(context, queue, fnName, opName, vecSizes[ index ], verifyFn, d ) != 0 ) - { - log_error( " Vector double%d FAILED\n", vecSizes[ index ] ); - retVal = -1; - } - } - - return retVal; -} - -bool isequal_verify_fn_double( double valueA, double valueB ) -{ - if( isnan( valueA ) || isnan( valueB ) ) - return false; - return valueA == valueB; -} - -int test_relational_isequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed(gRandomSeed); - return test_equiv_kernel_set_double( device, context, queue, "isequal", "==", isequal_verify_fn_double, seed ); -} - -bool isnotequal_verify_fn_double( double valueA, double valueB ) -{ - if( isnan( valueA ) || isnan( valueB ) ) - return true; - return valueA != valueB; -} - -int test_relational_isnotequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed(gRandomSeed); - return test_equiv_kernel_set_double( device, context, queue, "isnotequal", "!=", isnotequal_verify_fn_double, seed ); -} - -bool isgreater_verify_fn_double( double valueA, double valueB ) -{ - if( isnan( valueA ) || isnan( valueB ) ) - return false; - return valueA > valueB; -} - -int test_relational_isgreater_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed(gRandomSeed); - return test_equiv_kernel_set_double( device, context, queue, "isgreater", ">", isgreater_verify_fn_double, seed ); -} - -bool isgreaterequal_verify_fn_double( double valueA, double valueB ) -{ - if( isnan( valueA ) || isnan( valueB ) ) - return false; - return valueA >= valueB; -} - -int test_relational_isgreaterequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed(gRandomSeed); - return test_equiv_kernel_set_double( device, context, queue, "isgreaterequal", ">=", isgreaterequal_verify_fn_double, seed ); -} - -bool isless_verify_fn_double( double valueA, double valueB ) -{ - if( isnan( valueA ) || isnan( valueB ) ) - return false; - return valueA < valueB; -} - -int test_relational_isless_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed(gRandomSeed); - return test_equiv_kernel_set_double( device, context, queue, "isless", "<", isless_verify_fn_double, seed ); -} - -bool islessequal_verify_fn_double( double valueA, double valueB ) -{ - if( isnan( valueA ) || isnan( valueB ) ) - return false; - return valueA <= valueB; -} - -int test_relational_islessequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed(gRandomSeed); - return test_equiv_kernel_set_double( device, context, queue, "islessequal", "<=", islessequal_verify_fn_double, seed ); -} - -bool islessgreater_verify_fn_double( double valueA, double valueB ) -{ - if( isnan( valueA ) || isnan( valueB ) ) - return false; - return ( valueA < valueB ) || ( valueA > valueB ); -} - -int test_relational_islessgreater_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed(gRandomSeed); - return test_equiv_kernel_set_double( device, context, queue, "islessgreater", "<>", islessgreater_verify_fn_double, seed ); -} - - diff --git a/test_conformance/relationals/test_comparisons_float.cpp b/test_conformance/relationals/test_comparisons_float.cpp deleted file mode 100644 index 274cd71b..00000000 --- a/test_conformance/relationals/test_comparisons_float.cpp +++ /dev/null @@ -1,362 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "testBase.h" -#include "harness/conversions.h" -#include "harness/typeWrappers.h" - -#define TEST_SIZE 512 - -const char *equivTestKernelPattern_float = -"__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n" -"{\n" -" int tid = get_global_id(0);\n" -" destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n" -" destValuesB[tid] = sourceA[tid] %s sourceB[tid];\n" -"\n" -"}\n"; - -const char *equivTestKernelPatternLessGreater_float = -"__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n" -"{\n" -" int tid = get_global_id(0);\n" -" destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n" -" destValuesB[tid] = (sourceA[tid] < sourceB[tid]) | (sourceA[tid] > sourceB[tid]);\n" -"\n" -"}\n"; - - -const char *equivTestKernelPattern_float3 = -"__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n" -"{\n" -" int tid = get_global_id(0);\n" -" float3 sampA = vload3(tid, (__global float *)sourceA);\n" -" float3 sampB = vload3(tid, (__global float *)sourceB);\n" -" vstore3(%s( sampA, sampB ), tid, (__global int *)destValues);\n" -" vstore3(( sampA %s sampB ), tid, (__global int *)destValuesB);\n" -"\n" -"}\n"; - -const char *equivTestKernelPatternLessGreater_float3 = -"__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n" -"{\n" -" int tid = get_global_id(0);\n" -" float3 sampA = vload3(tid, (__global float *)sourceA);\n" -" float3 sampB = vload3(tid, (__global float *)sourceB);\n" -" vstore3(%s( sampA, sampB ), tid, (__global int *)destValues);\n" -" vstore3(( sampA < sampB ) | (sampA > sampB), tid, (__global int *)destValuesB);\n" -"\n" -"}\n"; - -typedef bool (*equivVerifyFn)( float inDataA, float inDataB ); - -int IsFloatInfinity(float x) -{ - return isinf(x); -} - -int IsFloatNaN(float x) -{ - return isnan(x); -} - -void verify_equiv_values_float( unsigned int vecSize, float *inDataA, float *inDataB, int *outData, equivVerifyFn verifyFn ) -{ - unsigned int i; - int trueResult; - bool result; - - trueResult = ( vecSize == 1 ) ? 1 : -1; - for( i = 0; i < vecSize; i++ ) - { - result = verifyFn( inDataA[ i ], inDataB[ i ] ); - outData[ i ] = result ? trueResult : 0; - } -} - -void generate_equiv_test_data_float( float *outData, unsigned int vecSize, bool alpha, MTdata d ) -{ - unsigned int i; - - generate_random_data( kFloat, vecSize * TEST_SIZE, d, outData ); - - // Fill the first few vectors with NAN in each vector element (or the second set if we're alpha, so we can test either case) - if( alpha ) - outData += vecSize * vecSize; - for( i = 0; i < vecSize; i++ ) - { - outData[ 0 ] = NAN; - outData += vecSize + 1; - } - // Make sure the third set is filled regardless, to test the case where both have NANs - if( !alpha ) - outData += vecSize * vecSize; - for( i = 0; i < vecSize; i++ ) - { - outData[ 0 ] = NAN; - outData += vecSize + 1; - } -} - -int test_equiv_kernel_float(cl_context context, cl_command_queue queue, const char *fnName, const char *opName, - unsigned int vecSize, equivVerifyFn verifyFn, MTdata d ) -{ - clProgramWrapper program; - clKernelWrapper kernel; - clMemWrapper streams[4]; - float inDataA[TEST_SIZE * 16], inDataB[ TEST_SIZE * 16 ]; - int outData[TEST_SIZE * 16], expected[16]; - int error, i, j; - size_t threads[1], localThreads[1]; - char kernelSource[10240]; - char *programPtr; - char sizeName[4]; - - - /* Create the source */ - if( vecSize == 1 ) - sizeName[ 0 ] = 0; - else - sprintf( sizeName, "%d", vecSize ); - - - if(DENSE_PACK_VECS && vecSize == 3) { - if (strcmp(fnName, "islessgreater")) { - sprintf( kernelSource, equivTestKernelPattern_float3, sizeName, sizeName, sizeName, sizeName, fnName, opName ); - } else { - sprintf( kernelSource, equivTestKernelPatternLessGreater_float3, sizeName, sizeName, sizeName, sizeName, fnName ); - } - } else { - if (strcmp(fnName, "islessgreater")) { - sprintf( kernelSource, equivTestKernelPattern_float, sizeName, sizeName, sizeName, sizeName, fnName, opName ); - } else { - sprintf( kernelSource, equivTestKernelPatternLessGreater_float, sizeName, sizeName, sizeName, sizeName, fnName ); - } - } - - /* Create kernels */ - programPtr = kernelSource; - if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) ) - { - return -1; - } - - /* Generate some streams */ - generate_equiv_test_data_float( inDataA, vecSize, true, d ); - generate_equiv_test_data_float( inDataB, vecSize, false, d ); - - streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, - sizeof(cl_float) * vecSize * TEST_SIZE, - &inDataA, &error); - if( streams[0] == NULL ) - { - print_error( error, "Creating input array A failed!\n"); - return -1; - } - streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, - sizeof(cl_float) * vecSize * TEST_SIZE, - &inDataB, &error); - if( streams[1] == NULL ) - { - print_error( error, "Creating input array A failed!\n"); - return -1; - } - streams[2] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizeof( cl_int ) * vecSize * TEST_SIZE, NULL, &error); - if( streams[2] == NULL ) - { - print_error( error, "Creating output array failed!\n"); - return -1; - } - streams[3] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizeof( cl_int ) * vecSize * TEST_SIZE, NULL, &error); - if( streams[3] == NULL ) - { - print_error( error, "Creating output array failed!\n"); - return -1; - } - - - /* Assign streams and execute */ - error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] ); - test_error( error, "Unable to set indexed kernel arguments" ); - error = clSetKernelArg( kernel, 1, sizeof( streams[1] ), &streams[1] ); - test_error( error, "Unable to set indexed kernel arguments" ); - error = clSetKernelArg( kernel, 2, sizeof( streams[2] ), &streams[2] ); - test_error( error, "Unable to set indexed kernel arguments" ); - error = clSetKernelArg( kernel, 3, sizeof( streams[3] ), &streams[3] ); - test_error( error, "Unable to set indexed kernel arguments" ); - - - /* Run the kernel */ - threads[0] = TEST_SIZE; - - error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] ); - test_error( error, "Unable to get work group size to use" ); - - error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL ); - test_error( error, "Unable to execute test kernel" ); - - /* Now get the results */ - error = clEnqueueReadBuffer( queue, streams[2], true, 0, sizeof( int ) * TEST_SIZE * vecSize, outData, 0, NULL, NULL ); - test_error( error, "Unable to read output array!" ); - - /* And verify! */ - for( i = 0; i < TEST_SIZE; i++ ) - { - verify_equiv_values_float( vecSize, &inDataA[ i * vecSize ], &inDataB[ i * vecSize ], expected, verifyFn); - - for( j = 0; j < (int)vecSize; j++ ) - { - if( expected[ j ] != outData[ i * vecSize + j ] ) - { - log_error( "ERROR: Data sample %d:%d at size %d does not validate! Expected %d, got %d, source %f,%f\n", - i, j, vecSize, expected[ j ], outData[ i * vecSize + j ], inDataA[i*vecSize + j], inDataB[i*vecSize + j] ); - return -1; - } - } - } - - /* Now get the results */ - error = clEnqueueReadBuffer( queue, streams[3], true, 0, sizeof( int ) * TEST_SIZE * vecSize, outData, 0, NULL, NULL ); - test_error( error, "Unable to read output array!" ); - - /* And verify! */ - int fail = 0; - for( i = 0; i < TEST_SIZE; i++ ) - { - verify_equiv_values_float( vecSize, &inDataA[ i * vecSize ], &inDataB[ i * vecSize ], expected, verifyFn); - - for( j = 0; j < (int)vecSize; j++ ) - { - if( expected[ j ] != outData[ i * vecSize + j ] ) - { - if (gInfNanSupport == 0) - { - if (IsFloatNaN(inDataA[i*vecSize + j]) || IsFloatNaN (inDataB[i*vecSize + j])) - { - fail = 0; - } - else - fail = 1; - } - if (fail) - { - log_error( "ERROR: Data sample %d:%d at size %d does not validate! Expected %d, got %d, source %f,%f\n", - i, j, vecSize, expected[ j ], outData[ i * vecSize + j ], inDataA[i*vecSize + j], inDataB[i*vecSize + j] ); - return -1; - } - } - } - } - - return 0; -} - -int test_equiv_kernel_set_float(cl_context context, cl_command_queue queue, const char *fnName, const char *opName, equivVerifyFn verifyFn, MTdata d ) -{ - unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 }; - unsigned int index; - int retVal = 0; - - for( index = 0; vecSizes[ index ] != 0; index++ ) - { - // Test! - if( test_equiv_kernel_float(context, queue, fnName, opName, vecSizes[ index ], verifyFn, d ) != 0 ) - { - log_error( " Vector float%d FAILED\n", vecSizes[ index ] ); - retVal = -1; - } - } - - return retVal; -} - -bool isequal_verify_fn_float( float valueA, float valueB ) -{ - return valueA == valueB; -} - -int test_relational_isequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed( gRandomSeed ); - return test_equiv_kernel_set_float( context, queue, "isequal", "==", isequal_verify_fn_float, seed ); -} - -bool isnotequal_verify_fn_float( float valueA, float valueB ) -{ - return valueA != valueB; -} - -int test_relational_isnotequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed( gRandomSeed ); - return test_equiv_kernel_set_float( context, queue, "isnotequal", "!=", isnotequal_verify_fn_float, seed ); -} - -bool isgreater_verify_fn_float( float valueA, float valueB ) -{ - return valueA > valueB; -} - -int test_relational_isgreater_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed( gRandomSeed ); - return test_equiv_kernel_set_float( context, queue, "isgreater", ">", isgreater_verify_fn_float, seed ); -} - -bool isgreaterequal_verify_fn_float( float valueA, float valueB ) -{ - return valueA >= valueB; -} - -int test_relational_isgreaterequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed( gRandomSeed ); - return test_equiv_kernel_set_float( context, queue, "isgreaterequal", ">=", isgreaterequal_verify_fn_float, seed ); -} - -bool isless_verify_fn_float( float valueA, float valueB ) -{ - return valueA < valueB; -} - -int test_relational_isless_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed( gRandomSeed ); - return test_equiv_kernel_set_float( context, queue, "isless", "<", isless_verify_fn_float, seed ); -} - -bool islessequal_verify_fn_float( float valueA, float valueB ) -{ - return valueA <= valueB; -} - -int test_relational_islessequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed( gRandomSeed ); - return test_equiv_kernel_set_float( context, queue, "islessequal", "<=", islessequal_verify_fn_float, seed ); -} - -bool islessgreater_verify_fn_float( float valueA, float valueB ) -{ - return ( valueA < valueB ) || ( valueA > valueB ); -} - -int test_relational_islessgreater_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - RandomSeed seed( gRandomSeed ); - return test_equiv_kernel_set_float( context, queue, "islessgreater", "<>", islessgreater_verify_fn_float, seed ); -} - - diff --git a/test_conformance/relationals/test_comparisons_fp.cpp b/test_conformance/relationals/test_comparisons_fp.cpp new file mode 100644 index 00000000..580b7422 --- /dev/null +++ b/test_conformance/relationals/test_comparisons_fp.cpp @@ -0,0 +1,661 @@ +// +// Copyright (c) 2022 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + +#include +#include +#include +#include +#include + +#include + +#include "test_comparisons_fp.h" + +#define TEST_SIZE 512 + +static char ftype[32] = { 0 }; +static char ftype_vec[32] = { 0 }; +static char itype[32] = { 0 }; +static char itype_vec[32] = { 0 }; +static char extension[128] = { 0 }; + +// clang-format off +// for readability sake keep this section unformatted +const char* equivTestKernPat[] = { +extension, +"__kernel void sample_test(__global ", ftype_vec, " *sourceA, __global ", ftype_vec, +" *sourceB, __global ", itype_vec, " *destValues, __global ", itype_vec, " *destValuesB)\n" +"{\n" +" int tid = get_global_id(0);\n" +" destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n" +" destValuesB[tid] = sourceA[tid] %s sourceB[tid];\n" +"}\n"}; + +const char* equivTestKernPatLessGreater[] = { +extension, +"__kernel void sample_test(__global ", ftype_vec, " *sourceA, __global ", ftype_vec, +" *sourceB, __global ", itype_vec, " *destValues, __global ", itype_vec, " *destValuesB)\n" +"{\n" +" int tid = get_global_id(0);\n" +" destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n" +" destValuesB[tid] = (sourceA[tid] < sourceB[tid]) | (sourceA[tid] > sourceB[tid]);\n" +"}\n"}; + +const char* equivTestKerPat_3[] = { +extension, +"__kernel void sample_test(__global ", ftype_vec, " *sourceA, __global ", ftype_vec, +" *sourceB, __global ", itype_vec, " *destValues, __global ", itype_vec, " *destValuesB)\n" +"{\n" +" int tid = get_global_id(0);\n" +" ",ftype_vec," sampA = vload3(tid, (__global ",ftype," *)sourceA);\n" +" ",ftype_vec," sampB = vload3(tid, (__global ",ftype," *)sourceB);\n" +" vstore3(%s( sampA, sampB ), tid, (__global ",itype," *)destValues);\n" +" vstore3(( sampA %s sampB ), tid, (__global ",itype," *)destValuesB);\n" +"}\n"}; + +const char* equivTestKerPatLessGreater_3[] = { +extension, +"__kernel void sample_test(__global ", ftype_vec, " *sourceA, __global ", ftype_vec, +" *sourceB, __global ", itype_vec, " *destValues, __global ", itype_vec, " *destValuesB)\n" +"{\n" +" int tid = get_global_id(0);\n" +" ", ftype_vec, " sampA = vload3(tid, (__global ", ftype, " *)sourceA);\n" +" ", ftype_vec, " sampB = vload3(tid, (__global ", ftype, " *)sourceB);\n" +" vstore3(%s( sampA, sampB ), tid, (__global ", itype, " *)destValues);\n" +" vstore3(( sampA < sampB ) | (sampA > sampB), tid, (__global ", itype, " *)destValuesB);\n" +"}\n" +}; +// clang-format on + + +std::string concat_kernel(const char* sstr[], int num) +{ + std::string res; + for (int i = 0; i < num; i++) res += std::string(sstr[i]); + return res; +} + +template +std::string string_format(const std::string& format, Args... args) +{ + int size_s = std::snprintf(nullptr, 0, format.c_str(), args...) + + 1; // Extra space for '\0' + if (size_s <= 0) + { + throw std::runtime_error("Error during formatting."); + } + auto size = static_cast(size_s); + std::unique_ptr buf(new char[size]); + std::snprintf(buf.get(), size, format.c_str(), args...); + return std::string(buf.get(), + buf.get() + size - 1); // We don't want the '\0' inside +} + +template bool verify(const T& A, const T& B) +{ + return F()(A, B); +} + +RelationalsFPTest::RelationalsFPTest(cl_context context, cl_device_id device, + cl_command_queue queue, const char* fn, + const char* op) + : context(context), device(device), queue(queue), fnName(fn), opName(op), + halfFlushDenormsToZero(0) +{ + // hardcoded for now, to be changed into typeid().name solution in future + // for now C++ spec doesn't guarantee human readable type name + + eqTypeNames = { { kHalf, "short" }, + { kFloat, "int" }, + { kDouble, "long" } }; +} + +template +void RelationalsFPTest::generate_equiv_test_data(T* outData, + unsigned int vecSize, + bool alpha, + const RelTestParams& param, + const MTdata& d) +{ + unsigned int i; + + generate_random_data(param.dataType, vecSize * TEST_SIZE, d, outData); + + // Fill the first few vectors with NAN in each vector element (or the second + // set if we're alpha, so we can test either case) + if (alpha) outData += vecSize * vecSize; + for (i = 0; i < vecSize; i++) + { + outData[0] = param.nan; + outData += vecSize + 1; + } + // Make sure the third set is filled regardless, to test the case where both + // have NANs + if (!alpha) outData += vecSize * vecSize; + for (i = 0; i < vecSize; i++) + { + outData[0] = param.nan; + outData += vecSize + 1; + } +} + +template +void RelationalsFPTest::verify_equiv_values(unsigned int vecSize, + const T* const inDataA, + const T* const inDataB, + U* const outData, + const VerifyFunc& verifyFn) +{ + unsigned int i; + int trueResult; + bool result; + + trueResult = (vecSize == 1) ? 1 : -1; + for (i = 0; i < vecSize; i++) + { + result = verifyFn(inDataA[i], inDataB[i]); + outData[i] = result ? trueResult : 0; + } +} + +template +int RelationalsFPTest::test_equiv_kernel(unsigned int vecSize, + const RelTestParams& param, + const MTdata& d) +{ + clProgramWrapper program; + clKernelWrapper kernel; + clMemWrapper streams[4]; + T inDataA[TEST_SIZE * 16], inDataB[TEST_SIZE * 16]; + + // support half, float, double equivalents - otherwise assert + typedef typename std::conditional< + (sizeof(T) == sizeof(std::int16_t)), std::int16_t, + typename std::conditional<(sizeof(T) == sizeof(std::int32_t)), + std::int32_t, std::int64_t>::type>::type U; + + U outData[TEST_SIZE * 16], expected[16]; + int error, i, j; + size_t threads[1], localThreads[1]; + std::string kernelSource; + char sizeName[4]; + + /* Create the source */ + if (vecSize == 1) + sizeName[0] = 0; + else + sprintf(sizeName, "%d", vecSize); + + if (eqTypeNames.find(param.dataType) == eqTypeNames.end()) + log_error( + "RelationalsFPTest::test_equiv_kernel: unsupported fp data type"); + + sprintf(ftype, "%s", get_explicit_type_name(param.dataType)); + sprintf(ftype_vec, "%s%s", get_explicit_type_name(param.dataType), + sizeName); + + sprintf(itype, "%s", eqTypeNames[param.dataType].c_str()); + sprintf(itype_vec, "%s%s", eqTypeNames[param.dataType].c_str(), sizeName); + + if (std::is_same::value) + strcpy(extension, "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"); + else if (std::is_same::value) + strcpy(extension, "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"); + else + extension[0] = '\0'; + + if (DENSE_PACK_VECS && vecSize == 3) + { + if (strcmp(fnName.c_str(), "islessgreater")) + { + auto str = + concat_kernel(equivTestKerPat_3, + sizeof(equivTestKerPat_3) / sizeof(const char*)); + kernelSource = string_format(str, fnName.c_str(), opName.c_str()); + } + else + { + auto str = concat_kernel(equivTestKerPatLessGreater_3, + sizeof(equivTestKerPatLessGreater_3) + / sizeof(const char*)); + kernelSource = string_format(str, fnName.c_str()); + } + } + else + { + if (strcmp(fnName.c_str(), "islessgreater")) + { + auto str = + concat_kernel(equivTestKernPat, + sizeof(equivTestKernPat) / sizeof(const char*)); + kernelSource = string_format(str, fnName.c_str(), opName.c_str()); + } + else + { + auto str = concat_kernel(equivTestKernPatLessGreater, + sizeof(equivTestKernPatLessGreater) + / sizeof(const char*)); + kernelSource = string_format(str, fnName.c_str()); + } + } + + /* Create kernels */ + const char* programPtr = kernelSource.c_str(); + if (create_single_kernel_helper(context, &program, &kernel, 1, + (const char**)&programPtr, "sample_test")) + { + return -1; + } + + /* Generate some streams */ + generate_equiv_test_data(inDataA, vecSize, true, param, d); + generate_equiv_test_data(inDataB, vecSize, false, param, d); + + streams[0] = + clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, + sizeof(T) * vecSize * TEST_SIZE, &inDataA, &error); + if (streams[0] == NULL) + { + print_error(error, "Creating input array A failed!\n"); + return -1; + } + streams[1] = + clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, + sizeof(T) * vecSize * TEST_SIZE, &inDataB, &error); + if (streams[1] == NULL) + { + print_error(error, "Creating input array A failed!\n"); + return -1; + } + streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(U) * vecSize * TEST_SIZE, NULL, &error); + if (streams[2] == NULL) + { + print_error(error, "Creating output array failed!\n"); + return -1; + } + streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, + sizeof(U) * vecSize * TEST_SIZE, NULL, &error); + if (streams[3] == NULL) + { + print_error(error, "Creating output array failed!\n"); + return -1; + } + + /* Assign streams and execute */ + error = clSetKernelArg(kernel, 0, sizeof(streams[0]), &streams[0]); + test_error(error, "Unable to set indexed kernel arguments"); + error = clSetKernelArg(kernel, 1, sizeof(streams[1]), &streams[1]); + test_error(error, "Unable to set indexed kernel arguments"); + error = clSetKernelArg(kernel, 2, sizeof(streams[2]), &streams[2]); + test_error(error, "Unable to set indexed kernel arguments"); + error = clSetKernelArg(kernel, 3, sizeof(streams[3]), &streams[3]); + test_error(error, "Unable to set indexed kernel arguments"); + + /* Run the kernel */ + threads[0] = TEST_SIZE; + + error = get_max_common_work_group_size(context, kernel, threads[0], + &localThreads[0]); + test_error(error, "Unable to get work group size to use"); + + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, + localThreads, 0, NULL, NULL); + test_error(error, "Unable to execute test kernel"); + + /* Now get the results */ + error = clEnqueueReadBuffer(queue, streams[2], true, 0, + sizeof(U) * TEST_SIZE * vecSize, outData, 0, + NULL, NULL); + test_error(error, "Unable to read output array!"); + + auto verror_msg = [](const int& i, const int& j, const unsigned& vs, + const U& e, const U& o, const T& iA, const T& iB) { + std::stringstream sstr; + sstr << "ERROR: Data sample " << i << ":" << j << " at size " << vs + << " does not validate! Expected " << e << ", got " << o + << ", source " << iA << ":" << iB << std::endl; + log_error(sstr.str().c_str()); + }; + + /* And verify! */ + for (i = 0; i < TEST_SIZE; i++) + { + verify_equiv_values(vecSize, &inDataA[i * vecSize], + &inDataB[i * vecSize], expected, + param.verifyFn); + + for (j = 0; j < (int)vecSize; j++) + { + if (expected[j] != outData[i * vecSize + j]) + { + bool acceptFail = true; + if (std::is_same::value) + { + bool in_denorm = IsHalfSubnormal(inDataA[i * vecSize + j]) + || IsHalfSubnormal(inDataB[i * vecSize + j]); + + if (halfFlushDenormsToZero && in_denorm) + { + acceptFail = false; + } + } + + if (acceptFail) + { + verror_msg( + i, j, vecSize, expected[j], outData[i * vecSize + j], + inDataA[i * vecSize + j], inDataB[i * vecSize + j]); + return -1; + } + } + } + } + + /* Now get the results */ + error = clEnqueueReadBuffer(queue, streams[3], true, 0, + sizeof(U) * TEST_SIZE * vecSize, outData, 0, + NULL, NULL); + test_error(error, "Unable to read output array!"); + + /* And verify! */ + int fail = 0; + for (i = 0; i < TEST_SIZE; i++) + { + verify_equiv_values(vecSize, &inDataA[i * vecSize], + &inDataB[i * vecSize], expected, + param.verifyFn); + + for (j = 0; j < (int)vecSize; j++) + { + if (expected[j] != outData[i * vecSize + j]) + { + if (std::is_same::value) + { + if (gInfNanSupport == 0) + { + if (isnan(inDataA[i * vecSize + j]) + || isnan(inDataB[i * vecSize + j])) + fail = 0; + else + fail = 1; + } + if (fail) + { + verror_msg(i, j, vecSize, expected[j], + outData[i * vecSize + j], + inDataA[i * vecSize + j], + inDataB[i * vecSize + j]); + return -1; + } + } + else if (std::is_same::value) + { + bool in_denorm = IsHalfSubnormal(inDataA[i * vecSize + j]) + || IsHalfSubnormal(inDataB[i * vecSize + j]); + + if (!(halfFlushDenormsToZero && in_denorm)) + { + verror_msg(i, j, vecSize, expected[j], + outData[i * vecSize + j], + inDataA[i * vecSize + j], + inDataB[i * vecSize + j]); + return -1; + } + } + else + { + verror_msg( + i, j, vecSize, expected[j], outData[i * vecSize + j], + inDataA[i * vecSize + j], inDataB[i * vecSize + j]); + return -1; + } + } + } + } + return 0; +} + +template +int RelationalsFPTest::test_relational(int numElements, + const RelTestParams& param) +{ + RandomSeed seed(gRandomSeed); + unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 }; + unsigned int index; + int retVal = 0; + + for (index = 0; vecSizes[index] != 0; index++) + { + // Test! + if (test_equiv_kernel(vecSizes[index], param, seed) != 0) + { + log_error(" Vector %s%d FAILED\n", ftype, vecSizes[index]); + retVal = -1; + } + } + return retVal; +} + +cl_int RelationalsFPTest::SetUp(int elements) +{ + if (is_extension_available(device, "cl_khr_fp16")) + { + cl_device_fp_config config = 0; + cl_int error = clGetDeviceInfo(device, CL_DEVICE_HALF_FP_CONFIG, + sizeof(config), &config, NULL); + test_error(error, "Unable to get device CL_DEVICE_HALF_FP_CONFIG"); + + halfFlushDenormsToZero = (0 == (config & CL_FP_DENORM)); + log_info("Supports half precision denormals: %s\n", + halfFlushDenormsToZero ? "NO" : "YES"); + } + + return CL_SUCCESS; +} + +cl_int RelationalsFPTest::Run() +{ + cl_int error = CL_SUCCESS; + for (auto&& param : params) + { + switch (param->dataType) + { + case kHalf: + error = test_relational( + num_elements, *((RelTestParams*)param.get())); + break; + case kFloat: + error = test_relational( + num_elements, *((RelTestParams*)param.get())); + break; + case kDouble: + error = test_relational( + num_elements, *((RelTestParams*)param.get())); + break; + default: + test_error(-1, "RelationalsFPTest::Run: incorrect fp type"); + break; + } + test_error(error, "RelationalsFPTest::Run: test_relational failed"); + } + return CL_SUCCESS; +} + +cl_int IsEqualFPTest::SetUp(int elements) +{ + num_elements = elements; + if (is_extension_available(device, "cl_khr_fp16")) + params.emplace_back(new RelTestParams( + &verify, kHalf, HALF_NAN)); + + params.emplace_back(new RelTestParams( + &verify>, kFloat, NAN)); + + if (is_extension_available(device, "cl_khr_fp64")) + params.emplace_back(new RelTestParams( + &verify>, kDouble, NAN)); + + return RelationalsFPTest::SetUp(elements); +} + +cl_int IsNotEqualFPTest::SetUp(int elements) +{ + num_elements = elements; + if (is_extension_available(device, "cl_khr_fp16")) + params.emplace_back(new RelTestParams( + &verify, kHalf, HALF_NAN)); + + params.emplace_back(new RelTestParams( + &verify>, kFloat, NAN)); + + if (is_extension_available(device, "cl_khr_fp64")) + params.emplace_back(new RelTestParams( + &verify>, kDouble, NAN)); + + return RelationalsFPTest::SetUp(elements); +} + +cl_int IsGreaterFPTest::SetUp(int elements) +{ + num_elements = elements; + if (is_extension_available(device, "cl_khr_fp16")) + params.emplace_back(new RelTestParams( + &verify, kHalf, HALF_NAN)); + + params.emplace_back(new RelTestParams( + &verify>, kFloat, NAN)); + + if (is_extension_available(device, "cl_khr_fp64")) + params.emplace_back(new RelTestParams( + &verify>, kDouble, NAN)); + + return RelationalsFPTest::SetUp(elements); +} + +cl_int IsGreaterEqualFPTest::SetUp(int elements) +{ + num_elements = elements; + if (is_extension_available(device, "cl_khr_fp16")) + params.emplace_back(new RelTestParams( + &verify, kHalf, HALF_NAN)); + + params.emplace_back(new RelTestParams( + &verify>, kFloat, NAN)); + + if (is_extension_available(device, "cl_khr_fp64")) + params.emplace_back(new RelTestParams( + &verify>, kDouble, NAN)); + + return RelationalsFPTest::SetUp(elements); +} + +cl_int IsLessFPTest::SetUp(int elements) +{ + num_elements = elements; + if (is_extension_available(device, "cl_khr_fp16")) + params.emplace_back(new RelTestParams( + &verify, kHalf, HALF_NAN)); + + params.emplace_back(new RelTestParams( + &verify>, kFloat, NAN)); + + if (is_extension_available(device, "cl_khr_fp64")) + params.emplace_back(new RelTestParams( + &verify>, kDouble, NAN)); + + return RelationalsFPTest::SetUp(elements); +} + +cl_int IsLessEqualFPTest::SetUp(int elements) +{ + num_elements = elements; + if (is_extension_available(device, "cl_khr_fp16")) + params.emplace_back(new RelTestParams( + &verify, kHalf, HALF_NAN)); + + params.emplace_back(new RelTestParams( + &verify>, kFloat, NAN)); + + if (is_extension_available(device, "cl_khr_fp64")) + params.emplace_back(new RelTestParams( + &verify>, kDouble, NAN)); + + return RelationalsFPTest::SetUp(elements); +} + +cl_int IsLessGreaterFPTest::SetUp(int elements) +{ + num_elements = elements; + if (is_extension_available(device, "cl_khr_fp16")) + params.emplace_back(new RelTestParams( + &verify, kHalf, HALF_NAN)); + + params.emplace_back(new RelTestParams( + &verify>, kFloat, NAN)); + + if (is_extension_available(device, "cl_khr_fp64")) + params.emplace_back(new RelTestParams( + &verify>, kDouble, NAN)); + + return RelationalsFPTest::SetUp(elements); +} + +int test_relational_isequal(cl_device_id device, cl_context context, + cl_command_queue queue, int numElements) +{ + return MakeAndRunTest(device, context, queue, numElements); +} + +int test_relational_isnotequal(cl_device_id device, cl_context context, + cl_command_queue queue, int numElements) +{ + return MakeAndRunTest(device, context, queue, + numElements); +} + +int test_relational_isgreater(cl_device_id device, cl_context context, + cl_command_queue queue, int numElements) +{ + return MakeAndRunTest(device, context, queue, numElements); +} + +int test_relational_isgreaterequal(cl_device_id device, cl_context context, + cl_command_queue queue, int numElements) +{ + return MakeAndRunTest(device, context, queue, + numElements); +} + +int test_relational_isless(cl_device_id device, cl_context context, + cl_command_queue queue, int numElements) +{ + return MakeAndRunTest(device, context, queue, numElements); +} + +int test_relational_islessequal(cl_device_id device, cl_context context, + cl_command_queue queue, int numElements) +{ + return MakeAndRunTest(device, context, queue, + numElements); +} + +int test_relational_islessgreater(cl_device_id device, cl_context context, + cl_command_queue queue, int numElements) +{ + return MakeAndRunTest(device, context, queue, + numElements); +} diff --git a/test_conformance/relationals/test_comparisons_fp.h b/test_conformance/relationals/test_comparisons_fp.h new file mode 100644 index 00000000..7faca1c5 --- /dev/null +++ b/test_conformance/relationals/test_comparisons_fp.h @@ -0,0 +1,227 @@ +// +// Copyright (c) 2022 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + +#ifndef _TEST_COMPARISONS_FP_H +#define _TEST_COMPARISONS_FP_H + +#include +#include +#include +#include + +#include + +#include "testBase.h" + +#define HALF_NAN 0x7e00 +template using VerifyFunc = bool (*)(const T &, const T &); + +struct RelTestBase +{ + explicit RelTestBase(const ExplicitTypes &dt): dataType(dt) {} + ExplicitTypes dataType; +}; + +template struct RelTestParams : public RelTestBase +{ + RelTestParams(const VerifyFunc &vfn, const ExplicitTypes &dt, + const T &nan_) + : RelTestBase(dt), verifyFn(vfn), nan(nan_) + {} + + VerifyFunc verifyFn; + T nan; +}; + +struct RelationalsFPTest +{ + RelationalsFPTest(cl_context context, cl_device_id device, + cl_command_queue queue, const char *fn, const char *op); + + virtual cl_int SetUp(int elements); + + // Test body returning an OpenCL error code + virtual cl_int Run(); + + template + void generate_equiv_test_data(T *, unsigned int, bool, + const RelTestParams &, const MTdata &); + + template + void verify_equiv_values(unsigned int, const T *const, const T *const, + U *const, const VerifyFunc &); + + template + int test_equiv_kernel(unsigned int vecSize, const RelTestParams ¶m, + const MTdata &d); + + template + int test_relational(int numElements, const RelTestParams ¶m); + +protected: + cl_context context; + cl_device_id device; + cl_command_queue queue; + + std::string fnName; + std::string opName; + + std::vector> params; + std::map eqTypeNames; + size_t num_elements; + + int halfFlushDenormsToZero; +}; + +struct IsEqualFPTest : public RelationalsFPTest +{ + IsEqualFPTest(cl_device_id d, cl_context c, cl_command_queue q) + : RelationalsFPTest(c, d, q, "isequal", "==") + {} + cl_int SetUp(int elements) override; + + // for correct handling nan/inf we need fp value + struct half_equals_to + { + bool operator()(const cl_half &lhs, const cl_half &rhs) const + { + return cl_half_to_float(lhs) == cl_half_to_float(rhs); + } + }; +}; + +struct IsNotEqualFPTest : public RelationalsFPTest +{ + IsNotEqualFPTest(cl_device_id d, cl_context c, cl_command_queue q) + : RelationalsFPTest(c, d, q, "isnotequal", "!=") + {} + cl_int SetUp(int elements) override; + + // for correct handling nan/inf we need fp value + struct half_not_equals_to + { + bool operator()(const cl_half &lhs, const cl_half &rhs) const + { + return cl_half_to_float(lhs) != cl_half_to_float(rhs); + } + }; +}; + +struct IsGreaterFPTest : public RelationalsFPTest +{ + IsGreaterFPTest(cl_device_id d, cl_context c, cl_command_queue q) + : RelationalsFPTest(c, d, q, "isgreater", ">") + {} + cl_int SetUp(int elements) override; + + struct half_greater + { + bool operator()(const cl_half &lhs, const cl_half &rhs) const + { + return cl_half_to_float(lhs) > cl_half_to_float(rhs); + } + }; +}; + +struct IsGreaterEqualFPTest : public RelationalsFPTest +{ + IsGreaterEqualFPTest(cl_device_id d, cl_context c, cl_command_queue q) + : RelationalsFPTest(c, d, q, "isgreaterequal", ">=") + {} + cl_int SetUp(int elements) override; + + struct half_greater_equal + { + bool operator()(const cl_half &lhs, const cl_half &rhs) const + { + return cl_half_to_float(lhs) >= cl_half_to_float(rhs); + } + }; +}; + +struct IsLessFPTest : public RelationalsFPTest +{ + IsLessFPTest(cl_device_id d, cl_context c, cl_command_queue q) + : RelationalsFPTest(c, d, q, "isless", "<") + {} + cl_int SetUp(int elements) override; + + struct half_less + { + bool operator()(const cl_half &lhs, const cl_half &rhs) const + { + return cl_half_to_float(lhs) < cl_half_to_float(rhs); + } + }; +}; + +struct IsLessEqualFPTest : public RelationalsFPTest +{ + IsLessEqualFPTest(cl_device_id d, cl_context c, cl_command_queue q) + : RelationalsFPTest(c, d, q, "islessequal", "<=") + {} + cl_int SetUp(int elements) override; + + struct half_less_equal + { + bool operator()(const cl_half &lhs, const cl_half &rhs) const + { + return cl_half_to_float(lhs) <= cl_half_to_float(rhs); + } + }; +}; + +struct IsLessGreaterFPTest : public RelationalsFPTest +{ + IsLessGreaterFPTest(cl_device_id d, cl_context c, cl_command_queue q) + : RelationalsFPTest(c, d, q, "islessgreater", "<>") + {} + cl_int SetUp(int elements) override; + + struct half_less_greater + { + bool operator()(const cl_half &lhs, const cl_half &rhs) const + { + float flhs = cl_half_to_float(lhs), frhs = cl_half_to_float(rhs); + return (flhs < frhs) || (flhs > frhs); + } + }; + + template struct less_greater + { + bool operator()(const T &lhs, const T &rhs) const + { + return (lhs < rhs) || (lhs > rhs); + } + }; +}; + +template +int MakeAndRunTest(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) +{ + auto test_fixture = T(device, context, queue); + + cl_int error = test_fixture.SetUp(num_elements); + test_error_ret(error, "Error in test initialization", TEST_FAIL); + + error = test_fixture.Run(); + test_error_ret(error, "Test Failed", TEST_FAIL); + + return TEST_PASS; +} + +#endif // _TEST_COMPARISONS_FP_H diff --git a/test_conformance/relationals/test_relationals.cpp b/test_conformance/relationals/test_relationals.cpp index 5a874af7..d744fb2a 100644 --- a/test_conformance/relationals/test_relationals.cpp +++ b/test_conformance/relationals/test_relationals.cpp @@ -18,8 +18,11 @@ #include "harness/typeWrappers.h" #include "harness/testHarness.h" +// clang-format off + const char *anyAllTestKernelPattern = "%s\n" // optional pragma +"%s\n" // optional pragma "__kernel void sample_test(__global %s%s *sourceA, __global int *destValues)\n" "{\n" " int tid = get_global_id(0);\n" @@ -29,6 +32,7 @@ const char *anyAllTestKernelPattern = const char *anyAllTestKernelPatternVload = "%s\n" // optional pragma +"%s\n" // optional pragma "__kernel void sample_test(__global %s%s *sourceA, __global int *destValues)\n" "{\n" " int tid = get_global_id(0);\n" @@ -36,6 +40,8 @@ const char *anyAllTestKernelPatternVload = "\n" "}\n"; +// clang-format on + #define TEST_SIZE 512 typedef int (*anyAllVerifyFn)( ExplicitType vecType, unsigned int vecSize, void *inData ); @@ -67,14 +73,22 @@ int test_any_all_kernel(cl_context context, cl_command_queue queue, get_explicit_type_name( vecType ), sizeName); if(DENSE_PACK_VECS && vecSize == 3) { // anyAllTestKernelPatternVload - sprintf( kernelSource, anyAllTestKernelPatternVload, - vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", - get_explicit_type_name( vecType ), sizeName, fnName, - get_explicit_type_name(vecType)); + sprintf( + kernelSource, anyAllTestKernelPatternVload, + vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" + : "", + vecType == kHalf ? "#pragma OPENCL EXTENSION cl_khr_fp16 : enable" + : "", + get_explicit_type_name(vecType), sizeName, fnName, + get_explicit_type_name(vecType)); } else { - sprintf( kernelSource, anyAllTestKernelPattern, - vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", - get_explicit_type_name( vecType ), sizeName, fnName ); + sprintf( + kernelSource, anyAllTestKernelPattern, + vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" + : "", + vecType == kHalf ? "#pragma OPENCL EXTENSION cl_khr_fp16 : enable" + : "", + get_explicit_type_name(vecType), sizeName, fnName); } /* Create kernels */ programPtr = kernelSource; @@ -282,8 +296,11 @@ int test_relational_all(cl_device_id device, cl_context context, cl_command_queu return retVal; } +// clang-format off + const char *selectTestKernelPattern = "%s\n" // optional pragma +"%s\n" // optional pragma "__kernel void sample_test(__global %s%s *sourceA, __global %s%s *sourceB, __global %s%s *sourceC, __global %s%s *destValues)\n" "{\n" " int tid = get_global_id(0);\n" @@ -294,6 +311,7 @@ const char *selectTestKernelPattern = const char *selectTestKernelPatternVload = "%s\n" // optional pragma +"%s\n" // optional pragma "__kernel void sample_test(__global %s%s *sourceA, __global %s%s *sourceB, __global %s%s *sourceC, __global %s%s *destValues)\n" "{\n" " int tid = get_global_id(0);\n" @@ -302,6 +320,8 @@ const char *selectTestKernelPatternVload = "\n" "}\n"; +// clang-format on + typedef void (*selectVerifyFn)( ExplicitType vecType, ExplicitType testVecType, unsigned int vecSize, void *inDataA, void *inDataB, void *inDataTest, void *outData ); int test_select_kernel(cl_context context, cl_command_queue queue, const char *fnName, @@ -335,26 +355,34 @@ int test_select_kernel(cl_context context, cl_command_queue queue, const char *f if(DENSE_PACK_VECS && vecSize == 3) { // anyAllTestKernelPatternVload - sprintf( kernelSource, selectTestKernelPatternVload, - (vecType == kDouble || testVecType == kDouble) ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", - get_explicit_type_name( vecType ), sizeName, - get_explicit_type_name( vecType ), sizeName, - get_explicit_type_name( testVecType ), sizeName, - get_explicit_type_name( vecType ), outSizeName, - get_explicit_type_name( vecType ), sizeName, - fnName, - get_explicit_type_name( vecType ), - get_explicit_type_name( vecType ), - get_explicit_type_name( vecType ), - get_explicit_type_name( testVecType ) ); + sprintf(kernelSource, selectTestKernelPatternVload, + (vecType == kDouble || testVecType == kDouble) + ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" + : "", + (vecType == kHalf || testVecType == kHalf) + ? "#pragma OPENCL EXTENSION cl_khr_fp16 : enable" + : "", + get_explicit_type_name(vecType), sizeName, + get_explicit_type_name(vecType), sizeName, + get_explicit_type_name(testVecType), sizeName, + get_explicit_type_name(vecType), outSizeName, + get_explicit_type_name(vecType), sizeName, fnName, + get_explicit_type_name(vecType), + get_explicit_type_name(vecType), + get_explicit_type_name(vecType), + get_explicit_type_name(testVecType)); } else { - sprintf( kernelSource, selectTestKernelPattern, - (vecType == kDouble || testVecType == kDouble) ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", - get_explicit_type_name( vecType ), sizeName, - get_explicit_type_name( vecType ), sizeName, - get_explicit_type_name( testVecType ), sizeName, - get_explicit_type_name( vecType ), outSizeName, - fnName ); + sprintf(kernelSource, selectTestKernelPattern, + (vecType == kDouble || testVecType == kDouble) + ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" + : "", + (vecType == kHalf || testVecType == kHalf) + ? "#pragma OPENCL EXTENSION cl_khr_fp16 : enable" + : "", + get_explicit_type_name(vecType), sizeName, + get_explicit_type_name(vecType), sizeName, + get_explicit_type_name(testVecType), sizeName, + get_explicit_type_name(vecType), outSizeName, fnName); } /* Create kernels */ @@ -500,14 +528,17 @@ void bitselect_verify_fn( ExplicitType vecType, ExplicitType testVecType, unsign int test_relational_bitselect(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { - ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kDouble }; + constexpr ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, + kInt, kUInt, kLong, kULong, + kHalf, kFloat, kDouble }; + constexpr auto vecTypeSize = sizeof(vecType) / sizeof(ExplicitType); unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 }; unsigned int index, typeIndex; int retVal = 0; RandomSeed seed( gRandomSeed ); - for( typeIndex = 0; typeIndex < 10; typeIndex++ ) + for (typeIndex = 0; typeIndex < vecTypeSize; typeIndex++) { if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong) && !gHasLong) continue; @@ -522,6 +553,19 @@ int test_relational_bitselect(cl_device_id device, cl_context context, cl_comman else log_info("Testing doubles.\n"); } + + if (vecType[typeIndex] == kHalf) + { + if (!is_extension_available(device, "cl_khr_fp16")) + { + log_info("Extension cl_khr_fp16 not supported; skipping half " + "tests.\n"); + continue; + } + else + log_info("Testing halfs.\n"); + } + for( index = 0; vecSizes[ index ] != 0; index++ ) { // Test! @@ -584,14 +628,18 @@ void select_signed_verify_fn( ExplicitType vecType, ExplicitType testVecType, un int test_relational_select_signed(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { - ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kDouble }; + constexpr ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, + kInt, kUInt, kLong, kULong, + kHalf, kFloat, kDouble }; + constexpr auto vecTypeSize = sizeof(vecType) / sizeof(ExplicitType); + ExplicitType testVecType[] = { kChar, kShort, kInt, kLong, kNumExplicitTypes }; unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 }; unsigned int index, typeIndex, testTypeIndex; int retVal = 0; RandomSeed seed( gRandomSeed ); - for( typeIndex = 0; typeIndex < 10; typeIndex++ ) + for (typeIndex = 0; typeIndex < vecTypeSize; typeIndex++) { if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong) && !gHasLong) continue; @@ -604,6 +652,19 @@ int test_relational_select_signed(cl_device_id device, cl_context context, cl_co log_info("Testing doubles.\n"); } } + if (vecType[typeIndex] == kHalf) + { + if (!is_extension_available(device, "cl_khr_fp16")) + { + log_info("Extension cl_khr_fp16 not supported; skipping half " + "tests.\n"); + continue; + } + else + { + log_info("Testing halfs.\n"); + } + } for( testTypeIndex = 0; testVecType[ testTypeIndex ] != kNumExplicitTypes; testTypeIndex++ ) { if( testVecType[ testTypeIndex ] != vecType[ typeIndex ] ) @@ -673,7 +734,11 @@ void select_unsigned_verify_fn( ExplicitType vecType, ExplicitType testVecType, int test_relational_select_unsigned(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { - ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kDouble }; + constexpr ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, + kInt, kUInt, kLong, kULong, + kHalf, kFloat, kDouble }; + constexpr auto vecTypeSize = sizeof(vecType) / sizeof(ExplicitType); + ExplicitType testVecType[] = { kUChar, kUShort, kUInt, kULong, kNumExplicitTypes }; unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 }; unsigned int index, typeIndex, testTypeIndex; @@ -681,7 +746,7 @@ int test_relational_select_unsigned(cl_device_id device, cl_context context, cl_ RandomSeed seed(gRandomSeed); - for( typeIndex = 0; typeIndex < 10; typeIndex++ ) + for (typeIndex = 0; typeIndex < vecTypeSize; typeIndex++) { if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong) && !gHasLong) continue; @@ -694,6 +759,19 @@ int test_relational_select_unsigned(cl_device_id device, cl_context context, cl_ log_info("Testing doubles.\n"); } } + if (vecType[typeIndex] == kHalf) + { + if (!is_extension_available(device, "cl_khr_fp16")) + { + log_info("Extension cl_khr_fp16 not supported; skipping half " + "tests.\n"); + continue; + } + else + { + log_info("Testing halfs.\n"); + } + } for( testTypeIndex = 0; testVecType[ testTypeIndex ] != kNumExplicitTypes; testTypeIndex++ ) { if( testVecType[ testTypeIndex ] != vecType[ typeIndex ] ) @@ -714,85 +792,3 @@ int test_relational_select_unsigned(cl_device_id device, cl_context context, cl_ return retVal; } - - - -extern int test_relational_isequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_isnotequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_isgreater_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_isgreaterequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_isless_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_islessequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_islessgreater_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_isequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_isnotequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_isgreater_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_isgreaterequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_isless_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_islessequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); -extern int test_relational_islessgreater_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); - - -int test_relational_isequal(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - int err = 0; - err |= test_relational_isequal_float( device, context, queue, numElements ); - err |= test_relational_isequal_double( device, context, queue, numElements ); - return err; -} - - -int test_relational_isnotequal(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - int err = 0; - err |= test_relational_isnotequal_float( device, context, queue, numElements ); - err |= test_relational_isnotequal_double( device, context, queue, numElements ); - return err; -} - - -int test_relational_isgreater(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - int err = 0; - err |= test_relational_isgreater_float( device, context, queue, numElements ); - err |= test_relational_isgreater_double( device, context, queue, numElements ); - return err; -} - - -int test_relational_isgreaterequal(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - int err = 0; - err |= test_relational_isgreaterequal_float( device, context, queue, numElements ); - err |= test_relational_isgreaterequal_double( device, context, queue, numElements ); - return err; -} - - -int test_relational_isless(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - int err = 0; - err |= test_relational_isless_float( device, context, queue, numElements ); - err |= test_relational_isless_double( device, context, queue, numElements ); - return err; -} - - -int test_relational_islessequal(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - int err = 0; - err |= test_relational_islessequal_float( device, context, queue, numElements ); - err |= test_relational_islessequal_double( device, context, queue, numElements ); - return err; -} - - -int test_relational_islessgreater(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) -{ - int err = 0; - err |= test_relational_islessgreater_float( device, context, queue, numElements ); - err |= test_relational_islessgreater_double( device, context, queue, numElements ); - return err; -} - -