From b7e7a3eb65d80d6847bd522f66f876fd5f6fe938 Mon Sep 17 00:00:00 2001 From: Marco Antognini Date: Tue, 13 Apr 2021 15:58:44 +0100 Subject: [PATCH] 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 * Coalesce if-statements Signed-off-by: Marco Antognini * Keep else branch Address comments. Signed-off-by: Marco Antognini --- .../math_brute_force/binary_double.cpp | 89 +-------------- .../math_brute_force/binary_float.cpp | 89 +-------------- .../math_brute_force/binary_i_double.cpp | 90 +-------------- .../math_brute_force/binary_i_float.cpp | 89 +-------------- .../binary_operator_double.cpp | 89 +-------------- .../binary_operator_float.cpp | 89 +-------------- .../binary_two_results_i_double.cpp | 90 +-------------- .../binary_two_results_i_float.cpp | 92 +--------------- .../math_brute_force/i_unary_double.cpp | 67 ----------- .../math_brute_force/i_unary_float.cpp | 68 ------------ .../math_brute_force/macro_binary_double.cpp | 87 --------------- .../math_brute_force/macro_binary_float.cpp | 85 -------------- .../math_brute_force/macro_unary_double.cpp | 68 ------------ .../math_brute_force/macro_unary_float.cpp | 68 ------------ .../math_brute_force/mad_double.cpp | 104 +----------------- .../math_brute_force/mad_float.cpp | 102 +---------------- test_conformance/math_brute_force/main.cpp | 77 +------------ .../math_brute_force/ternary_double.cpp | 102 +---------------- .../math_brute_force/ternary_float.cpp | 102 +---------------- .../math_brute_force/unary_double.cpp | 79 +------------ .../math_brute_force/unary_float.cpp | 79 +------------ .../unary_two_results_double.cpp | 77 +------------ .../unary_two_results_float.cpp | 78 +------------ .../unary_two_results_i_double.cpp | 78 +------------ .../unary_two_results_i_float.cpp | 77 +------------ .../math_brute_force/unary_u_double.cpp | 70 +----------- .../math_brute_force/unary_u_float.cpp | 78 +------------ test_conformance/math_brute_force/utility.h | 6 - 28 files changed, 43 insertions(+), 2226 deletions(-) diff --git a/test_conformance/math_brute_force/binary_double.cpp b/test_conformance/math_brute_force/binary_double.cpp index 15a48a66..e51327df 100644 --- a/test_conformance/math_brute_force/binary_double.cpp +++ b/test_conformance/math_brute_force/binary_double.cpp @@ -436,95 +436,10 @@ int TestFunc_Double_Double_Double(const Func *f, MTdata d, bool relaxedMode) 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 fb56c642..5a37a407 100644 --- a/test_conformance/math_brute_force/binary_float.cpp +++ b/test_conformance/math_brute_force/binary_float.cpp @@ -428,95 +428,10 @@ int TestFunc_Float_Float_Float(const Func *f, MTdata d, bool relaxedMode) 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 686e32b4..c3c3eb13 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 @@ int TestFunc_Double_Double_Int(const Func *f, MTdata d, bool relaxedMode) 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 d99382f1..93cb910b 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 @@ int TestFunc_Float_Float_Int(const Func *f, MTdata d, bool relaxedMode) 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 a02e53ba..09540b03 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 @@ int TestFunc_Double_Double_Double_Operator(const Func *f, MTdata d, 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 39070cb2..a59be163 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 @@ int TestFunc_Float_Float_Float_Operator(const Func *f, MTdata d, 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 d71585e6..f548486f 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 @@ int TestFunc_DoubleI_Double_Double(const Func *f, MTdata d, bool relaxedMode) 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 b6f1f1bd..01c3ca7d 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 @@ int TestFunc_FloatI_Float_Float(const Func *f, MTdata d, bool relaxedMode) 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 c3822f36..040f8e65 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 @@ int TestFunc_Int_Double(const Func *f, MTdata d, bool relaxedMode) 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 d1fb867d..61ce9a23 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 @@ int TestFunc_Int_Float(const Func *f, MTdata d, bool relaxedMode) 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 cbcecb33..7c40a18b 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 @@ int TestMacro_Int_Double_Double(const Func *f, MTdata d, bool relaxedMode) 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 bbccbbe9..12c14dc5 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 @@ int TestMacro_Int_Float_Float(const Func *f, MTdata d, bool relaxedMode) 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 154d8ecb..4fadbfc3 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 @@ int TestMacro_Int_Double(const Func *f, MTdata d, bool relaxedMode) 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 725a8316..4e72748a 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 @@ int TestMacro_Int_Float(const Func *f, MTdata d, bool relaxedMode) 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 5e7dba98..a5f59183 100644 --- a/test_conformance/math_brute_force/mad_double.cpp +++ b/test_conformance/math_brute_force/mad_double.cpp @@ -285,111 +285,11 @@ int TestFunc_mad_Double(const Func *f, MTdata d, bool relaxedMode) 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 453e43e7..ff4e99d0 100644 --- a/test_conformance/math_brute_force/mad_float.cpp +++ b/test_conformance/math_brute_force/mad_float.cpp @@ -284,109 +284,11 @@ int TestFunc_mad_Float(const Func *f, MTdata d, bool relaxedMode) 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 0e1b40a4..c2a376f1 100644 --- a/test_conformance/math_brute_force/main.cpp +++ b/test_conformance/math_brute_force/main.cpp @@ -62,8 +62,6 @@ static int32_t gEndTestNumber = -1; 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 gCheckTininessBeforeRounding = 1; 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 @@ int main(int argc, const char *argv[]) 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 @@ static int ParseArgs(int argc, const char **argv) optionFound = 1; switch (*arg) { - case 'a': gReportAverageTimes ^= 1; break; - case 'c': gToggleCorrectlyRoundedDivideSqrt ^= 1; break; case 'd': gHasDouble ^= 1; break; @@ -927,8 +906,6 @@ static int ParseArgs(int argc, const char **argv) case 's': gStopOnError ^= 1; break; - case 't': gMeasureTimes ^= 1; break; - case 'v': gVerboseBruteForce ^= 1; break; case 'w': // wimpy mode @@ -970,7 +947,6 @@ static int ParseArgs(int argc, const char **argv) gMinVectorSizeIndex = 4; gMaxVectorSizeIndex = gMinVectorSizeIndex + 1; break; - break; default: vlog(" <-- unknown flag: %c (0x%2.2x)\n)", *arg, *arg); @@ -1053,9 +1029,8 @@ static void PrintFunctions(void) static void PrintUsage(void) { - vlog("%s [-acglstz]: \n", appName); + vlog("%s [-cglsz]: \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 @@ static void PrintUsage(void) "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 @@ float Abs_Error(float test, double reference) return fabs((float)(reference - (double)test)); } -#if defined(__APPLE__) -#include -#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 96a0e7bb..6e0327a2 100644 --- a/test_conformance/math_brute_force/ternary_double.cpp +++ b/test_conformance/math_brute_force/ternary_double.cpp @@ -723,109 +723,11 @@ int TestFunc_Double_Double_Double_Double(const Func *f, MTdata d, 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 20b3eb71..9ef59887 100644 --- a/test_conformance/math_brute_force/ternary_float.cpp +++ b/test_conformance/math_brute_force/ternary_float.cpp @@ -858,109 +858,11 @@ int TestFunc_Float_Float_Float_Float(const Func *f, MTdata d, bool relaxedMode) 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 615d0fb9..3dc1548c 100644 --- a/test_conformance/math_brute_force/unary_double.cpp +++ b/test_conformance/math_brute_force/unary_double.cpp @@ -290,85 +290,10 @@ int TestFunc_Double_Double(const Func *f, MTdata d, bool relaxedMode) 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 36665892..73365d19 100644 --- a/test_conformance/math_brute_force/unary_float.cpp +++ b/test_conformance/math_brute_force/unary_float.cpp @@ -315,85 +315,10 @@ int TestFunc_Float_Float(const Func *f, MTdata d, bool relaxedMode) 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 77e66ca5..9402e24c 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 @@ int TestFunc_Double2_Double(const Func *f, MTdata d, bool relaxedMode) 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 70b24c7a..d3fbf20b 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 @@ int TestFunc_Float2_Float(const Func *f, MTdata d, bool relaxedMode) 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 150e5d8f..66bf414e 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 @@ int TestFunc_DoubleI_Double(const Func *f, MTdata d, bool relaxedMode) 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 1312be01..f6647b22 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 @@ int TestFunc_FloatI_Float(const Func *f, MTdata d, bool relaxedMode) 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 d104fc88..c7034577 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 @@ int TestFunc_Double_ULong(const Func *f, MTdata d, bool relaxedMode) 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 4b453c87..86389e4a 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 @@ int TestFunc_Float_UInt(const Func *f, MTdata d, bool relaxedMode) 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 894a0675..9ab915c7 100644 --- a/test_conformance/math_brute_force/utility.h +++ b/test_conformance/math_brute_force/utility.h @@ -59,8 +59,6 @@ extern cl_mem gOutBuffer2[VECTOR_SIZE_COUNT]; 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 Abs_Error(float test, double reference); 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 @@ static inline double DoubleFromUInt32(uint32_t bits) 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 +