From a53917a37e6c1a38e62ac513ecba4a013b6f876e Mon Sep 17 00:00:00 2001 From: Marco Antognini Date: Tue, 9 Mar 2021 22:55:33 +0000 Subject: [PATCH] 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 --- .../math_brute_force/binary_double.cpp | 8 +- .../math_brute_force/binary_float.cpp | 79 +-- .../math_brute_force/binary_i_double.cpp | 8 +- .../math_brute_force/binary_i_float.cpp | 72 +-- .../binary_operator_double.cpp | 4 +- .../binary_operator_float.cpp | 76 +-- .../math_brute_force/macro_binary_double.cpp | 8 +- .../math_brute_force/macro_binary_float.cpp | 58 +-- .../math_brute_force/ternary_float.cpp | 2 +- .../math_brute_force/unary_double.cpp | 478 +++++++++--------- 10 files changed, 398 insertions(+), 395 deletions(-) diff --git a/test_conformance/math_brute_force/binary_double.cpp b/test_conformance/math_brute_force/binary_double.cpp index fad03ade..cbb186ed 100644 --- a/test_conformance/math_brute_force/binary_double.cpp +++ b/test_conformance/math_brute_force/binary_double.cpp @@ -186,8 +186,8 @@ static const double specialValuesDouble[] = { 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 @@ static const double specialValuesDouble[] = { 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 a31bfb2f..8dfb9f40 100644 --- a/test_conformance/math_brute_force/binary_float.cpp +++ b/test_conformance/math_brute_force/binary_float.cpp @@ -126,6 +126,45 @@ static cl_int BuildKernel_FloatFn(cl_uint job_id, cl_uint thread_id UNUSED, 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 @@ static const float specialValuesFloat[] = { 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 6839dfb9..eb2c59e0 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 @@ static const double specialValuesDouble[] = { 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 @@ static const double specialValuesDouble[] = { 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 ceb79ddf..019c96a5 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 @@ static cl_int BuildKernel_FloatFn(cl_uint job_id, cl_uint thread_id UNUSED, 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 @@ static const float specialValuesFloat[] = { 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 const int specialValuesInt[] = { 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 939ea6d6..a7fb3cd0 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 @@ static const double specialValuesDouble[] = { 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 efef4fe5..a9d3b7cf 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 @@ static cl_int BuildKernel_FloatFn(cl_uint job_id, cl_uint thread_id UNUSED, 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 @@ static const float specialValuesFloat[] = { 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 81bc4d0b..2ea785e6 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 @@ static const double specialValuesDouble[] = { 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 @@ static const double specialValuesDouble[] = { 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 1b5dc33a..a61ab6b3 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 @@ static cl_int BuildKernel_FloatFn(cl_uint job_id, cl_uint thread_id UNUSED, 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 @@ static const float specialValuesFloat[] = { 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 3b2adf80..1fcdc4a2 100644 --- a/test_conformance/math_brute_force/ternary_float.cpp +++ b/test_conformance/math_brute_force/ternary_float.cpp @@ -208,7 +208,7 @@ static const float specialValuesFloat[] = { 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 1ff3d9c8..99959ae3 100644 --- a/test_conformance/math_brute_force/unary_double.cpp +++ b/test_conformance/math_brute_force/unary_double.cpp @@ -160,244 +160,7 @@ typedef struct TestInfo // 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 @@ exit: 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; +}