Reduce differences by using common names (#1187)

Improve format.

The binary_operator tests are left untouched by this commit as they
require some non-automatic changes.

Signed-off-by: Marco Antognini <marco.antognini@arm.com>
This commit is contained in:
Marco Antognini
2021-03-11 09:44:38 +00:00
committed by GitHub
parent 68ee30fb4b
commit ee600e89d7
26 changed files with 247 additions and 293 deletions

View File

@@ -22,9 +22,8 @@
const double twoToMinus1022 = MAKE_HEX_DOUBLE(0x1p-1022, 1, -1022);
static int BuildKernelDouble(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, bool relaxedMode)
{
const char *c[] = { "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
"__kernel void math_kernel",
@@ -120,14 +119,12 @@ typedef struct BuildKernelInfo
bool relaxedMode; // Whether to build with -cl-fast-relaxed-math.
} BuildKernelInfo;
static cl_int BuildKernel_DoubleFn(cl_uint job_id, cl_uint thread_id UNUSED,
void *p)
static cl_int BuildKernelFn(cl_uint job_id, cl_uint thread_id UNUSED, void *p)
{
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,
info->relaxedMode);
return BuildKernel(info->nameInCode, i, info->kernel_count,
info->kernels[i], info->programs + i, info->relaxedMode);
}
// Thread specific data for a worker thread
@@ -170,7 +167,7 @@ typedef struct TestInfo
} TestInfo;
// A table of more difficult cases to get right
static const double specialValuesDouble[] = {
static const double specialValues[] = {
-NAN,
-INFINITY,
-DBL_MAX,
@@ -280,10 +277,10 @@ static const double specialValuesDouble[] = {
+0.0,
};
static size_t specialValuesDoubleCount =
sizeof(specialValuesDouble) / sizeof(specialValuesDouble[0]);
static size_t specialValuesCount =
sizeof(specialValues) / sizeof(specialValues[0]);
static cl_int TestDouble(cl_uint job_id, cl_uint thread_id, void *p);
static cl_int Test(cl_uint job_id, cl_uint thread_id, void *data);
int TestFunc_Double_Double_Double(const Func *f, MTdata d, bool relaxedMode)
{
@@ -411,7 +408,7 @@ int TestFunc_Double_Double_Double(const Func *f, MTdata d, bool relaxedMode)
gMinVectorSizeIndex, test_info.threadCount, test_info.k,
test_info.programs, f->nameInCode, relaxedMode
};
if ((error = ThreadPool_Do(BuildKernel_DoubleFn,
if ((error = ThreadPool_Do(BuildKernelFn,
gMaxVectorSizeIndex - gMinVectorSizeIndex,
&build_info)))
goto exit;
@@ -420,7 +417,7 @@ int TestFunc_Double_Double_Double(const Func *f, MTdata d, bool relaxedMode)
// Run the kernels
if (!gSkipCorrectnessTesting)
{
error = ThreadPool_Do(TestDouble, test_info.jobCount, &test_info);
error = ThreadPool_Do(Test, test_info.jobCount, &test_info);
// Accumulate the arithmetic errors
for (i = 0; i < test_info.threadCount; i++)
@@ -562,7 +559,7 @@ exit:
return error;
}
static cl_int TestDouble(cl_uint job_id, cl_uint thread_id, void *data)
static cl_int Test(cl_uint job_id, cl_uint thread_id, void *data)
{
const TestInfo *job = (const TestInfo *)data;
size_t buffer_elements = job->subBufferSize;
@@ -608,8 +605,7 @@ static cl_int TestDouble(cl_uint job_id, cl_uint thread_id, void *data)
cl_ulong *p = (cl_ulong *)gIn + thread_id * buffer_elements;
cl_ulong *p2 = (cl_ulong *)gIn2 + thread_id * buffer_elements;
j = 0;
int totalSpecialValueCount =
specialValuesDoubleCount * specialValuesDoubleCount;
int totalSpecialValueCount = specialValuesCount * specialValuesCount;
int indx = (totalSpecialValueCount - 1) / buffer_elements;
if (job_id <= (cl_uint)indx)
@@ -618,18 +614,18 @@ static cl_int TestDouble(cl_uint job_id, cl_uint thread_id, void *data)
cl_double *fp2 = (cl_double *)p2;
uint32_t x, y;
x = (job_id * buffer_elements) % specialValuesDoubleCount;
y = (job_id * buffer_elements) / specialValuesDoubleCount;
x = (job_id * buffer_elements) % specialValuesCount;
y = (job_id * buffer_elements) / specialValuesCount;
for (; j < buffer_elements; j++)
{
fp[j] = specialValuesDouble[x];
fp2[j] = specialValuesDouble[y];
if (++x >= specialValuesDoubleCount)
fp[j] = specialValues[x];
fp2[j] = specialValues[y];
if (++x >= specialValuesCount)
{
x = 0;
y++;
if (y >= specialValuesDoubleCount) break;
if (y >= specialValuesCount) break;
}
}
}