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

Remove gTestFastRelaxed Dependencies in Brute Force (#807)

Browse Source 
* The global variable `gTestFastRelaxed` has state which is used to
control the behaviour of the compiler flag `-cl-fast-relaxed-math` and
the precision testing of relaxed, fp32 and fp64 types. This is confusing
since the global variable is being set and read in different translation
units, making it very difficult to reason about the logic of the brute
force framework. It is particular difficult to follow since the global
variables is cached and then turned off in the case of fp32 and f64 in
order to use the same code path as relaxed testing, after it is then
turned back on.

* Remove uses of the global variable outside of `main.cpp` (the global
variable remains in use within `main.cpp` since it is a command line
option and used to turn of relaxed testing completely). Replace all uses
of the global variable with boolean `relaxedMode` which is passed as a
function paramter but replaces `gTestFastRelaxed` semantically.
This commit is contained in:
Jack Frankland
2020-06-04 10:44:39 +01:00
committed by GitHub
parent 519f658b51
commit ecee2c18d2
16 changed files with 480 additions and 288 deletions
Download Patch File Download Diff File Expand all files Collapse all files

94
test_conformance/math_brute_force/binary.cpp
View File

@@ -18,12 +18,16 @@
#include <string.h>
#include "FunctionList.h"
int TestFunc_Float_Float_Float(const Func *f, MTdata);
int TestFunc_Double_Double_Double(const Func *f, MTdata);
int TestFunc_Float_Float_Float_nextafter(const Func *f, MTdata);
int TestFunc_Double_Double_Double_nextafter(const Func *f, MTdata);
int TestFunc_Float_Float_Float_common(const Func *f, MTdata, int isNextafter);
int TestFunc_Double_Double_Double_common(const Func *f, MTdata, int isNextafter);
int TestFunc_Float_Float_Float(const Func *f, MTdata, bool relaxedMode);
int TestFunc_Double_Double_Double(const Func *f, MTdata, bool relaxedMode);
int TestFunc_Float_Float_Float_nextafter(const Func *f, MTdata,
bool relaxedMode);
int TestFunc_Double_Double_Double_nextafter(const Func *f, MTdata,
bool relaxedMode);
int TestFunc_Float_Float_Float_common(const Func *f, MTdata, int isNextafter,
bool relaxedMode);
int TestFunc_Double_Double_Double_common(const Func *f, MTdata, int isNextafter,
bool relaxedMode);
const float twoToMinus126 = MAKE_HEX_FLOAT(0x1p-126f, 1, -126);
const double twoToMinus1022 = MAKE_HEX_DOUBLE(0x1p-1022, 1, -1022);
@@ -36,9 +40,11 @@ extern const vtbl _binary_nextafter = {
TestFunc_Double_Double_Double_nextafter
};
static int BuildKernel( const char *name, int vectorSize, cl_uint kernel_count, cl_kernel *k, cl_program *p );
static int BuildKernel(const char *name, int vectorSize, cl_uint kernel_count,
cl_kernel *k, cl_program *p, bool relaxedMode);
static int BuildKernel( const char *name, int vectorSize, cl_uint kernel_count, cl_kernel *k, cl_program *p )
static int BuildKernel(const char *name, int vectorSize, cl_uint kernel_count,
cl_kernel *k, cl_program *p, bool relaxedMode)
{
const char *c[] = { "__kernel void math_kernel", sizeNames[vectorSize], "( __global float", sizeNames[vectorSize], "* out, __global float", sizeNames[vectorSize], "* in1, __global float", sizeNames[vectorSize], "* in2 )\n"
"{\n"
@@ -98,10 +104,13 @@ static int BuildKernel( const char *name, int vectorSize, cl_uint kernel_count,
char testName[32];
snprintf( testName, sizeof( testName ) -1, "math_kernel%s", sizeNames[vectorSize] );
return MakeKernels(kern, (cl_uint) kernSize, testName, kernel_count, k, p);
return MakeKernels(kern, (cl_uint)kernSize, testName, kernel_count, k, p,
relaxedMode);
}
static int BuildKernelDouble( const char *name, int vectorSize, cl_uint kernel_count, cl_kernel *k, cl_program *p )
static int BuildKernelDouble(const char *name, int vectorSize,
cl_uint kernel_count, cl_kernel *k, cl_program *p,
bool relaxedMode)
{
const char *c[] = { "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
"__kernel void math_kernel", sizeNames[vectorSize], "( __global double", sizeNames[vectorSize], "* out, __global double", sizeNames[vectorSize], "* in1, __global double", sizeNames[vectorSize], "* in2 )\n"
@@ -163,7 +172,8 @@ static int BuildKernelDouble( const char *name, int vectorSize, cl_uint kernel_c
char testName[32];
snprintf( testName, sizeof( testName ) -1, "math_kernel%s", sizeNames[vectorSize] );
return MakeKernels(kern, (cl_uint) kernSize, testName, kernel_count, k, p);
return MakeKernels(kern, (cl_uint)kernSize, testName, kernel_count, k, p,
relaxedMode);
}
// A table of more difficult cases to get right
@@ -192,6 +202,7 @@ typedef struct BuildKernelInfo
cl_kernel **kernels;
cl_program *programs;
const char *nameInCode;
bool relaxedMode; // Whether to build with -cl-fast-relaxed-math.
}BuildKernelInfo;
static cl_int BuildKernel_FloatFn( cl_uint job_id, cl_uint thread_id UNUSED, void *p );
@@ -199,7 +210,8 @@ static cl_int BuildKernel_FloatFn( cl_uint job_id, cl_uint thread_id UNUSED, voi
{
BuildKernelInfo *info = (BuildKernelInfo*) p;
cl_uint i = info->offset + job_id;
return BuildKernel( info->nameInCode, i, info->kernel_count, info->kernels[i], info->programs + i );
return BuildKernel(info->nameInCode, i, info->kernel_count,
info->kernels[i], info->programs + i, info->relaxedMode);
}
static cl_int BuildKernel_DoubleFn( cl_uint job_id, cl_uint thread_id UNUSED, void *p );
@@ -207,7 +219,9 @@ static cl_int BuildKernel_DoubleFn( cl_uint job_id, cl_uint thread_id UNUSED, vo
{
BuildKernelInfo *info = (BuildKernelInfo*) p;
cl_uint i = info->offset + job_id;
return BuildKernelDouble( info->nameInCode, i, info->kernel_count, info->kernels[i], info->programs + i );
return BuildKernelDouble(info->nameInCode, i, info->kernel_count,
info->kernels[i], info->programs + i,
info->relaxedMode);
}
//Thread specific data for a worker thread
@@ -240,11 +254,14 @@ typedef struct TestInfo
int isFDim;
int skipNanInf;
int isNextafter;
}TestInfo;
bool relaxedMode; // True if test is running in relaxed mode, false
// otherwise.
} TestInfo;
static cl_int TestFloat( cl_uint job_id, cl_uint thread_id, void *p );
int TestFunc_Float_Float_Float_common(const Func *f, MTdata d, int isNextafter)
int TestFunc_Float_Float_Float_common(const Func *f, MTdata d, int isNextafter,
bool relaxedMode)
{
TestInfo test_info;
cl_int error;
@@ -254,7 +271,7 @@ int TestFunc_Float_Float_Float_common(const Func *f, MTdata d, int isNextafter)
double maxErrorVal2 = 0.0;
int skipTestingRelaxed = 0;
logFunctionInfo(f->name,sizeof(cl_float),gTestFastRelaxed);
logFunctionInfo(f->name, sizeof(cl_float), relaxedMode);
// Init test_info
memset( &test_info, 0, sizeof( test_info ) );
@@ -284,6 +301,7 @@ int TestFunc_Float_Float_Float_common(const Func *f, MTdata d, int isNextafter)
test_info.isFDim = 0 == strcmp( "fdim", f->nameInCode );
test_info.skipNanInf = test_info.isFDim && ! gInfNanSupport;
test_info.isNextafter = isNextafter;
test_info.relaxedMode = relaxedMode;
// cl_kernels aren't thread safe, so we make one for each vector size for every thread
for( i = gMinVectorSizeIndex; i < gMaxVectorSizeIndex; i++ )
{
@@ -342,7 +360,10 @@ int TestFunc_Float_Float_Float_common(const Func *f, MTdata d, int isNextafter)
// Init the kernels
{
BuildKernelInfo build_info = { gMinVectorSizeIndex, test_info.threadCount, test_info.k, test_info.programs, f->nameInCode };
BuildKernelInfo build_info = {
gMinVectorSizeIndex, test_info.threadCount, test_info.k,
test_info.programs, f->nameInCode, relaxedMode
};
if( (error = ThreadPool_Do( BuildKernel_FloatFn, gMaxVectorSizeIndex - gMinVectorSizeIndex, &build_info ) ))
goto exit;
}
@@ -482,6 +503,7 @@ static cl_int TestFloat( cl_uint job_id, cl_uint thread_id, void *data )
float ulps = job->ulps;
fptr func = job->f->func;
int ftz = job->ftz;
bool relaxedMode = job->relaxedMode;
MTdata d = tinfo->d;
cl_uint j, k;
cl_int error;
@@ -496,7 +518,7 @@ static cl_int TestFloat( cl_uint job_id, cl_uint thread_id, void *data )
RoundingMode oldRoundMode;
int skipVerification = 0;
if(gTestFastRelaxed)
if (relaxedMode)
{
if (strcmp(name,"pow")==0 && gFastRelaxedDerived)
{
@@ -710,7 +732,7 @@ static cl_int TestFloat( cl_uint job_id, cl_uint thread_id, void *data )
// As per OpenCL 2.0 spec, section 5.8.4.3, enabling fast-relaxed-math mode also enables
// -cl-finite-math-only optimization. This optimization allows to assume that arguments and
// results are not NaNs or +/-INFs. Hence, accept any result if inputs or results are NaNs or INFs.
if ( gTestFastRelaxed || skipNanInf)
if (relaxedMode || skipNanInf)
{
if( skipNanInf && overflow[j])
continue;
@@ -772,7 +794,7 @@ static cl_int TestFloat( cl_uint job_id, cl_uint thread_id, void *data )
// As per OpenCL 2.0 spec, section 5.8.4.3, enabling fast-relaxed-math mode also enables
// -cl-finite-math-only optimization. This optimization allows to assume that arguments and
// results are not NaNs or +/-INFs. Hence, accept any result if inputs or results are NaNs or INFs.
if( gTestFastRelaxed || skipNanInf )
if (relaxedMode || skipNanInf)
{
if( fetestexcept(FE_OVERFLOW) && skipNanInf )
continue;
@@ -817,7 +839,7 @@ static cl_int TestFloat( cl_uint job_id, cl_uint thread_id, void *data )
// As per OpenCL 2.0 spec, section 5.8.4.3, enabling fast-relaxed-math mode also enables
// -cl-finite-math-only optimization. This optimization allows to assume that arguments and
// results are not NaNs or +/-INFs. Hence, accept any result if inputs or results are NaNs or INFs.
if( gTestFastRelaxed || skipNanInf )
if (relaxedMode || skipNanInf)
{
if( fetestexcept(FE_OVERFLOW) && skipNanInf )
continue;
@@ -870,7 +892,7 @@ static cl_int TestFloat( cl_uint job_id, cl_uint thread_id, void *data )
// As per OpenCL 2.0 spec, section 5.8.4.3, enabling fast-relaxed-math mode also enables
// -cl-finite-math-only optimization. This optimization allows to assume that arguments and
// results are not NaNs or +/-INFs. Hence, accept any result if inputs or results are NaNs or INFs.
if ( gTestFastRelaxed || skipNanInf )
if (relaxedMode || skipNanInf)
{
// Note: no double rounding here. Reference functions calculate in single precision.
if( overflow[j] && skipNanInf)
@@ -977,7 +999,8 @@ static size_t specialValuesDoubleCount = sizeof( specialValuesDouble ) / sizeof(
static cl_int TestDouble( cl_uint job_id, cl_uint thread_id, void *p );
int TestFunc_Double_Double_Double_common(const Func *f, MTdata d, int isNextafter)
int TestFunc_Double_Double_Double_common(const Func *f, MTdata d,
int isNextafter, bool relaxedMode)
{
TestInfo test_info;
cl_int error;
@@ -986,7 +1009,7 @@ int TestFunc_Double_Double_Double_common(const Func *f, MTdata d, int isNextafte
double maxErrorVal = 0.0;
double maxErrorVal2 = 0.0;
logFunctionInfo(f->name,sizeof(cl_double),gTestFastRelaxed);
logFunctionInfo(f->name, sizeof(cl_double), relaxedMode);
// Init test_info
memset( &test_info, 0, sizeof( test_info ) );
@@ -1075,7 +1098,10 @@ int TestFunc_Double_Double_Double_common(const Func *f, MTdata d, int isNextafte
// Init the kernels
{
BuildKernelInfo build_info = { gMinVectorSizeIndex, test_info.threadCount, test_info.k, test_info.programs, f->nameInCode };
BuildKernelInfo build_info = {
gMinVectorSizeIndex, test_info.threadCount, test_info.k,
test_info.programs, f->nameInCode, relaxedMode
};
if( (error = ThreadPool_Do( BuildKernel_DoubleFn, gMaxVectorSizeIndex - gMinVectorSizeIndex, &build_info ) ))
goto exit;
}
@@ -1534,23 +1560,25 @@ exit:
}
int TestFunc_Float_Float_Float(const Func *f, MTdata d)
int TestFunc_Float_Float_Float(const Func *f, MTdata d, bool relaxedMode)
{
return TestFunc_Float_Float_Float_common(f, d, 0);
return TestFunc_Float_Float_Float_common(f, d, 0, relaxedMode);
}
int TestFunc_Double_Double_Double(const Func *f, MTdata d)
int TestFunc_Double_Double_Double(const Func *f, MTdata d, bool relaxedMode)
{
return TestFunc_Double_Double_Double_common(f, d, 0);
return TestFunc_Double_Double_Double_common(f, d, 0, relaxedMode);
}
int TestFunc_Float_Float_Float_nextafter(const Func *f, MTdata d)
int TestFunc_Float_Float_Float_nextafter(const Func *f, MTdata d,
bool relaxedMode)
{
return TestFunc_Float_Float_Float_common(f, d, 1);
return TestFunc_Float_Float_Float_common(f, d, 1, relaxedMode);
}
int TestFunc_Double_Double_Double_nextafter(const Func *f, MTdata d)
int TestFunc_Double_Double_Double_nextafter(const Func *f, MTdata d,
bool relaxedMode)
{
return TestFunc_Double_Double_Double_common(f, d, 1);
return TestFunc_Double_Double_Double_common(f, d, 1, relaxedMode);
}