aboutsummaryrefslogtreecommitdiff
path: root/test_conformance/math_brute_force/ternary_double.cpp
diff options
context:
space:
mode:
Diffstat (limited to 'test_conformance/math_brute_force/ternary_double.cpp')
-rw-r--r--test_conformance/math_brute_force/ternary_double.cpp140
1 files changed, 32 insertions, 108 deletions
diff --git a/test_conformance/math_brute_force/ternary_double.cpp b/test_conformance/math_brute_force/ternary_double.cpp
index 606fdc5a..b5f1ab09 100644
--- a/test_conformance/math_brute_force/ternary_double.cpp
+++ b/test_conformance/math_brute_force/ternary_double.cpp
@@ -14,127 +14,49 @@
// limitations under the License.
//
+#include "common.h"
#include "function_list.h"
#include "test_functions.h"
#include "utility.h"
+#include <cinttypes>
#include <cstring>
#define CORRECTLY_ROUNDED 0
#define FLUSHED 1
-static 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",
- sizeNames[vectorSize],
- "( __global double",
- sizeNames[vectorSize],
- "* out, __global double",
- sizeNames[vectorSize],
- "* in1, __global double",
- sizeNames[vectorSize],
- "* in2, __global double",
- sizeNames[vectorSize],
- "* in3 )\n"
- "{\n"
- " size_t i = get_global_id(0);\n"
- " out[i] = ",
- name,
- "( in1[i], in2[i], in3[i] );\n"
- "}\n" };
-
- const char *c3[] = {
- "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
- "__kernel void math_kernel",
- sizeNames[vectorSize],
- "( __global double* out, __global double* in, __global double* in2, "
- "__global double* in3)\n"
- "{\n"
- " size_t i = get_global_id(0);\n"
- " if( i + 1 < get_global_size(0) )\n"
- " {\n"
- " double3 d0 = vload3( 0, in + 3 * i );\n"
- " double3 d1 = vload3( 0, in2 + 3 * i );\n"
- " double3 d2 = vload3( 0, in3 + 3 * i );\n"
- " d0 = ",
- name,
- "( d0, d1, d2 );\n"
- " vstore3( d0, 0, out + 3*i );\n"
- " }\n"
- " else\n"
- " {\n"
- " size_t parity = i & 1; // Figure out how many elements are "
- "left over after BUFFER_SIZE % (3*sizeof(float)). Assume power of two "
- "buffer size \n"
- " double3 d0;\n"
- " double3 d1;\n"
- " double3 d2;\n"
- " switch( parity )\n"
- " {\n"
- " case 1:\n"
- " d0 = (double3)( in[3*i], NAN, NAN ); \n"
- " d1 = (double3)( in2[3*i], NAN, NAN ); \n"
- " d2 = (double3)( in3[3*i], NAN, NAN ); \n"
- " break;\n"
- " case 0:\n"
- " d0 = (double3)( in[3*i], in[3*i+1], NAN ); \n"
- " d1 = (double3)( in2[3*i], in2[3*i+1], NAN ); \n"
- " d2 = (double3)( in3[3*i], in3[3*i+1], NAN ); \n"
- " break;\n"
- " }\n"
- " d0 = ",
- name,
- "( d0, d1, d2 );\n"
- " switch( parity )\n"
- " {\n"
- " case 0:\n"
- " out[3*i+1] = d0.y; \n"
- " // fall through\n"
- " case 1:\n"
- " out[3*i] = d0.x; \n"
- " break;\n"
- " }\n"
- " }\n"
- "}\n"
- };
-
- const char **kern = c;
- size_t kernSize = sizeof(c) / sizeof(c[0]);
-
- if (sizeValues[vectorSize] == 3)
- {
- kern = c3;
- kernSize = sizeof(c3) / sizeof(c3[0]);
- }
-
- char testName[32];
- snprintf(testName, sizeof(testName) - 1, "math_kernel%s",
- sizeNames[vectorSize]);
+namespace {
- return MakeKernel(kern, (cl_uint)kernSize, testName, k, p, relaxedMode);
+int BuildKernel(const char *name, int vectorSize, cl_kernel *k, cl_program *p,
+ bool relaxedMode)
+{
+ auto kernel_name = GetKernelName(vectorSize);
+ auto source = GetTernaryKernel(kernel_name, name, ParameterType::Double,
+ ParameterType::Double, ParameterType::Double,
+ ParameterType::Double, vectorSize);
+ std::array<const char *, 1> sources{ source.c_str() };
+ return MakeKernel(sources.data(), sources.size(), kernel_name.c_str(), 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);
}
// A table of more difficult cases to get right
-static const double specialValues[] = {
+const double specialValues[] = {
-NAN,
-INFINITY,
-DBL_MAX,
@@ -202,14 +124,16 @@ static const double specialValues[] = {
+0.0,
};
-static const size_t specialValuesCount =
+constexpr size_t specialValuesCount =
sizeof(specialValues) / sizeof(specialValues[0]);
+} // anonymous namespace
+
int TestFunc_Double_Double_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;
int ftz = f->ftz || gForceFTZ;
@@ -224,8 +148,8 @@ int TestFunc_Double_Double_Double_Double(const Func *f, MTdata d,
// 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)))
@@ -387,7 +311,7 @@ int TestFunc_Double_Double_Double_Double(const Func *f, MTdata d,
float err = Bruteforce_Ulp_Error_Double(test, correct);
int fail = !(fabsf(err) <= f->double_ulps);
- if (fail && ftz)
+ if (fail && (ftz || relaxedMode))
{
// retry per section 6.5.3.2
if (IsDoubleSubnormal(correct))
@@ -704,8 +628,9 @@ int TestFunc_Double_Double_Double_Double(const Func *f, MTdata d,
{
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
{
@@ -733,7 +658,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-10 04:42:15 +0000