Remove unsupported code (#1211)
* Remove code for runtime measurement
The GetTime() and associated functions are not fully implemented on
Linux. This functionality is assumed to be untested, or unused at best.
Reduce differences between tests by removing this unnecessary feature.
It can be (re-)implemented later, if desired, once the math_brute_force
component is in better shape.
Signed-off-by: Marco Antognini <marco.antognini@arm.com>
* Coalesce if-statements
Signed-off-by: Marco Antognini <marco.antognini@arm.com>
* Keep else branch
Address comments.
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 15a48a6..e51327d 100644
--- a/test_conformance/math_brute_force/binary_double.cpp
+++ b/test_conformance/math_brute_force/binary_double.cpp
@@ -436,95 +436,10 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input arrays
- double *p = (double *)gIn;
- double *p2 = (double *)gIn2;
- for (j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
- {
- p[j] = DoubleFromUInt32(genrand_int32(d));
- p2[j] = DoubleFromUInt32(genrand_int32(d));
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- BUFFER_SIZE, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_double) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 2,
- sizeof(gInBuffer2), &gInBuffer2)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t%8.2f @ {%a, %a}", maxError, maxErrorVal, maxErrorVal2);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/binary_float.cpp b/test_conformance/math_brute_force/binary_float.cpp
index fb56c64..5a37a40 100644
--- a/test_conformance/math_brute_force/binary_float.cpp
+++ b/test_conformance/math_brute_force/binary_float.cpp
@@ -428,95 +428,10 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input arrays
- cl_uint *p = (cl_uint *)gIn;
- cl_uint *p2 = (cl_uint *)gIn2;
- for (j = 0; j < BUFFER_SIZE / sizeof(float); j++)
- {
- p[j] = (genrand_int32(d) & ~0x40000000) | 0x20000000;
- p2[j] = 0x3fc00000;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- BUFFER_SIZE, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_float) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 2,
- sizeof(gInBuffer2), &gInBuffer2)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t%8.2f @ {%a, %a}", maxError, maxErrorVal, maxErrorVal2);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/binary_i_double.cpp b/test_conformance/math_brute_force/binary_i_double.cpp
index 686e32b..c3c3eb1 100644
--- a/test_conformance/math_brute_force/binary_i_double.cpp
+++ b/test_conformance/math_brute_force/binary_i_double.cpp
@@ -437,96 +437,10 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input arrays
- double *p = (double *)gIn;
- cl_int *p2 = (cl_int *)gIn2;
- for (j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
- {
- p[j] = DoubleFromUInt32(genrand_int32(d));
- p2[j] = 3;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
- if ((error =
- clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- BUFFER_SIZE / 2, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_double) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 2,
- sizeof(gInBuffer2), &gInBuffer2)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t%8.2f @ {%a, %d}", maxError, maxErrorVal, maxErrorVal2);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/binary_i_float.cpp b/test_conformance/math_brute_force/binary_i_float.cpp
index d99382f..93cb910 100644
--- a/test_conformance/math_brute_force/binary_i_float.cpp
+++ b/test_conformance/math_brute_force/binary_i_float.cpp
@@ -430,95 +430,10 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input arrays
- uint32_t *p = (uint32_t *)gIn;
- uint32_t *p2 = (uint32_t *)gIn2;
- for (j = 0; j < BUFFER_SIZE / sizeof(float); j++)
- {
- p[j] = (genrand_int32(d) & ~0x40000000) | 0x38000000;
- p2[j] = 3;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- BUFFER_SIZE, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_float) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 2,
- sizeof(gInBuffer2), &gInBuffer2)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t%8.2f @ {%a, %d}", maxError, maxErrorVal, maxErrorVal2);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/binary_operator_double.cpp b/test_conformance/math_brute_force/binary_operator_double.cpp
index a02e53b..09540b0 100644
--- a/test_conformance/math_brute_force/binary_operator_double.cpp
+++ b/test_conformance/math_brute_force/binary_operator_double.cpp
@@ -430,95 +430,10 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input arrays
- double *p = (double *)gIn;
- double *p2 = (double *)gIn2;
- for (j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
- {
- p[j] = DoubleFromUInt32(genrand_int32(d));
- p2[j] = DoubleFromUInt32(genrand_int32(d));
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- BUFFER_SIZE, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_double) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 2,
- sizeof(gInBuffer2), &gInBuffer2)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t%8.2f @ {%a, %a}", maxError, maxErrorVal, maxErrorVal2);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/binary_operator_float.cpp b/test_conformance/math_brute_force/binary_operator_float.cpp
index 39070cb..a59be16 100644
--- a/test_conformance/math_brute_force/binary_operator_float.cpp
+++ b/test_conformance/math_brute_force/binary_operator_float.cpp
@@ -422,95 +422,10 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input arrays
- uint32_t *p = (uint32_t *)gIn;
- uint32_t *p2 = (uint32_t *)gIn2;
- for (j = 0; j < BUFFER_SIZE / sizeof(float); j++)
- {
- p[j] = (genrand_int32(d) & ~0x40000000) | 0x20000000;
- p2[j] = 0x3fc00000;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- BUFFER_SIZE, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_float) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 2,
- sizeof(gInBuffer2), &gInBuffer2)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t%8.2f @ {%a, %a}", maxError, maxErrorVal, maxErrorVal2);
+ }
+
vlog("\n");
exit:
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 d71585e..f548486 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
@@ -565,96 +565,10 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input array
- double *p = (double *)gIn;
- for (j = 0; j < bufferSize / sizeof(cl_double); j++)
- p[j] = DoubleFromUInt32(genrand_int32(d));
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- bufferSize, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_double) * sizeValues[j];
- size_t localCount = (bufferSize + vectorSize - 1)
- / vectorSize; // bufferSize / vectorSize rounded up
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gOutBuffer2[j]),
- &gOutBuffer2[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 3, sizeof(gInBuffer2),
- &gInBuffer2)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t{%8.2f, %lld} @ %a", maxError, maxError2, maxErrorVal);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/binary_two_results_i_float.cpp b/test_conformance/math_brute_force/binary_two_results_i_float.cpp
index b6f1f1b..01c3ca7 100644
--- a/test_conformance/math_brute_force/binary_two_results_i_float.cpp
+++ b/test_conformance/math_brute_force/binary_two_results_i_float.cpp
@@ -550,98 +550,10 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input array
- cl_uint *p = (cl_uint *)gIn;
- for (j = 0; j < bufferSize / sizeof(float); j++)
- {
- p[j] = genrand_int32(d);
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- bufferSize, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_float) * sizeValues[j];
- size_t localCount = (bufferSize + vectorSize - 1)
- / vectorSize; // bufferSize / vectorSize rounded up
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gOutBuffer2[j]),
- &gOutBuffer2[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 3, sizeof(gInBuffer2),
- &gInBuffer2)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t{%8.2f, %lld} @ %a", maxError, maxError2, maxErrorVal);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/i_unary_double.cpp b/test_conformance/math_brute_force/i_unary_double.cpp
index c3822f3..040f8e6 100644
--- a/test_conformance/math_brute_force/i_unary_double.cpp
+++ b/test_conformance/math_brute_force/i_unary_double.cpp
@@ -289,73 +289,6 @@
vlog("passed");
}
- if (gMeasureTimes)
- {
- // Init input array
- double *p = (double *)gIn;
- for (j = 0; j < bufferSize / sizeof(cl_double); j++)
- p[j] = DoubleFromUInt32(genrand_int32(d));
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeValues[j] * sizeof(cl_double);
- size_t localCount = (bufferSize + vectorSize - 1) / vectorSize;
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/i_unary_float.cpp b/test_conformance/math_brute_force/i_unary_float.cpp
index d1fb867..61ce9a2 100644
--- a/test_conformance/math_brute_force/i_unary_float.cpp
+++ b/test_conformance/math_brute_force/i_unary_float.cpp
@@ -286,74 +286,6 @@
vlog("passed");
}
- if (gMeasureTimes)
- {
- // Init input array
- uint32_t *p = (uint32_t *)gIn;
- for (j = 0; j < bufferSize / sizeof(float); j++)
- p[j] = genrand_int32(d);
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_float) * sizeValues[j];
- size_t localCount = (bufferSize + vectorSize - 1)
- / vectorSize; // bufferSize / vectorSize rounded up
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/macro_binary_double.cpp b/test_conformance/math_brute_force/macro_binary_double.cpp
index cbcecb3..7c40a18 100644
--- a/test_conformance/math_brute_force/macro_binary_double.cpp
+++ b/test_conformance/math_brute_force/macro_binary_double.cpp
@@ -406,93 +406,6 @@
vlog("passed");
}
- if (gMeasureTimes)
- {
- // Init input arrays
- cl_ulong *p = (cl_ulong *)gIn;
- cl_ulong *p2 = (cl_ulong *)gIn2;
- for (j = 0; j < BUFFER_SIZE / sizeof(double); j++)
- {
- p[j] =
- (cl_ulong)genrand_int32(d) | ((cl_ulong)genrand_int32(d) << 32);
- p2[j] =
- (cl_ulong)genrand_int32(d) | ((cl_ulong)genrand_int32(d) << 32);
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- BUFFER_SIZE, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_double) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 2,
- sizeof(gInBuffer2), &gInBuffer2)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/macro_binary_float.cpp b/test_conformance/math_brute_force/macro_binary_float.cpp
index bbccbbe..12c14dc 100644
--- a/test_conformance/math_brute_force/macro_binary_float.cpp
+++ b/test_conformance/math_brute_force/macro_binary_float.cpp
@@ -397,91 +397,6 @@
vlog("passed");
}
- if (gMeasureTimes)
- {
- // Init input arrays
- cl_uint *p = (cl_uint *)gIn;
- cl_uint *p2 = (cl_uint *)gIn2;
- for (j = 0; j < BUFFER_SIZE / sizeof(float); j++)
- {
- p[j] = genrand_int32(d);
- p2[j] = genrand_int32(d);
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- BUFFER_SIZE, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_float) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 2,
- sizeof(gInBuffer2), &gInBuffer2)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/macro_unary_double.cpp b/test_conformance/math_brute_force/macro_unary_double.cpp
index 154d8ec..4fadbfc 100644
--- a/test_conformance/math_brute_force/macro_unary_double.cpp
+++ b/test_conformance/math_brute_force/macro_unary_double.cpp
@@ -272,74 +272,6 @@
vlog("passed");
}
- if (gMeasureTimes)
- {
- // Init input array
- cl_ulong *p = (cl_ulong *)gIn;
- for (j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
- p[j] = DoubleFromUInt32(genrand_int32(d));
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_double) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/macro_unary_float.cpp b/test_conformance/math_brute_force/macro_unary_float.cpp
index 725a831..4e72748 100644
--- a/test_conformance/math_brute_force/macro_unary_float.cpp
+++ b/test_conformance/math_brute_force/macro_unary_float.cpp
@@ -272,74 +272,6 @@
vlog("passed");
}
- if (gMeasureTimes)
- {
- // Init input array
- cl_uint *p = (cl_uint *)gIn;
- for (j = 0; j < BUFFER_SIZE / sizeof(float); j++)
- p[j] = genrand_int32(d);
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_float) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/mad_double.cpp b/test_conformance/math_brute_force/mad_double.cpp
index 5e7dba9..a5f5918 100644
--- a/test_conformance/math_brute_force/mad_double.cpp
+++ b/test_conformance/math_brute_force/mad_double.cpp
@@ -285,111 +285,11 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input array
- double *p = (double *)gIn;
- double *p2 = (double *)gIn2;
- double *p3 = (double *)gIn3;
- for (j = 0; j < bufferSize / sizeof(double); j++)
- {
- p[j] = DoubleFromUInt32(genrand_int32(d));
- p2[j] = DoubleFromUInt32(genrand_int32(d));
- p3[j] = DoubleFromUInt32(genrand_int32(d));
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- bufferSize, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer3, CL_FALSE, 0,
- bufferSize, gIn3, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer3 ***\n", error);
- return error;
- }
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_double) * sizeValues[j];
- size_t localCount = (bufferSize + vectorSize - 1)
- / vectorSize; // bufferSize / vectorSize rounded up
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer2),
- &gInBuffer2)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 3, sizeof(gInBuffer3),
- &gInBuffer3)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t%8.2f @ {%a, %a, %a}", maxError, maxErrorVal, maxErrorVal2,
maxErrorVal3);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/mad_float.cpp b/test_conformance/math_brute_force/mad_float.cpp
index 453e43e..ff4e99d 100644
--- a/test_conformance/math_brute_force/mad_float.cpp
+++ b/test_conformance/math_brute_force/mad_float.cpp
@@ -284,109 +284,11 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input array
- cl_uint *p = (cl_uint *)gIn;
- cl_uint *p2 = (cl_uint *)gIn2;
- cl_uint *p3 = (cl_uint *)gIn3;
- for (j = 0; j < bufferSize / sizeof(float); j++)
- {
- p[j] = genrand_int32(d);
- p2[j] = genrand_int32(d);
- p3[j] = genrand_int32(d);
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- bufferSize, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer3, CL_FALSE, 0,
- bufferSize, gIn3, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer3 ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_float) * sizeValues[j];
- size_t localCount = (bufferSize + vectorSize - 1)
- / vectorSize; // bufferSize / vectorSize rounded up
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer2),
- &gInBuffer2)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 3, sizeof(gInBuffer3),
- &gInBuffer3)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t%8.2f @ {%a, %a, %a}", maxError, maxErrorVal, maxErrorVal2,
maxErrorVal3);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/main.cpp b/test_conformance/math_brute_force/main.cpp
index 0e1b40a..c2a376f 100644
--- a/test_conformance/math_brute_force/main.cpp
+++ b/test_conformance/math_brute_force/main.cpp
@@ -62,8 +62,6 @@
int gSkipCorrectnessTesting = 0;
int gStopOnError = 0;
static bool gSkipRestOfTests;
-int gMeasureTimes = 0;
-int gReportAverageTimes = 0;
int gForceFTZ = 0;
int gWimpyMode = 0;
int gHasDouble = 0;
@@ -87,7 +85,6 @@
int gIsInRTZMode = 0;
uint32_t gMaxVectorSizeIndex = VECTOR_SIZE_COUNT;
uint32_t gMinVectorSizeIndex = 0;
-const char *method[] = { "Best", "Average" };
void *gIn = NULL;
void *gIn2 = NULL;
void *gIn3 = NULL;
@@ -813,24 +810,8 @@
else if (gStopOnError)
vlog("Stopping at first error.\n");
- if (gMeasureTimes)
- {
- vlog("%s times are reported at right (cycles per element):\n",
- method[gReportAverageTimes]);
- vlog("\n");
- if (gSkipCorrectnessTesting)
- vlog(" \t ");
- else
- vlog(" \t ");
- if (gWimpyMode) vlog(" ");
- for (int i = gMinVectorSizeIndex; i < gMaxVectorSizeIndex; i++)
- vlog("\t float%s", sizeNames[i]);
- }
- else
- {
- vlog(" \t ");
- if (gWimpyMode) vlog(" ");
- }
+ vlog(" \t ");
+ if (gWimpyMode) vlog(" ");
if (!gSkipCorrectnessTesting) vlog("\t max_ulps");
vlog("\n-------------------------------------------------------------------"
@@ -905,8 +886,6 @@
optionFound = 1;
switch (*arg)
{
- case 'a': gReportAverageTimes ^= 1; break;
-
case 'c': gToggleCorrectlyRoundedDivideSqrt ^= 1; break;
case 'd': gHasDouble ^= 1; break;
@@ -927,8 +906,6 @@
case 's': gStopOnError ^= 1; break;
- case 't': gMeasureTimes ^= 1; break;
-
case 'v': gVerboseBruteForce ^= 1; break;
case 'w': // wimpy mode
@@ -970,7 +947,6 @@
gMinVectorSizeIndex = 4;
gMaxVectorSizeIndex = gMinVectorSizeIndex + 1;
break;
- break;
default:
vlog(" <-- unknown flag: %c (0x%2.2x)\n)", *arg, *arg);
@@ -1053,9 +1029,8 @@
static void PrintUsage(void)
{
- vlog("%s [-acglstz]: <optional: math function names>\n", appName);
+ vlog("%s [-cglsz]: <optional: math function names>\n", appName);
vlog("\toptions:\n");
- vlog("\t\t-a\tReport average times instead of best times\n");
vlog("\t\t-c\tToggle test fp correctly rounded divide and sqrt (Default: "
"off)\n");
vlog("\t\t-d\tToggle double precision testing. (Default: on iff khr_fp_64 "
@@ -1070,7 +1045,6 @@
"accuracy checks.)\n");
vlog("\t\t-m\tToggle run multi-threaded. (Default: on) )\n");
vlog("\t\t-s\tStop on error\n");
- vlog("\t\t-t\tToggle timing (on by default)\n");
vlog("\t\t-w\tToggle Wimpy Mode, * Not a valid test * \n");
vlog("\t\t-[2^n]\tSet wimpy reduction factor, recommended range of n is "
"1-10, default factor(%u)\n",
@@ -1809,51 +1783,6 @@
return fabs((float)(reference - (double)test));
}
-#if defined(__APPLE__)
-#include <mach/mach_time.h>
-#endif
-
-uint64_t GetTime(void)
-{
-#if defined(__APPLE__)
- return mach_absolute_time();
-#elif defined(_WIN32) && defined(_MSC_VER)
- return ReadTime();
-#else
-// mach_absolute_time is a high precision timer with precision < 1 microsecond.
-#warning need accurate clock here. Times are invalid.
- return 0;
-#endif
-}
-
-
-#if defined(_WIN32) && defined(_MSC_VER)
-/* function is defined in "compat.h" */
-#else
-double SubtractTime(uint64_t endTime, uint64_t startTime)
-{
- uint64_t diff = endTime - startTime;
- static double conversion = 0.0;
-
- if (0.0 == conversion)
- {
-#if defined(__APPLE__)
- mach_timebase_info_data_t info = { 0, 0 };
- kern_return_t err = mach_timebase_info(&info);
- if (0 == err)
- conversion = 1e-9 * (double)info.numer / (double)info.denom;
-#else
-// This function consumes output from GetTime() above, and converts the time to
-// secionds.
-#warning need accurate ticks to seconds conversion factor here. Times are invalid.
-#endif
- }
-
- // strictly speaking we should also be subtracting out timer latency here
- return conversion * (double)diff;
-}
-#endif
-
cl_uint RoundUpToNextPowerOfTwo(cl_uint x)
{
if (0 == (x & (x - 1))) return x;
diff --git a/test_conformance/math_brute_force/ternary_double.cpp b/test_conformance/math_brute_force/ternary_double.cpp
index 96a0e7b..6e0327a 100644
--- a/test_conformance/math_brute_force/ternary_double.cpp
+++ b/test_conformance/math_brute_force/ternary_double.cpp
@@ -723,109 +723,11 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input array
- double *p = (double *)gIn;
- double *p2 = (double *)gIn2;
- double *p3 = (double *)gIn3;
- for (j = 0; j < bufferSize / sizeof(double); j++)
- {
- p[j] = DoubleFromUInt32(genrand_int32(d));
- p2[j] = DoubleFromUInt32(genrand_int32(d));
- p3[j] = DoubleFromUInt32(genrand_int32(d));
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- bufferSize, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer3, CL_FALSE, 0,
- bufferSize, gIn3, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer3 ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_double) * sizeValues[j];
- size_t localCount = (bufferSize + vectorSize - 1)
- / vectorSize; // bufferSize / vectorSize rounded up
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer2),
- &gInBuffer2)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 3, sizeof(gInBuffer3),
- &gInBuffer3)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t%8.2f @ {%a, %a, %a}", maxError, maxErrorVal, maxErrorVal2,
maxErrorVal3);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/ternary_float.cpp b/test_conformance/math_brute_force/ternary_float.cpp
index 20b3eb7..9ef5988 100644
--- a/test_conformance/math_brute_force/ternary_float.cpp
+++ b/test_conformance/math_brute_force/ternary_float.cpp
@@ -858,109 +858,11 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input array
- cl_uint *p = (cl_uint *)gIn;
- cl_uint *p2 = (cl_uint *)gIn2;
- cl_uint *p3 = (cl_uint *)gIn3;
- for (j = 0; j < bufferSize / sizeof(float); j++)
- {
- p[j] = genrand_int32(d);
- p2[j] = genrand_int32(d);
- p3[j] = genrand_int32(d);
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
- bufferSize, gIn2, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
- return error;
- }
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer3, CL_FALSE, 0,
- bufferSize, gIn3, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer3 ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_float) * sizeValues[j];
- size_t localCount = (bufferSize + vectorSize - 1)
- / vectorSize; // bufferSize / vectorSize rounded up
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer2),
- &gInBuffer2)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 3, sizeof(gInBuffer3),
- &gInBuffer3)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t%8.2f @ {%a, %a, %a}", maxError, maxErrorVal, maxErrorVal2,
maxErrorVal3);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/unary_double.cpp b/test_conformance/math_brute_force/unary_double.cpp
index 615d0fb..3dc1548 100644
--- a/test_conformance/math_brute_force/unary_double.cpp
+++ b/test_conformance/math_brute_force/unary_double.cpp
@@ -290,85 +290,10 @@
vlog("Wimp pass");
else
vlog("passed");
+
+ vlog("\t%8.2f @ %a", maxError, maxErrorVal);
}
- if (gMeasureTimes)
- {
- // Init input array
- double *p = (double *)gIn;
-
- if (strstr(f->name, "exp"))
- for (j = 0; j < BUFFER_SIZE / sizeof(double); j++)
- p[j] = (double)genrand_real1(d);
- else if (strstr(f->name, "log"))
- for (j = 0; j < BUFFER_SIZE / sizeof(double); j++)
- p[j] = fabs(DoubleFromUInt32(genrand_int32(d)));
- else
- for (j = 0; j < BUFFER_SIZE / sizeof(double); j++)
- p[j] = DoubleFromUInt32(genrand_int32(d));
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_double) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting) vlog("\t%8.2f @ %a", maxError, maxErrorVal);
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/unary_float.cpp b/test_conformance/math_brute_force/unary_float.cpp
index 3666589..73365d1 100644
--- a/test_conformance/math_brute_force/unary_float.cpp
+++ b/test_conformance/math_brute_force/unary_float.cpp
@@ -315,85 +315,10 @@
vlog(" (rlx skip correctness testing)\n");
goto exit;
}
+
+ vlog("\t%8.2f @ %a", maxError, maxErrorVal);
}
- if (gMeasureTimes)
- {
- // Init input array
- uint32_t *p = (uint32_t *)gIn;
- if (strstr(f->name, "exp") || strstr(f->name, "sin")
- || strstr(f->name, "cos") || strstr(f->name, "tan"))
- for (j = 0; j < BUFFER_SIZE / sizeof(float); j++)
- ((float *)p)[j] = (float)genrand_real1(d);
- else if (strstr(f->name, "log"))
- for (j = 0; j < BUFFER_SIZE / sizeof(float); j++)
- p[j] = genrand_int32(d) & 0x7fffffff;
- else
- for (j = 0; j < BUFFER_SIZE / sizeof(float); j++)
- p[j] = genrand_int32(d);
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- BUFFER_SIZE, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeof(cl_float) * sizeValues[j];
- size_t localCount = (BUFFER_SIZE + vectorSize - 1)
- / vectorSize; // BUFFER_SIZE / vectorSize rounded up
- if ((error = clSetKernelArg(test_info.k[j][0], 0,
- sizeof(gOutBuffer[j]), &gOutBuffer[j])))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(test_info.k[j][0], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(test_info.programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (i = 0; i < PERF_LOOP_COUNT; i++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, test_info.k[j][0],
- 1, NULL, &localCount, NULL,
- 0, NULL, NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (BUFFER_SIZE / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting) vlog("\t%8.2f @ %a", maxError, maxErrorVal);
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/unary_two_results_double.cpp b/test_conformance/math_brute_force/unary_two_results_double.cpp
index 77e66ca..9402e24 100644
--- a/test_conformance/math_brute_force/unary_two_results_double.cpp
+++ b/test_conformance/math_brute_force/unary_two_results_double.cpp
@@ -429,84 +429,11 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input array
- double *p = (double *)gIn;
- for (j = 0; j < bufferSize / sizeof(double); j++)
- p[j] = DoubleFromUInt32(genrand_int32(d));
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeValues[j] * sizeof(cl_double);
- size_t localCount = (bufferSize + vectorSize - 1) / vectorSize;
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gOutBuffer2[j]),
- &gOutBuffer2[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t{%8.2f, %8.2f} @ {%a, %a}", maxError0, maxError1, maxErrorVal0,
maxErrorVal1);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/unary_two_results_float.cpp b/test_conformance/math_brute_force/unary_two_results_float.cpp
index 70b24c7..d3fbf20 100644
--- a/test_conformance/math_brute_force/unary_two_results_float.cpp
+++ b/test_conformance/math_brute_force/unary_two_results_float.cpp
@@ -564,85 +564,11 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input array
- uint32_t *p = (uint32_t *)gIn;
- for (j = 0; j < bufferSize / sizeof(float); j++)
- p[j] = genrand_int32(d);
-
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeValues[j] * sizeof(cl_float);
- size_t localCount = (bufferSize + vectorSize - 1) / vectorSize;
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gOutBuffer2[j]),
- &gOutBuffer2[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t{%8.2f, %8.2f} @ {%a, %a}", maxError0, maxError1, maxErrorVal0,
maxErrorVal1);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/unary_two_results_i_double.cpp b/test_conformance/math_brute_force/unary_two_results_i_double.cpp
index 150e5d8..66bf414 100644
--- a/test_conformance/math_brute_force/unary_two_results_i_double.cpp
+++ b/test_conformance/math_brute_force/unary_two_results_i_double.cpp
@@ -401,84 +401,10 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input array
- double *p = (double *)gIn;
-
- for (j = 0; j < bufferSize / sizeof(double); j++)
- p[j] = DoubleFromUInt32(genrand_int32(d));
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeValues[j] * sizeof(cl_double);
- size_t localCount = (bufferSize + vectorSize - 1) / vectorSize;
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gOutBuffer2[j]),
- &gOutBuffer2[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t{%8.2f, %lld} @ %a", maxError, maxError2, maxErrorVal);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/unary_two_results_i_float.cpp b/test_conformance/math_brute_force/unary_two_results_i_float.cpp
index 1312be0..f6647b2 100644
--- a/test_conformance/math_brute_force/unary_two_results_i_float.cpp
+++ b/test_conformance/math_brute_force/unary_two_results_i_float.cpp
@@ -400,83 +400,10 @@
vlog("Wimp pass");
else
vlog("passed");
- }
- if (gMeasureTimes)
- {
- // Init input array
- uint32_t *p = (uint32_t *)gIn;
- for (j = 0; j < bufferSize / sizeof(float); j++)
- p[j] = genrand_int32(d);
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeValues[j] * sizeof(cl_float);
- size_t localCount = (bufferSize + vectorSize - 1) / vectorSize;
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gOutBuffer2[j]),
- &gOutBuffer2[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 2, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting)
vlog("\t{%8.2f, %lld} @ %a", maxError, maxError2, maxErrorVal);
+ }
+
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/unary_u_double.cpp b/test_conformance/math_brute_force/unary_u_double.cpp
index d104fc8..c703457 100644
--- a/test_conformance/math_brute_force/unary_u_double.cpp
+++ b/test_conformance/math_brute_force/unary_u_double.cpp
@@ -299,76 +299,10 @@
vlog("Wimp pass");
else
vlog("passed");
+
+ vlog("\t%8.2f @ %a", maxError, maxErrorVal);
}
- if (gMeasureTimes)
- {
- // Init input array
- double *p = (double *)gIn;
-
- for (j = 0; j < bufferSize / sizeof(double); j++) p[j] = random64(d);
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeValues[j] * sizeof(cl_double);
- size_t localCount = (bufferSize + vectorSize - 1) / vectorSize;
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(double));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sD%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting) vlog("\t%8.2f @ %a", maxError, maxErrorVal);
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/unary_u_float.cpp b/test_conformance/math_brute_force/unary_u_float.cpp
index 4b453c8..86389e4 100644
--- a/test_conformance/math_brute_force/unary_u_float.cpp
+++ b/test_conformance/math_brute_force/unary_u_float.cpp
@@ -332,84 +332,10 @@
vlog("Wimp pass");
else
vlog("passed");
+
+ vlog("\t%8.2f @ %a", maxError, maxErrorVal);
}
- if (gMeasureTimes)
- {
- // Init input array
- uint32_t *p = (uint32_t *)gIn;
- if (strstr(f->name, "exp") || strstr(f->name, "sin")
- || strstr(f->name, "cos") || strstr(f->name, "tan"))
- for (j = 0; j < bufferSize / sizeof(float); j++)
- ((float *)p)[j] = (float)genrand_real1(d);
- else if (strstr(f->name, "log"))
- for (j = 0; j < bufferSize / sizeof(float); j++)
- p[j] = genrand_int32(d) & 0x7fffffff;
- else
- for (j = 0; j < bufferSize / sizeof(float); j++)
- p[j] = genrand_int32(d);
- if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
- bufferSize, gIn, 0, NULL, NULL)))
- {
- vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
- return error;
- }
-
-
- // Run the kernels
- for (j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
- {
- size_t vectorSize = sizeValues[j] * sizeof(cl_float);
- size_t localCount = (bufferSize + vectorSize - 1) / vectorSize;
- if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
- &gOutBuffer[j])))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
- if ((error = clSetKernelArg(kernels[j], 1, sizeof(gInBuffer),
- &gInBuffer)))
- {
- LogBuildError(programs[j]);
- goto exit;
- }
-
- double sum = 0.0;
- double bestTime = INFINITY;
- for (k = 0; k < PERF_LOOP_COUNT; k++)
- {
- uint64_t startTime = GetTime();
- if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
- &localCount, NULL, 0, NULL,
- NULL)))
- {
- vlog_error("FAILED -- could not execute kernel\n");
- goto exit;
- }
-
- // Make sure OpenCL is done
- if ((error = clFinish(gQueue)))
- {
- vlog_error("Error %d at clFinish\n", error);
- goto exit;
- }
-
- uint64_t endTime = GetTime();
- double time = SubtractTime(endTime, startTime);
- sum += time;
- if (time < bestTime) bestTime = time;
- }
-
- if (gReportAverageTimes) bestTime = sum / PERF_LOOP_COUNT;
- double clocksPerOp = bestTime * (double)gDeviceFrequency
- * gComputeDevices * gSimdSize * 1e6
- / (bufferSize / sizeof(float));
- vlog_perf(clocksPerOp, LOWER_IS_BETTER, "clocks / element", "%sf%s",
- f->name, sizeNames[j]);
- }
- }
-
- if (!gSkipCorrectnessTesting) vlog("\t%8.2f @ %a", maxError, maxErrorVal);
vlog("\n");
exit:
diff --git a/test_conformance/math_brute_force/utility.h b/test_conformance/math_brute_force/utility.h
index 894a067..9ab915c 100644
--- a/test_conformance/math_brute_force/utility.h
+++ b/test_conformance/math_brute_force/utility.h
@@ -59,8 +59,6 @@
extern uint32_t gComputeDevices;
extern uint32_t gSimdSize;
extern int gSkipCorrectnessTesting;
-extern int gMeasureTimes;
-extern int gReportAverageTimes;
extern int gForceFTZ;
extern int gFastRelaxedDerived;
extern int gWimpyMode;
@@ -91,8 +89,6 @@
float Ulp_Error(float test, double reference);
float Bruteforce_Ulp_Error_Double(double test, long double reference);
-uint64_t GetTime(void);
-double SubtractTime(uint64_t endTime, uint64_t startTime);
int MakeKernel(const char **c, cl_uint count, const char *name, cl_kernel *k,
cl_program *p, bool relaxedMode);
int MakeKernels(const char **c, cl_uint count, const char *name,
@@ -123,8 +119,6 @@
void _LogBuildError(cl_program p, int line, const char *file);
#define LogBuildError(program) _LogBuildError(program, __LINE__, __FILE__)
-#define PERF_LOOP_COUNT 100
-
// The spec is fairly clear that we may enforce a hard cutoff to prevent
// premature flushing to zero.
// However, to avoid conflict for 1.0, we are letting results at TYPE_MIN +
