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.cpp236
1 files changed, 78 insertions, 158 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 bbfd707b..0dc5b9f9 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
@@ -25,115 +25,18 @@
namespace {
-int BuildKernel(const char *name, int vectorSize, cl_kernel *k, cl_program *p,
- bool relaxedMode)
+cl_int BuildKernelFn(cl_uint job_id, cl_uint thread_id UNUSED, void *p)
{
- const char *c[] = { "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
- "__kernel void math_kernel",
- sizeNames[vectorSize],
- "( __global double",
- sizeNames[vectorSize],
- "* out, __global int",
- sizeNames[vectorSize],
- "* out2, __global double",
- sizeNames[vectorSize],
- "* in1, __global double",
- sizeNames[vectorSize],
- "* in2 )\n"
- "{\n"
- " size_t i = get_global_id(0);\n"
- " out[i] = ",
- name,
- "( in1[i], in2[i], out2 + i );\n"
- "}\n" };
-
- const char *c3[] = {
- "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
- "__kernel void math_kernel",
- sizeNames[vectorSize],
- "( __global double* out, __global int* out2, __global double* in, "
- "__global double* in2)\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"
- " int3 i0 = 0xdeaddead;\n"
- " d0 = ",
- name,
- "( d0, d1, &i0 );\n"
- " vstore3( d0, 0, out + 3*i );\n"
- " vstore3( i0, 0, out2 + 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"
- " int3 i0 = 0xdeaddead;\n"
- " switch( parity )\n"
- " {\n"
- " case 1:\n"
- " d0 = (double3)( in[3*i], NAN, NAN ); \n"
- " d1 = (double3)( in2[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"
- " break;\n"
- " }\n"
- " d0 = ",
- name,
- "( d0, d1, &i0 );\n"
- " switch( parity )\n"
- " {\n"
- " case 0:\n"
- " out[3*i+1] = d0.y; \n"
- " out2[3*i+1] = i0.y; \n"
- " // fall through\n"
- " case 1:\n"
- " out[3*i] = d0.x; \n"
- " out2[3*i] = i0.x; \n"
- " break;\n"
- " }\n"
- " }\n"
- "}\n"
+ BuildKernelInfo &info = *(BuildKernelInfo *)p;
+ auto generator = [](const std::string &kernel_name, const char *builtin,
+ cl_uint vector_size_index) {
+ return GetBinaryKernel(kernel_name, builtin, ParameterType::Double,
+ ParameterType::Int, ParameterType::Double,
+ ParameterType::Double, vector_size_index);
};
-
- 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]);
-
- return MakeKernel(kern, (cl_uint)kernSize, testName, k, p, relaxedMode);
+ return BuildKernels(info, job_id, generator);
}
-struct BuildKernelInfo2
-{
- cl_kernel *kernels;
- Programs &programs;
- const char *nameInCode;
- bool relaxedMode; // Whether to build with -cl-fast-relaxed-math.
-};
-
-cl_int BuildKernelFn(cl_uint job_id, cl_uint thread_id UNUSED, void *p)
-{
- BuildKernelInfo2 *info = (BuildKernelInfo2 *)p;
- cl_uint vectorSize = gMinVectorSizeIndex + job_id;
- return BuildKernel(info->nameInCode, vectorSize, info->kernels + vectorSize,
- &(info->programs[vectorSize]), info->relaxedMode);
-}
struct ComputeReferenceInfoD
{
@@ -174,7 +77,8 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
{
int error;
Programs programs;
- cl_kernel kernels[VECTOR_SIZE_COUNT];
+ const unsigned thread_id = 0; // Test is currently not multithreaded.
+ KernelMatrix kernels;
float maxError = 0.0f;
int64_t maxError2 = 0;
int ftz = f->ftz || gForceFTZ;
@@ -191,14 +95,12 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
int testingRemquo = !strcmp(f->name, "remquo");
// Init the kernels
- {
- BuildKernelInfo2 build_info{ kernels, programs, f->nameInCode,
- relaxedMode };
- if ((error = ThreadPool_Do(BuildKernelFn,
- gMaxVectorSizeIndex - gMinVectorSizeIndex,
- &build_info)))
- return error;
- }
+ BuildKernelInfo build_info{ 1, kernels, programs, f->nameInCode,
+ relaxedMode };
+ if ((error = ThreadPool_Do(BuildKernelFn,
+ gMaxVectorSizeIndex - gMinVectorSizeIndex,
+ &build_info)))
+ return error;
for (uint64_t i = 0; i < (1ULL << 32); i += step)
{
@@ -225,28 +127,53 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
return error;
}
- // write garbage into output arrays
+ // Write garbage into output arrays
for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
{
uint32_t pattern = 0xffffdead;
- memset_pattern4(gOut[j], &pattern, BUFFER_SIZE);
- if ((error =
- clEnqueueWriteBuffer(gQueue, gOutBuffer[j], CL_FALSE, 0,
- BUFFER_SIZE, gOut[j], 0, NULL, NULL)))
+ if (gHostFill)
{
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2(%d) ***\n",
- error, j);
- goto exit;
- }
+ memset_pattern4(gOut[j], &pattern, BUFFER_SIZE);
+ if ((error = clEnqueueWriteBuffer(gQueue, gOutBuffer[j],
+ CL_FALSE, 0, BUFFER_SIZE,
+ gOut[j], 0, NULL, NULL)))
+ {
+ vlog_error(
+ "\n*** Error %d in clEnqueueWriteBuffer2(%d) ***\n",
+ error, j);
+ return error;
+ }
- memset_pattern4(gOut2[j], &pattern, BUFFER_SIZE);
- if ((error = clEnqueueWriteBuffer(gQueue, gOutBuffer2[j], CL_FALSE,
- 0, BUFFER_SIZE, gOut2[j], 0, NULL,
- NULL)))
+ memset_pattern4(gOut2[j], &pattern, BUFFER_SIZE);
+ if ((error = clEnqueueWriteBuffer(gQueue, gOutBuffer2[j],
+ CL_FALSE, 0, BUFFER_SIZE,
+ gOut2[j], 0, NULL, NULL)))
+ {
+ vlog_error(
+ "\n*** Error %d in clEnqueueWriteBuffer2b(%d) ***\n",
+ error, j);
+ return error;
+ }
+ }
+ else
{
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2b(%d) ***\n",
- error, j);
- goto exit;
+ if ((error = clEnqueueFillBuffer(gQueue, gOutBuffer[j],
+ &pattern, sizeof(pattern), 0,
+ BUFFER_SIZE, 0, NULL, NULL)))
+ {
+ vlog_error("Error: clEnqueueFillBuffer 1 failed! err: %d\n",
+ error);
+ return error;
+ }
+
+ if ((error = clEnqueueFillBuffer(gQueue, gOutBuffer2[j],
+ &pattern, sizeof(pattern), 0,
+ BUFFER_SIZE, 0, NULL, NULL)))
+ {
+ vlog_error("Error: clEnqueueFillBuffer 2 failed! err: %d\n",
+ error);
+ return error;
+ }
}
}
@@ -256,37 +183,38 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
size_t vectorSize = sizeof(cl_double) * sizeValues[j];
size_t localCount = (BUFFER_SIZE + vectorSize - 1)
/ vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
+ if ((error = clSetKernelArg(kernels[j][thread_id], 0,
+ sizeof(gOutBuffer[j]), &gOutBuffer[j])))
{
LogBuildError(programs[j]);
- goto exit;
+ return error;
}
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gOutBuffer2[j]),
- &gOutBuffer2[j])))
+ if ((error =
+ clSetKernelArg(kernels[j][thread_id], 1,
+ sizeof(gOutBuffer2[j]), &gOutBuffer2[j])))
{
LogBuildError(programs[j]);
- goto exit;
+ return error;
}
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer),
- &gInBuffer)))
+ if ((error = clSetKernelArg(kernels[j][thread_id], 2,
+ sizeof(gInBuffer), &gInBuffer)))
{
LogBuildError(programs[j]);
- goto exit;
+ return error;
}
- if ((error = clSetKernelArg(kernels[j], 3, sizeof(gInBuffer2),
- &gInBuffer2)))
+ if ((error = clSetKernelArg(kernels[j][thread_id], 3,
+ sizeof(gInBuffer2), &gInBuffer2)))
{
LogBuildError(programs[j]);
- goto exit;
+ return error;
}
- if ((error =
- clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL, NULL)))
+ if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j][thread_id],
+ 1, NULL, &localCount, NULL, 0,
+ NULL, NULL)))
{
vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
+ return error;
}
}
@@ -325,14 +253,14 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
BUFFER_SIZE, gOut[j], 0, NULL, NULL)))
{
vlog_error("ReadArray failed %d\n", error);
- goto exit;
+ return error;
}
if ((error =
clEnqueueReadBuffer(gQueue, gOutBuffer2[j], CL_TRUE, 0,
BUFFER_SIZE, gOut2[j], 0, NULL, NULL)))
{
vlog_error("ReadArray2 failed %d\n", error);
- goto exit;
+ return error;
}
}
@@ -542,8 +470,7 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
((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;
+ return -1;
}
}
}
@@ -577,12 +504,5 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode)
vlog("\n");
-exit:
- // Release
- for (auto k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
- {
- clReleaseKernel(kernels[k]);
- }
-
- return error;
+ return CL_SUCCESS;
}
generated by cgit v1.2.3 (git 2.25.1) at 2024-06-13 05:39:34 +0000