ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Added cl_half support for test_relationals (#1623)

Browse Source 
* 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)
This commit is contained in:
Marcin Hajder
2023-05-16 17:43:47 +02:00
committed by GitHub
parent 1884042f5d
commit f31b2f029c
7 changed files with 1008 additions and 845 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
test_conformance/printf/test_printf.cpp
View File

@@ -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

3
test_conformance/relationals/CMakeLists.txt
View File

@@ -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
)

363
test_conformance/relationals/test_comparisons_double.cpp
View File

@@ -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 );
}

362
test_conformance/relationals/test_comparisons_float.cpp
View File

@@ -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 );
}

661
test_conformance/relationals/test_comparisons_fp.cpp Normal file
View File

@@ -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 <iostream>
#include <map>
#include <memory>
#include <stdexcept>
#include <vector>
#include <CL/cl_half.h>
#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 <typename... Args>
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_t>(size_s);
std::unique_ptr<char[]> 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 <typename T, typename F> 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 <typename T>
void RelationalsFPTest::generate_equiv_test_data(T* outData,
unsigned int vecSize,
bool alpha,
const RelTestParams<T>& 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 <typename T, typename U>
void RelationalsFPTest::verify_equiv_values(unsigned int vecSize,
const T* const inDataA,
const T* const inDataB,
U* const outData,
const VerifyFunc<T>& 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 <typename T>
int RelationalsFPTest::test_equiv_kernel(unsigned int vecSize,
const RelTestParams<T>& 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<T, double>::value)
strcpy(extension, "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n");
else if (std::is_same<T, cl_half>::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<T>(inDataA, vecSize, true, param, d);
generate_equiv_test_data<T>(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<T, U>(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<T, cl_half>::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<T, U>(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<T, float>::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<T, cl_half>::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 <typename T>
int RelationalsFPTest::test_relational(int numElements,
const RelTestParams<T>& 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<T>(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<cl_half>(
num_elements, *((RelTestParams<cl_half>*)param.get()));
break;
case kFloat:
error = test_relational<float>(
num_elements, *((RelTestParams<float>*)param.get()));
break;
case kDouble:
error = test_relational<double>(
num_elements, *((RelTestParams<double>*)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<cl_half>(
&verify<cl_half, half_equals_to>, kHalf, HALF_NAN));
params.emplace_back(new RelTestParams<float>(
&verify<float, std::equal_to<float>>, kFloat, NAN));
if (is_extension_available(device, "cl_khr_fp64"))
params.emplace_back(new RelTestParams<double>(
&verify<double, std::equal_to<double>>, 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<cl_half>(
&verify<cl_half, half_not_equals_to>, kHalf, HALF_NAN));
params.emplace_back(new RelTestParams<float>(
&verify<float, std::not_equal_to<float>>, kFloat, NAN));
if (is_extension_available(device, "cl_khr_fp64"))
params.emplace_back(new RelTestParams<double>(
&verify<double, std::not_equal_to<double>>, 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<cl_half>(
&verify<cl_half, half_greater>, kHalf, HALF_NAN));
params.emplace_back(new RelTestParams<float>(
&verify<float, std::greater<float>>, kFloat, NAN));
if (is_extension_available(device, "cl_khr_fp64"))
params.emplace_back(new RelTestParams<double>(
&verify<double, std::greater<double>>, 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<cl_half>(
&verify<cl_half, half_greater_equal>, kHalf, HALF_NAN));
params.emplace_back(new RelTestParams<float>(
&verify<float, std::greater_equal<float>>, kFloat, NAN));
if (is_extension_available(device, "cl_khr_fp64"))
params.emplace_back(new RelTestParams<double>(
&verify<double, std::greater_equal<double>>, 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<cl_half>(
&verify<cl_half, half_less>, kHalf, HALF_NAN));
params.emplace_back(new RelTestParams<float>(
&verify<float, std::less<float>>, kFloat, NAN));
if (is_extension_available(device, "cl_khr_fp64"))
params.emplace_back(new RelTestParams<double>(
&verify<double, std::less<double>>, 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<cl_half>(
&verify<cl_half, half_less_equal>, kHalf, HALF_NAN));
params.emplace_back(new RelTestParams<float>(
&verify<float, std::less_equal<float>>, kFloat, NAN));
if (is_extension_available(device, "cl_khr_fp64"))
params.emplace_back(new RelTestParams<double>(
&verify<double, std::less_equal<double>>, 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<cl_half>(
&verify<cl_half, half_less_greater>, kHalf, HALF_NAN));
params.emplace_back(new RelTestParams<float>(
&verify<float, less_greater<float>>, kFloat, NAN));
if (is_extension_available(device, "cl_khr_fp64"))
params.emplace_back(new RelTestParams<double>(
&verify<double, less_greater<double>>, 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<IsEqualFPTest>(device, context, queue, numElements);
}
int test_relational_isnotequal(cl_device_id device, cl_context context,
cl_command_queue queue, int numElements)
{
return MakeAndRunTest<IsNotEqualFPTest>(device, context, queue,
numElements);
}
int test_relational_isgreater(cl_device_id device, cl_context context,
cl_command_queue queue, int numElements)
{
return MakeAndRunTest<IsGreaterFPTest>(device, context, queue, numElements);
}
int test_relational_isgreaterequal(cl_device_id device, cl_context context,
cl_command_queue queue, int numElements)
{
return MakeAndRunTest<IsGreaterEqualFPTest>(device, context, queue,
numElements);
}
int test_relational_isless(cl_device_id device, cl_context context,
cl_command_queue queue, int numElements)
{
return MakeAndRunTest<IsLessFPTest>(device, context, queue, numElements);
}
int test_relational_islessequal(cl_device_id device, cl_context context,
cl_command_queue queue, int numElements)
{
return MakeAndRunTest<IsLessEqualFPTest>(device, context, queue,
numElements);
}
int test_relational_islessgreater(cl_device_id device, cl_context context,
cl_command_queue queue, int numElements)
{
return MakeAndRunTest<IsLessGreaterFPTest>(device, context, queue,
numElements);
}

227
test_conformance/relationals/test_comparisons_fp.h Normal file
View File

@@ -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 <map>
#include <memory>
#include <string>
#include <vector>
#include <CL/cl_half.h>
#include "testBase.h"
#define HALF_NAN 0x7e00
template <typename T> using VerifyFunc = bool (*)(const T &, const T &);
struct RelTestBase
{
explicit RelTestBase(const ExplicitTypes &dt): dataType(dt) {}
ExplicitTypes dataType;
};
template <typename T> struct RelTestParams : public RelTestBase
{
RelTestParams(const VerifyFunc<T> &vfn, const ExplicitTypes &dt,
const T &nan_)
: RelTestBase(dt), verifyFn(vfn), nan(nan_)
{}
VerifyFunc<T> 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 <typename T>
void generate_equiv_test_data(T *, unsigned int, bool,
const RelTestParams<T> &, const MTdata &);
template <typename T, typename U>
void verify_equiv_values(unsigned int, const T *const, const T *const,
U *const, const VerifyFunc<T> &);
template <typename T>
int test_equiv_kernel(unsigned int vecSize, const RelTestParams<T> &param,
const MTdata &d);
template <typename T>
int test_relational(int numElements, const RelTestParams<T> &param);
protected:
cl_context context;
cl_device_id device;
cl_command_queue queue;
std::string fnName;
std::string opName;
std::vector<std::unique_ptr<RelTestBase>> params;
std::map<ExplicitTypes, std::string> 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 <typename T> struct less_greater
{
bool operator()(const T &lhs, const T &rhs) const
{
return (lhs < rhs) || (lhs > rhs);
}
};
};
template <class T>
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

224
test_conformance/relationals/test_relationals.cpp
View File

@@ -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;
}