Move code around to reduce differences (#1185)
Code is moved to reduce the differences between tests for single- and
double-precision.
Improve consistency in double-literal.
Signed-off-by: Marco Antognini <marco.antognini@arm.com>
diff --git a/test_conformance/math_brute_force/binary_double.cpp b/test_conformance/math_brute_force/binary_double.cpp
index fad03ad..cbb186e 100644
--- a/test_conformance/math_brute_force/binary_double.cpp
+++ b/test_conformance/math_brute_force/binary_double.cpp
@@ -186,8 +186,8 @@
MAKE_HEX_DOUBLE(-0x1.0000000000001p31, -0x10000000000001LL, -21),
MAKE_HEX_DOUBLE(-0x1.0p31, -0x1LL, 31),
MAKE_HEX_DOUBLE(-0x1.fffffffffffffp30, -0x1fffffffffffffLL, -22),
- -1000.,
- -100.,
+ -1000.0,
+ -100.0,
-4.0,
-3.5,
-3.0,
@@ -240,8 +240,8 @@
MAKE_HEX_DOUBLE(+0x1.0000000000001p31, +0x10000000000001LL, -21),
MAKE_HEX_DOUBLE(+0x1.0p31, +0x1LL, 31),
MAKE_HEX_DOUBLE(+0x1.fffffffffffffp30, +0x1fffffffffffffLL, -22),
- +1000.,
- +100.,
+ +1000.0,
+ +100.0,
+4.0,
+3.5,
+3.0,
diff --git a/test_conformance/math_brute_force/binary_float.cpp b/test_conformance/math_brute_force/binary_float.cpp
index a31bfb2..8dfb9f4 100644
--- a/test_conformance/math_brute_force/binary_float.cpp
+++ b/test_conformance/math_brute_force/binary_float.cpp
@@ -126,6 +126,45 @@
info->kernels[i], info->programs + i, info->relaxedMode);
}
+// Thread specific data for a worker thread
+typedef struct ThreadInfo
+{
+ cl_mem inBuf; // input buffer for the thread
+ cl_mem inBuf2; // input buffer for the thread
+ cl_mem outBuf[VECTOR_SIZE_COUNT]; // output buffers for the thread
+ float maxError; // max error value. Init to 0.
+ double
+ maxErrorValue; // position of the max error value (param 1). Init to 0.
+ double maxErrorValue2; // position of the max error value (param 2). Init
+ // to 0.
+ MTdata d;
+ cl_command_queue tQueue; // per thread command queue to improve performance
+} ThreadInfo;
+
+typedef struct TestInfo
+{
+ size_t subBufferSize; // Size of the sub-buffer in elements
+ const Func *f; // A pointer to the function info
+ cl_program programs[VECTOR_SIZE_COUNT]; // programs for various vector sizes
+ cl_kernel
+ *k[VECTOR_SIZE_COUNT]; // arrays of thread-specific kernels for each
+ // worker thread: k[vector_size][thread_id]
+ ThreadInfo *
+ tinfo; // An array of thread specific information for each worker thread
+ cl_uint threadCount; // Number of worker threads
+ cl_uint jobCount; // Number of jobs
+ cl_uint step; // step between each chunk and the next.
+ cl_uint scale; // stride between individual test values
+ float ulps; // max_allowed ulps
+ int ftz; // non-zero if running in flush to zero mode
+
+ int isFDim;
+ int skipNanInf;
+ int isNextafter;
+ bool relaxedMode; // True if test is running in relaxed mode, false
+ // otherwise.
+} TestInfo;
+
// A table of more difficult cases to get right
static const float specialValuesFloat[] = {
-NAN,
@@ -226,50 +265,12 @@
MAKE_HEX_FLOAT(+0x0.000006p-126f, +0x0000006L, -150),
MAKE_HEX_FLOAT(+0x0.000004p-126f, +0x0000004L, -150),
MAKE_HEX_FLOAT(+0x0.000002p-126f, +0x0000002L, -150),
- +0.0f
+ +0.0f,
};
static const size_t specialValuesFloatCount =
sizeof(specialValuesFloat) / sizeof(specialValuesFloat[0]);
-// Thread specific data for a worker thread
-typedef struct ThreadInfo
-{
- cl_mem inBuf; // input buffer for the thread
- cl_mem inBuf2; // input buffer for the thread
- cl_mem outBuf[VECTOR_SIZE_COUNT]; // output buffers for the thread
- float maxError; // max error value. Init to 0.
- double
- maxErrorValue; // position of the max error value (param 1). Init to 0.
- double maxErrorValue2; // position of the max error value (param 2). Init
- // to 0.
- MTdata d;
- cl_command_queue tQueue; // per thread command queue to improve performance
-} ThreadInfo;
-
-typedef struct TestInfo
-{
- size_t subBufferSize; // Size of the sub-buffer in elements
- const Func *f; // A pointer to the function info
- cl_program programs[VECTOR_SIZE_COUNT]; // programs for various vector sizes
- cl_kernel
- *k[VECTOR_SIZE_COUNT]; // arrays of thread-specific kernels for each
- // worker thread: k[vector_size][thread_id]
- ThreadInfo *
- tinfo; // An array of thread specific information for each worker thread
- cl_uint threadCount; // Number of worker threads
- cl_uint jobCount; // Number of jobs
- cl_uint step; // step between each chunk and the next.
- cl_uint scale; // stride between individual test values
- float ulps; // max_allowed ulps
- int ftz; // non-zero if running in flush to zero mode
-
- int isFDim;
- int skipNanInf;
- int isNextafter;
- 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);
diff --git a/test_conformance/math_brute_force/binary_i_double.cpp b/test_conformance/math_brute_force/binary_i_double.cpp
index 6839dfb..eb2c59e 100644
--- a/test_conformance/math_brute_force/binary_i_double.cpp
+++ b/test_conformance/math_brute_force/binary_i_double.cpp
@@ -181,8 +181,8 @@
MAKE_HEX_DOUBLE(-0x1.0000000000001p31, -0x10000000000001LL, -21),
MAKE_HEX_DOUBLE(-0x1.0p31, -0x1LL, 31),
MAKE_HEX_DOUBLE(-0x1.fffffffffffffp30, -0x1fffffffffffffLL, -22),
- -1000.,
- -100.,
+ -1000.0,
+ -100.0,
-4.0,
-3.5,
-3.0,
@@ -235,8 +235,8 @@
MAKE_HEX_DOUBLE(+0x1.0000000000001p31, +0x10000000000001LL, -21),
MAKE_HEX_DOUBLE(+0x1.0p31, +0x1LL, 31),
MAKE_HEX_DOUBLE(+0x1.fffffffffffffp30, +0x1fffffffffffffLL, -22),
- +1000.,
- +100.,
+ +1000.0,
+ +100.0,
+4.0,
+3.5,
+3.0,
diff --git a/test_conformance/math_brute_force/binary_i_float.cpp b/test_conformance/math_brute_force/binary_i_float.cpp
index ceb79dd..019c96a 100644
--- a/test_conformance/math_brute_force/binary_i_float.cpp
+++ b/test_conformance/math_brute_force/binary_i_float.cpp
@@ -125,6 +125,41 @@
info->kernels[i], info->programs + i, info->relaxedMode);
}
+// Thread specific data for a worker thread
+typedef struct ThreadInfo
+{
+ cl_mem inBuf; // input buffer for the thread
+ cl_mem inBuf2; // input buffer for the thread
+ cl_mem outBuf[VECTOR_SIZE_COUNT]; // output buffers for the thread
+ float maxError; // max error value. Init to 0.
+ double
+ maxErrorValue; // position of the max error value (param 1). Init to 0.
+ cl_int maxErrorValue2; // position of the max error value (param 2). Init
+ // to 0.
+ MTdata d;
+ cl_command_queue tQueue; // per thread command queue to improve performance
+} ThreadInfo;
+
+typedef struct TestInfo
+{
+ size_t subBufferSize; // Size of the sub-buffer in elements
+ const Func *f; // A pointer to the function info
+ cl_program programs[VECTOR_SIZE_COUNT]; // programs for various vector sizes
+ cl_kernel
+ *k[VECTOR_SIZE_COUNT]; // arrays of thread-specific kernels for each
+ // worker thread: k[vector_size][thread_id]
+ ThreadInfo *
+ tinfo; // An array of thread specific information for each worker thread
+ cl_uint threadCount; // Number of worker threads
+ cl_uint jobCount; // Number of jobs
+ cl_uint step; // step between each chunk and the next.
+ cl_uint scale; // stride between individual test values
+ float ulps; // max_allowed ulps
+ int ftz; // non-zero if running in flush to zero mode
+
+ // no special values
+} TestInfo;
+
// A table of more difficult cases to get right
static const float specialValuesFloat[] = {
-NAN,
@@ -225,7 +260,7 @@
MAKE_HEX_FLOAT(+0x0.000006p-126f, +0x0000006L, -150),
MAKE_HEX_FLOAT(+0x0.000004p-126f, +0x0000004L, -150),
MAKE_HEX_FLOAT(+0x0.000002p-126f, +0x0000002L, -150),
- +0.0f
+ +0.0f,
};
static const size_t specialValuesFloatCount =
@@ -240,41 +275,6 @@
static size_t specialValuesIntCount =
sizeof(specialValuesInt) / sizeof(specialValuesInt[0]);
-// Thread specific data for a worker thread
-typedef struct ThreadInfo
-{
- cl_mem inBuf; // input buffer for the thread
- cl_mem inBuf2; // input buffer for the thread
- cl_mem outBuf[VECTOR_SIZE_COUNT]; // output buffers for the thread
- float maxError; // max error value. Init to 0.
- double
- maxErrorValue; // position of the max error value (param 1). Init to 0.
- cl_int maxErrorValue2; // position of the max error value (param 2). Init
- // to 0.
- MTdata d;
- cl_command_queue tQueue; // per thread command queue to improve performance
-} ThreadInfo;
-
-typedef struct TestInfo
-{
- size_t subBufferSize; // Size of the sub-buffer in elements
- const Func *f; // A pointer to the function info
- cl_program programs[VECTOR_SIZE_COUNT]; // programs for various vector sizes
- cl_kernel
- *k[VECTOR_SIZE_COUNT]; // arrays of thread-specific kernels for each
- // worker thread: k[vector_size][thread_id]
- ThreadInfo *
- tinfo; // An array of thread specific information for each worker thread
- cl_uint threadCount; // Number of worker threads
- cl_uint jobCount; // Number of jobs
- cl_uint step; // step between each chunk and the next.
- cl_uint scale; // stride between individual test values
- float ulps; // max_allowed ulps
- int ftz; // non-zero if running in flush to zero mode
-
- // no special values
-} TestInfo;
-
static cl_int TestFloat(cl_uint job_id, cl_uint thread_id, void *p);
int TestFunc_Float_Float_Int(const Func *f, MTdata d, bool relaxedMode)
diff --git a/test_conformance/math_brute_force/binary_operator_double.cpp b/test_conformance/math_brute_force/binary_operator_double.cpp
index 939ea6d..a7fb3cd 100644
--- a/test_conformance/math_brute_force/binary_operator_double.cpp
+++ b/test_conformance/math_brute_force/binary_operator_double.cpp
@@ -241,8 +241,8 @@
MAKE_HEX_DOUBLE(+0x1.0000000000001p31, +0x10000000000001LL, -21),
MAKE_HEX_DOUBLE(+0x1.0p31, +0x1LL, 31),
MAKE_HEX_DOUBLE(+0x1.fffffffffffffp30, +0x1fffffffffffffLL, -22),
- +1000.,
- +100.,
+ +1000.0,
+ +100.0,
+4.0,
+3.5,
+3.0,
diff --git a/test_conformance/math_brute_force/binary_operator_float.cpp b/test_conformance/math_brute_force/binary_operator_float.cpp
index efef4fe..a9d3b7c 100644
--- a/test_conformance/math_brute_force/binary_operator_float.cpp
+++ b/test_conformance/math_brute_force/binary_operator_float.cpp
@@ -130,6 +130,43 @@
info->kernels[i], info->programs + i, info->relaxedMode);
}
+// Thread specific data for a worker thread
+typedef struct ThreadInfo
+{
+ cl_mem inBuf; // input buffer for the thread
+ cl_mem inBuf2; // input buffer for the thread
+ cl_mem outBuf[VECTOR_SIZE_COUNT]; // output buffers for the thread
+ float maxError; // max error value. Init to 0.
+ double
+ maxErrorValue; // position of the max error value (param 1). Init to 0.
+ double maxErrorValue2; // position of the max error value (param 2). Init
+ // to 0.
+ MTdata d;
+ cl_command_queue tQueue; // per thread command queue to improve performance
+} ThreadInfo;
+
+typedef struct TestInfo
+{
+ size_t subBufferSize; // Size of the sub-buffer in elements
+ const Func *f; // A pointer to the function info
+ cl_program programs[VECTOR_SIZE_COUNT]; // programs for various vector sizes
+ cl_kernel
+ *k[VECTOR_SIZE_COUNT]; // arrays of thread-specific kernels for each
+ // worker thread: k[vector_size][thread_id]
+ ThreadInfo *
+ tinfo; // An array of thread specific information for each worker thread
+ cl_uint threadCount; // Number of worker threads
+ cl_uint jobCount; // Number of jobs
+ cl_uint step; // step between each chunk and the next.
+ cl_uint scale; // stride between individual test values
+ float ulps; // max_allowed ulps
+ int ftz; // non-zero if running in flush to zero mode
+ bool relaxedMode; // True if the test is being run in relaxed mode, false
+ // otherwise.
+
+ // no special fields
+} TestInfo;
+
// A table of more difficult cases to get right
static const float specialValuesFloat[] = {
-NAN,
@@ -230,49 +267,12 @@
MAKE_HEX_FLOAT(+0x0.000006p-126f, +0x0000006L, -150),
MAKE_HEX_FLOAT(+0x0.000004p-126f, +0x0000004L, -150),
MAKE_HEX_FLOAT(+0x0.000002p-126f, +0x0000002L, -150),
- +0.0f
+ +0.0f,
};
static const size_t specialValuesFloatCount =
sizeof(specialValuesFloat) / sizeof(specialValuesFloat[0]);
-// Thread specific data for a worker thread
-typedef struct ThreadInfo
-{
- cl_mem inBuf; // input buffer for the thread
- cl_mem inBuf2; // input buffer for the thread
- cl_mem outBuf[VECTOR_SIZE_COUNT]; // output buffers for the thread
- float maxError; // max error value. Init to 0.
- double
- maxErrorValue; // position of the max error value (param 1). Init to 0.
- double maxErrorValue2; // position of the max error value (param 2). Init
- // to 0.
- MTdata d;
- cl_command_queue tQueue; // per thread command queue to improve performance
-} ThreadInfo;
-
-typedef struct TestInfo
-{
- size_t subBufferSize; // Size of the sub-buffer in elements
- const Func *f; // A pointer to the function info
- cl_program programs[VECTOR_SIZE_COUNT]; // programs for various vector sizes
- cl_kernel
- *k[VECTOR_SIZE_COUNT]; // arrays of thread-specific kernels for each
- // worker thread: k[vector_size][thread_id]
- ThreadInfo *
- tinfo; // An array of thread specific information for each worker thread
- cl_uint threadCount; // Number of worker threads
- cl_uint jobCount; // Number of jobs
- cl_uint step; // step between each chunk and the next.
- cl_uint scale; // stride between individual test values
- float ulps; // max_allowed ulps
- int ftz; // non-zero if running in flush to zero mode
- bool relaxedMode; // True if the test is being run in relaxed mode, false
- // otherwise.
-
- // no special fields
-} TestInfo;
-
static cl_int TestFloat(cl_uint job_id, cl_uint thread_id, void *p);
int TestFunc_Float_Float_Float_Operator(const Func *f, MTdata d,
diff --git a/test_conformance/math_brute_force/macro_binary_double.cpp b/test_conformance/math_brute_force/macro_binary_double.cpp
index 81bc4d0..2ea785e 100644
--- a/test_conformance/math_brute_force/macro_binary_double.cpp
+++ b/test_conformance/math_brute_force/macro_binary_double.cpp
@@ -173,8 +173,8 @@
MAKE_HEX_DOUBLE(-0x1.0000000000001p31, -0x10000000000001LL, -21),
MAKE_HEX_DOUBLE(-0x1.0p31, -0x1LL, 31),
MAKE_HEX_DOUBLE(-0x1.fffffffffffffp30, -0x1fffffffffffffLL, -22),
- -1000.,
- -100.,
+ -1000.0,
+ -100.0,
-4.0,
-3.5,
-3.0,
@@ -227,8 +227,8 @@
MAKE_HEX_DOUBLE(+0x1.0000000000001p31, +0x10000000000001LL, -21),
MAKE_HEX_DOUBLE(+0x1.0p31, +0x1LL, 31),
MAKE_HEX_DOUBLE(+0x1.fffffffffffffp30, +0x1fffffffffffffLL, -22),
- +1000.,
- +100.,
+ +1000.0,
+ +100.0,
+4.0,
+3.5,
+3.0,
diff --git a/test_conformance/math_brute_force/macro_binary_float.cpp b/test_conformance/math_brute_force/macro_binary_float.cpp
index 1b5dc33..a61ab6b 100644
--- a/test_conformance/math_brute_force/macro_binary_float.cpp
+++ b/test_conformance/math_brute_force/macro_binary_float.cpp
@@ -124,6 +124,34 @@
info->kernels[i], info->programs + i, info->relaxedMode);
}
+// Thread specific data for a worker thread
+typedef struct ThreadInfo
+{
+ cl_mem inBuf; // input buffer for the thread
+ cl_mem inBuf2; // input buffer for the thread
+ cl_mem outBuf[VECTOR_SIZE_COUNT]; // output buffers for the thread
+ MTdata d;
+ cl_command_queue tQueue; // per thread command queue to improve performance
+} ThreadInfo;
+
+typedef struct TestInfo
+{
+ size_t subBufferSize; // Size of the sub-buffer in elements
+ const Func *f; // A pointer to the function info
+ cl_program programs[VECTOR_SIZE_COUNT]; // programs for various vector sizes
+ cl_kernel
+ *k[VECTOR_SIZE_COUNT]; // arrays of thread-specific kernels for each
+ // worker thread: k[vector_size][thread_id]
+ ThreadInfo *
+ tinfo; // An array of thread specific information for each worker thread
+ cl_uint threadCount; // Number of worker threads
+ cl_uint jobCount; // Number of jobs
+ cl_uint step; // step between each chunk and the next.
+ cl_uint scale; // stride between individual test values
+ int ftz; // non-zero if running in flush to zero mode
+
+} TestInfo;
+
// A table of more difficult cases to get right
static const float specialValuesFloat[] = {
-NAN,
@@ -224,40 +252,12 @@
MAKE_HEX_FLOAT(+0x0.000006p-126f, +0x0000006L, -150),
MAKE_HEX_FLOAT(+0x0.000004p-126f, +0x0000004L, -150),
MAKE_HEX_FLOAT(+0x0.000002p-126f, +0x0000002L, -150),
- +0.0f
+ +0.0f,
};
static const size_t specialValuesFloatCount =
sizeof(specialValuesFloat) / sizeof(specialValuesFloat[0]);
-// Thread specific data for a worker thread
-typedef struct ThreadInfo
-{
- cl_mem inBuf; // input buffer for the thread
- cl_mem inBuf2; // input buffer for the thread
- cl_mem outBuf[VECTOR_SIZE_COUNT]; // output buffers for the thread
- MTdata d;
- cl_command_queue tQueue; // per thread command queue to improve performance
-} ThreadInfo;
-
-typedef struct TestInfo
-{
- size_t subBufferSize; // Size of the sub-buffer in elements
- const Func *f; // A pointer to the function info
- cl_program programs[VECTOR_SIZE_COUNT]; // programs for various vector sizes
- cl_kernel
- *k[VECTOR_SIZE_COUNT]; // arrays of thread-specific kernels for each
- // worker thread: k[vector_size][thread_id]
- ThreadInfo *
- tinfo; // An array of thread specific information for each worker thread
- cl_uint threadCount; // Number of worker threads
- cl_uint jobCount; // Number of jobs
- cl_uint step; // step between each chunk and the next.
- cl_uint scale; // stride between individual test values
- int ftz; // non-zero if running in flush to zero mode
-
-} TestInfo;
-
static cl_int TestFloat(cl_uint job_id, cl_uint thread_id, void *p);
int TestMacro_Int_Float_Float(const Func *f, MTdata d, bool relaxedMode)
diff --git a/test_conformance/math_brute_force/ternary_float.cpp b/test_conformance/math_brute_force/ternary_float.cpp
index 3b2adf8..1fcdc4a 100644
--- a/test_conformance/math_brute_force/ternary_float.cpp
+++ b/test_conformance/math_brute_force/ternary_float.cpp
@@ -208,7 +208,7 @@
MAKE_HEX_FLOAT(+0x0.000006p-126f, +0x0000006L, -150),
MAKE_HEX_FLOAT(+0x0.000004p-126f, +0x0000004L, -150),
MAKE_HEX_FLOAT(+0x0.000002p-126f, +0x0000002L, -150),
- +0.0f
+ +0.0f,
};
static const size_t specialValuesFloatCount =
diff --git a/test_conformance/math_brute_force/unary_double.cpp b/test_conformance/math_brute_force/unary_double.cpp
index 1ff3d9c..99959ae 100644
--- a/test_conformance/math_brute_force/unary_double.cpp
+++ b/test_conformance/math_brute_force/unary_double.cpp
@@ -160,244 +160,7 @@
// otherwise.
} TestInfo;
-static cl_int TestDouble(cl_uint job_id, cl_uint thread_id, void *data)
-{
- const TestInfo *job = (const TestInfo *)data;
- size_t buffer_elements = job->subBufferSize;
- size_t buffer_size = buffer_elements * sizeof(cl_double);
- cl_uint scale = job->scale;
- cl_uint base = job_id * (cl_uint)job->step;
- ThreadInfo *tinfo = job->tinfo + thread_id;
- float ulps = job->ulps;
- dptr func = job->f->dfunc;
- cl_uint j, k;
- cl_int error;
- int ftz = job->ftz;
-
- Force64BitFPUPrecision();
-
- // start the map of the output arrays
- cl_event e[VECTOR_SIZE_COUNT];
- cl_ulong *out[VECTOR_SIZE_COUNT];
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- out[j] = (cl_ulong *)clEnqueueMapBuffer(
- tinfo->tQueue, tinfo->outBuf[j], CL_FALSE, CL_MAP_WRITE, 0,
- buffer_size, 0, NULL, e + j, &error);
- if (error || NULL == out[j])
- {
- vlog_error("Error: clEnqueueMapBuffer %d failed! err: %d\n", j,
- error);
- return error;
- }
- }
-
- // Get that moving
- if ((error = clFlush(tinfo->tQueue))) vlog("clFlush failed\n");
-
- // Write the new values to the input array
- cl_double *p = (cl_double *)gIn + thread_id * buffer_elements;
- for (j = 0; j < buffer_elements; j++)
- p[j] = DoubleFromUInt32(base + j * scale);
-
- if ((error = clEnqueueWriteBuffer(tinfo->tQueue, tinfo->inBuf, CL_FALSE, 0,
- buffer_size, p, 0, NULL, NULL)))
- {
- vlog_error("Error: clEnqueueWriteBuffer failed! err: %d\n", error);
- return error;
- }
-
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- // Wait for the map to finish
- if ((error = clWaitForEvents(1, e + j)))
- {
- vlog_error("Error: clWaitForEvents failed! err: %d\n", error);
- return error;
- }
- if ((error = clReleaseEvent(e[j])))
- {
- vlog_error("Error: clReleaseEvent failed! err: %d\n", error);
- return error;
- }
-
- // Fill the result buffer with garbage, so that old results don't carry
- // over
- uint32_t pattern = 0xffffdead;
- memset_pattern4(out[j], &pattern, buffer_size);
- if ((error = clEnqueueUnmapMemObject(tinfo->tQueue, tinfo->outBuf[j],
- out[j], 0, NULL, NULL)))
- {
- vlog_error("Error: clEnqueueMapBuffer failed! err: %d\n", error);
- return error;
- }
-
- // run the kernel
- size_t vectorCount =
- (buffer_elements + sizeValues[j] - 1) / sizeValues[j];
- cl_kernel kernel = job->k[j][thread_id]; // each worker thread has its
- // own copy of the cl_kernel
- cl_program program = job->programs[j];
-
- if ((error = clSetKernelArg(kernel, 0, sizeof(tinfo->outBuf[j]),
- &tinfo->outBuf[j])))
- {
- LogBuildError(program);
- return error;
- }
- if ((error = clSetKernelArg(kernel, 1, sizeof(tinfo->inBuf),
- &tinfo->inBuf)))
- {
- LogBuildError(program);
- return error;
- }
-
- if ((error = clEnqueueNDRangeKernel(tinfo->tQueue, kernel, 1, NULL,
- &vectorCount, NULL, 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- return error;
- }
- }
-
-
- // Get that moving
- if ((error = clFlush(tinfo->tQueue))) vlog("clFlush 2 failed\n");
-
- if (gSkipCorrectnessTesting) return CL_SUCCESS;
-
- // Calculate the correctly rounded reference result
- cl_double *r = (cl_double *)gOut_Ref + thread_id * buffer_elements;
- cl_double *s = (cl_double *)p;
- for (j = 0; j < buffer_elements; j++) r[j] = (cl_double)func.f_f(s[j]);
-
- // Read the data back -- no need to wait for the first N-1 buffers. This is
- // an in order queue.
- for (j = gMinVectorSizeIndex; j + 1 < gMaxVectorSizeIndex; j++)
- {
- out[j] = (cl_ulong *)clEnqueueMapBuffer(
- tinfo->tQueue, tinfo->outBuf[j], CL_FALSE, CL_MAP_READ, 0,
- buffer_size, 0, NULL, NULL, &error);
- if (error || NULL == out[j])
- {
- vlog_error("Error: clEnqueueMapBuffer %d failed! err: %d\n", j,
- error);
- return error;
- }
- }
- // Wait for the last buffer
- out[j] = (cl_ulong *)clEnqueueMapBuffer(tinfo->tQueue, tinfo->outBuf[j],
- CL_TRUE, CL_MAP_READ, 0,
- buffer_size, 0, NULL, NULL, &error);
- if (error || NULL == out[j])
- {
- vlog_error("Error: clEnqueueMapBuffer %d failed! err: %d\n", j, error);
- return error;
- }
-
-
- // Verify data
- cl_ulong *t = (cl_ulong *)r;
- for (j = 0; j < buffer_elements; j++)
- {
- for (k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
- {
- cl_ulong *q = out[k];
-
- // If we aren't getting the correctly rounded result
- if (t[j] != q[j])
- {
- cl_double test = ((cl_double *)q)[j];
- long double correct = func.f_f(s[j]);
- float err = Bruteforce_Ulp_Error_Double(test, correct);
- int fail = !(fabsf(err) <= ulps);
-
- if (fail)
- {
- if (ftz)
- {
- // retry per section 6.5.3.2
- if (IsDoubleResultSubnormal(correct, ulps))
- {
- fail = fail && (test != 0.0f);
- if (!fail) err = 0.0f;
- }
-
- // retry per section 6.5.3.3
- if (IsDoubleSubnormal(s[j]))
- {
- long double correct2 = func.f_f(0.0L);
- long double correct3 = func.f_f(-0.0L);
- float err2 =
- Bruteforce_Ulp_Error_Double(test, correct2);
- float err3 =
- Bruteforce_Ulp_Error_Double(test, correct3);
- fail = fail
- && ((!(fabsf(err2) <= ulps))
- && (!(fabsf(err3) <= ulps)));
- if (fabsf(err2) < fabsf(err)) err = err2;
- if (fabsf(err3) < fabsf(err)) err = err3;
-
- // retry per section 6.5.3.4
- if (IsDoubleResultSubnormal(correct2, ulps)
- || IsDoubleResultSubnormal(correct3, ulps))
- {
- fail = fail && (test != 0.0f);
- if (!fail) err = 0.0f;
- }
- }
- }
- }
- if (fabsf(err) > tinfo->maxError)
- {
- tinfo->maxError = fabsf(err);
- tinfo->maxErrorValue = s[j];
- }
- if (fail)
- {
- vlog_error("\nERROR: %s%s: %f ulp error at %.13la "
- "(0x%16.16llx): *%.13la vs. %.13la\n",
- job->f->name, sizeNames[k], err,
- ((cl_double *)gIn)[j], ((cl_ulong *)gIn)[j],
- ((cl_double *)gOut_Ref)[j], test);
- return -1;
- }
- }
- }
- }
-
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- if ((error = clEnqueueUnmapMemObject(tinfo->tQueue, tinfo->outBuf[j],
- out[j], 0, NULL, NULL)))
- {
- vlog_error("Error: clEnqueueUnmapMemObject %d failed 2! err: %d\n",
- j, error);
- return error;
- }
- }
-
- if ((error = clFlush(tinfo->tQueue))) vlog("clFlush 3 failed\n");
-
-
- if (0 == (base & 0x0fffffff))
- {
- if (gVerboseBruteForce)
- {
- vlog("base:%14u step:%10u scale:%10zd buf_elements:%10u ulps:%5.3f "
- "ThreadCount:%2u\n",
- base, job->step, buffer_elements, job->scale, job->ulps,
- job->threadCount);
- }
- else
- {
- vlog(".");
- }
- fflush(stdout);
- }
-
- return CL_SUCCESS;
-}
+static cl_int TestDouble(cl_uint job_id, cl_uint thread_id, void *data);
int TestFunc_Double_Double(const Func *f, MTdata d, bool relaxedMode)
{
@@ -660,3 +423,242 @@
return error;
}
+
+static cl_int TestDouble(cl_uint job_id, cl_uint thread_id, void *data)
+{
+ const TestInfo *job = (const TestInfo *)data;
+ size_t buffer_elements = job->subBufferSize;
+ size_t buffer_size = buffer_elements * sizeof(cl_double);
+ cl_uint scale = job->scale;
+ cl_uint base = job_id * (cl_uint)job->step;
+ ThreadInfo *tinfo = job->tinfo + thread_id;
+ float ulps = job->ulps;
+ dptr func = job->f->dfunc;
+ cl_uint j, k;
+ cl_int error;
+ int ftz = job->ftz;
+
+ Force64BitFPUPrecision();
+
+ // start the map of the output arrays
+ cl_event e[VECTOR_SIZE_COUNT];
+ cl_ulong *out[VECTOR_SIZE_COUNT];
+ for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
+ {
+ out[j] = (cl_ulong *)clEnqueueMapBuffer(
+ tinfo->tQueue, tinfo->outBuf[j], CL_FALSE, CL_MAP_WRITE, 0,
+ buffer_size, 0, NULL, e + j, &error);
+ if (error || NULL == out[j])
+ {
+ vlog_error("Error: clEnqueueMapBuffer %d failed! err: %d\n", j,
+ error);
+ return error;
+ }
+ }
+
+ // Get that moving
+ if ((error = clFlush(tinfo->tQueue))) vlog("clFlush failed\n");
+
+ // Write the new values to the input array
+ cl_double *p = (cl_double *)gIn + thread_id * buffer_elements;
+ for (j = 0; j < buffer_elements; j++)
+ p[j] = DoubleFromUInt32(base + j * scale);
+
+ if ((error = clEnqueueWriteBuffer(tinfo->tQueue, tinfo->inBuf, CL_FALSE, 0,
+ buffer_size, p, 0, NULL, NULL)))
+ {
+ vlog_error("Error: clEnqueueWriteBuffer failed! err: %d\n", error);
+ return error;
+ }
+
+ for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
+ {
+ // Wait for the map to finish
+ if ((error = clWaitForEvents(1, e + j)))
+ {
+ vlog_error("Error: clWaitForEvents failed! err: %d\n", error);
+ return error;
+ }
+ if ((error = clReleaseEvent(e[j])))
+ {
+ vlog_error("Error: clReleaseEvent failed! err: %d\n", error);
+ return error;
+ }
+
+ // Fill the result buffer with garbage, so that old results don't carry
+ // over
+ uint32_t pattern = 0xffffdead;
+ memset_pattern4(out[j], &pattern, buffer_size);
+ if ((error = clEnqueueUnmapMemObject(tinfo->tQueue, tinfo->outBuf[j],
+ out[j], 0, NULL, NULL)))
+ {
+ vlog_error("Error: clEnqueueMapBuffer failed! err: %d\n", error);
+ return error;
+ }
+
+ // run the kernel
+ size_t vectorCount =
+ (buffer_elements + sizeValues[j] - 1) / sizeValues[j];
+ cl_kernel kernel = job->k[j][thread_id]; // each worker thread has its
+ // own copy of the cl_kernel
+ cl_program program = job->programs[j];
+
+ if ((error = clSetKernelArg(kernel, 0, sizeof(tinfo->outBuf[j]),
+ &tinfo->outBuf[j])))
+ {
+ LogBuildError(program);
+ return error;
+ }
+ if ((error = clSetKernelArg(kernel, 1, sizeof(tinfo->inBuf),
+ &tinfo->inBuf)))
+ {
+ LogBuildError(program);
+ return error;
+ }
+
+ if ((error = clEnqueueNDRangeKernel(tinfo->tQueue, kernel, 1, NULL,
+ &vectorCount, NULL, 0, NULL, NULL)))
+ {
+ vlog_error("FAILED -- could not execute kernel\n");
+ return error;
+ }
+ }
+
+
+ // Get that moving
+ if ((error = clFlush(tinfo->tQueue))) vlog("clFlush 2 failed\n");
+
+ if (gSkipCorrectnessTesting) return CL_SUCCESS;
+
+ // Calculate the correctly rounded reference result
+ cl_double *r = (cl_double *)gOut_Ref + thread_id * buffer_elements;
+ cl_double *s = (cl_double *)p;
+ for (j = 0; j < buffer_elements; j++) r[j] = (cl_double)func.f_f(s[j]);
+
+ // Read the data back -- no need to wait for the first N-1 buffers. This is
+ // an in order queue.
+ for (j = gMinVectorSizeIndex; j + 1 < gMaxVectorSizeIndex; j++)
+ {
+ out[j] = (cl_ulong *)clEnqueueMapBuffer(
+ tinfo->tQueue, tinfo->outBuf[j], CL_FALSE, CL_MAP_READ, 0,
+ buffer_size, 0, NULL, NULL, &error);
+ if (error || NULL == out[j])
+ {
+ vlog_error("Error: clEnqueueMapBuffer %d failed! err: %d\n", j,
+ error);
+ return error;
+ }
+ }
+ // Wait for the last buffer
+ out[j] = (cl_ulong *)clEnqueueMapBuffer(tinfo->tQueue, tinfo->outBuf[j],
+ CL_TRUE, CL_MAP_READ, 0,
+ buffer_size, 0, NULL, NULL, &error);
+ if (error || NULL == out[j])
+ {
+ vlog_error("Error: clEnqueueMapBuffer %d failed! err: %d\n", j, error);
+ return error;
+ }
+
+
+ // Verify data
+ cl_ulong *t = (cl_ulong *)r;
+ for (j = 0; j < buffer_elements; j++)
+ {
+ for (k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
+ {
+ cl_ulong *q = out[k];
+
+ // If we aren't getting the correctly rounded result
+ if (t[j] != q[j])
+ {
+ cl_double test = ((cl_double *)q)[j];
+ long double correct = func.f_f(s[j]);
+ float err = Bruteforce_Ulp_Error_Double(test, correct);
+ int fail = !(fabsf(err) <= ulps);
+
+ if (fail)
+ {
+ if (ftz)
+ {
+ // retry per section 6.5.3.2
+ if (IsDoubleResultSubnormal(correct, ulps))
+ {
+ fail = fail && (test != 0.0f);
+ if (!fail) err = 0.0f;
+ }
+
+ // retry per section 6.5.3.3
+ if (IsDoubleSubnormal(s[j]))
+ {
+ long double correct2 = func.f_f(0.0L);
+ long double correct3 = func.f_f(-0.0L);
+ float err2 =
+ Bruteforce_Ulp_Error_Double(test, correct2);
+ float err3 =
+ Bruteforce_Ulp_Error_Double(test, correct3);
+ fail = fail
+ && ((!(fabsf(err2) <= ulps))
+ && (!(fabsf(err3) <= ulps)));
+ if (fabsf(err2) < fabsf(err)) err = err2;
+ if (fabsf(err3) < fabsf(err)) err = err3;
+
+ // retry per section 6.5.3.4
+ if (IsDoubleResultSubnormal(correct2, ulps)
+ || IsDoubleResultSubnormal(correct3, ulps))
+ {
+ fail = fail && (test != 0.0f);
+ if (!fail) err = 0.0f;
+ }
+ }
+ }
+ }
+ if (fabsf(err) > tinfo->maxError)
+ {
+ tinfo->maxError = fabsf(err);
+ tinfo->maxErrorValue = s[j];
+ }
+ if (fail)
+ {
+ vlog_error("\nERROR: %s%s: %f ulp error at %.13la "
+ "(0x%16.16llx): *%.13la vs. %.13la\n",
+ job->f->name, sizeNames[k], err,
+ ((cl_double *)gIn)[j], ((cl_ulong *)gIn)[j],
+ ((cl_double *)gOut_Ref)[j], test);
+ return -1;
+ }
+ }
+ }
+ }
+
+ for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
+ {
+ if ((error = clEnqueueUnmapMemObject(tinfo->tQueue, tinfo->outBuf[j],
+ out[j], 0, NULL, NULL)))
+ {
+ vlog_error("Error: clEnqueueUnmapMemObject %d failed 2! err: %d\n",
+ j, error);
+ return error;
+ }
+ }
+
+ if ((error = clFlush(tinfo->tQueue))) vlog("clFlush 3 failed\n");
+
+
+ if (0 == (base & 0x0fffffff))
+ {
+ if (gVerboseBruteForce)
+ {
+ vlog("base:%14u step:%10u scale:%10zd buf_elements:%10u ulps:%5.3f "
+ "ThreadCount:%2u\n",
+ base, job->step, buffer_elements, job->scale, job->ulps,
+ job->threadCount);
+ }
+ else
+ {
+ vlog(".");
+ }
+ fflush(stdout);
+ }
+
+ return CL_SUCCESS;
+}
