aboutsummaryrefslogtreecommitdiff
path: root/test_conformance/math_brute_force/binary_two_results_i_double.cpp
diff options
context:
space:
mode:
Diffstat (limited to 'test_conformance/math_brute_force/binary_two_results_i_double.cpp')
-rw-r--r--test_conformance/math_brute_force/binary_two_results_i_double.cpp76
1 files changed, 42 insertions, 34 deletions
diff --git a/test_conformance/math_brute_force/binary_two_results_i_double.cpp b/test_conformance/math_brute_force/binary_two_results_i_double.cpp
index 14f41092..bbfd707b 100644
--- a/test_conformance/math_brute_force/binary_two_results_i_double.cpp
+++ b/test_conformance/math_brute_force/binary_two_results_i_double.cpp
@@ -14,15 +14,19 @@
// limitations under the License.
//
+#include "common.h"
#include "function_list.h"
#include "test_functions.h"
#include "utility.h"
+#include <cinttypes>
#include <climits>
#include <cstring>
-static int BuildKernel(const char *name, int vectorSize, cl_kernel *k,
- cl_program *p, bool relaxedMode)
+namespace {
+
+int BuildKernel(const char *name, int vectorSize, cl_kernel *k, cl_program *p,
+ bool relaxedMode)
{
const char *c[] = { "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
"__kernel void math_kernel",
@@ -115,24 +119,23 @@ static int BuildKernel(const char *name, int vectorSize, cl_kernel *k,
return MakeKernel(kern, (cl_uint)kernSize, testName, k, p, relaxedMode);
}
-typedef struct BuildKernelInfo
+struct BuildKernelInfo2
{
- cl_uint offset; // the first vector size to build
cl_kernel *kernels;
- cl_program *programs;
+ Programs &programs;
const char *nameInCode;
bool relaxedMode; // Whether to build with -cl-fast-relaxed-math.
-} BuildKernelInfo;
+};
-static cl_int BuildKernelFn(cl_uint job_id, cl_uint thread_id UNUSED, void *p)
+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 BuildKernel(info->nameInCode, i, info->kernels + i,
- info->programs + i, info->relaxedMode);
+ BuildKernelInfo2 *info = (BuildKernelInfo2 *)p;
+ cl_uint vectorSize = gMinVectorSizeIndex + job_id;
+ return BuildKernel(info->nameInCode, vectorSize, info->kernels + vectorSize,
+ &(info->programs[vectorSize]), info->relaxedMode);
}
-typedef struct ComputeReferenceInfoD_
+struct ComputeReferenceInfoD
{
const double *x;
const double *y;
@@ -141,9 +144,9 @@ typedef struct ComputeReferenceInfoD_
long double (*f_ffpI)(long double, long double, int *);
cl_uint lim;
cl_uint count;
-} ComputeReferenceInfoD;
+};
-static cl_int ReferenceD(cl_uint jid, cl_uint tid, void *userInfo)
+cl_int ReferenceD(cl_uint jid, cl_uint tid, void *userInfo)
{
ComputeReferenceInfoD *cri = (ComputeReferenceInfoD *)userInfo;
cl_uint lim = cri->lim;
@@ -165,10 +168,12 @@ static cl_int ReferenceD(cl_uint jid, cl_uint tid, void *userInfo)
return CL_SUCCESS;
}
+} // anonymous namespace
+
int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
{
int error;
- cl_program programs[VECTOR_SIZE_COUNT];
+ Programs programs;
cl_kernel kernels[VECTOR_SIZE_COUNT];
float maxError = 0.0f;
int64_t maxError2 = 0;
@@ -187,8 +192,8 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
// Init the kernels
{
- BuildKernelInfo build_info = { gMinVectorSizeIndex, kernels, programs,
- f->nameInCode, relaxedMode };
+ BuildKernelInfo2 build_info{ kernels, programs, f->nameInCode,
+ relaxedMode };
if ((error = ThreadPool_Do(BuildKernelFn,
gMaxVectorSizeIndex - gMinVectorSizeIndex,
&build_info)))
@@ -375,7 +380,7 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
if (iptrUndefined) iErr = 0;
int fail = !(fabsf(err) <= f->double_ulps && iErr == 0);
- if (ftz && fail)
+ if ((ftz || relaxedMode) && fail)
{
// retry per section 6.5.3.2
if (IsDoubleResultSubnormal(correct, f->double_ulps))
@@ -523,17 +528,20 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
if (fail)
{
- vlog_error(
- "\nERROR: %sD%s: {%f, %lld} ulp error at {%.13la, "
- "%.13la} ({ 0x%16.16llx, 0x%16.16llx}): *{%.13la, "
- "%d} ({ 0x%16.16llx, 0x%8.8x}) vs. {%.13la, %d} ({ "
- "0x%16.16llx, 0x%8.8x})\n",
- f->name, sizeNames[k], err, iErr, ((double *)gIn)[j],
- ((double *)gIn2)[j], ((cl_ulong *)gIn)[j],
- ((cl_ulong *)gIn2)[j], ((double *)gOut_Ref)[j],
- ((int *)gOut_Ref2)[j], ((cl_ulong *)gOut_Ref)[j],
- ((cl_uint *)gOut_Ref2)[j], test, q2[j],
- ((cl_ulong *)q)[j], ((cl_uint *)q2)[j]);
+ vlog_error("\nERROR: %sD%s: {%f, %" PRId64
+ "} ulp error at {%.13la, "
+ "%.13la} ({ 0x%16.16" PRIx64 ", 0x%16.16" PRIx64
+ "}): *{%.13la, "
+ "%d} ({ 0x%16.16" PRIx64
+ ", 0x%8.8x}) vs. {%.13la, %d} ({ "
+ "0x%16.16" PRIx64 ", 0x%8.8x})\n",
+ f->name, sizeNames[k], err, iErr,
+ ((double *)gIn)[j], ((double *)gIn2)[j],
+ ((cl_ulong *)gIn)[j], ((cl_ulong *)gIn2)[j],
+ ((double *)gOut_Ref)[j], ((int *)gOut_Ref2)[j],
+ ((cl_ulong *)gOut_Ref)[j],
+ ((cl_uint *)gOut_Ref2)[j], test, q2[j],
+ ((cl_ulong *)q)[j], ((cl_uint *)q2)[j]);
error = -1;
goto exit;
}
@@ -544,8 +552,9 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
{
if (gVerboseBruteForce)
{
- vlog("base:%14u step:%10zu bufferSize:%10zd \n", i, step,
- BUFFER_SIZE);
+ vlog("base:%14" PRIu64 " step:%10" PRIu64
+ " bufferSize:%10d \n",
+ i, step, BUFFER_SIZE);
}
else
{
@@ -562,8 +571,8 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
else
vlog("passed");
- vlog("\t{%8.2f, %lld} @ {%a, %a}", maxError, maxError2, maxErrorVal,
- maxErrorVal2);
+ vlog("\t{%8.2f, %" PRId64 "} @ {%a, %a}", maxError, maxError2,
+ maxErrorVal, maxErrorVal2);
}
vlog("\n");
@@ -573,7 +582,6 @@ exit:
for (auto k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
{
clReleaseKernel(kernels[k]);
- clReleaseProgram(programs[k]);
}
return error;
generated by cgit v1.2.3 (git 2.25.1) at 2024-06-09 20:30:28 +0000