aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorGian Marco Iodice <gianmarco.iodice@arm.com>2019-06-10 16:45:40 +0100
committerGian Marco Iodice <gianmarco.iodice@arm.com>2019-06-11 14:42:21 +0000
commit82d9dd1c750d9a9a64650239b6ed3f955555a728 (patch)
treee2b56bd2381defc64e50a8e97d5e5c5f293bbb35
parent62404a58d4b6f5ddbf0e551764b314d1b2da2478 (diff)
downloadComputeLibrary-82d9dd1c750d9a9a64650239b6ed3f955555a728.tar.gz
COMPMID-2380: Create utility functions for is_one and is_zero with float
Change-Id: If5b968e19cf830d5472395a1b43bf72a456fd331 Signed-off-by: Gian Marco Iodice <gianmarco.iodice@arm.com> Reviewed-on: https://review.mlplatform.org/c/1322 Reviewed-by: Michele Di Giorgio <michele.digiorgio@arm.com> Comments-Addressed: Arm Jenkins <bsgcomp@arm.com> Tested-by: Arm Jenkins <bsgcomp@arm.com>
-rw-r--r--arm_compute/core/utils/helpers/float_ops.h34
-rw-r--r--src/core/CL/kernels/CLGEMMMatrixMultiplyKernel.cpp3
-rw-r--r--src/core/CL/kernels/CLGEMMMatrixMultiplyNativeKernel.cpp3
-rw-r--r--src/core/CL/kernels/CLGEMMMatrixMultiplyReshapedKernel.cpp3
-rw-r--r--src/core/CL/kernels/CLGEMMMatrixMultiplyReshapedOnlyRHSKernel.cpp11
-rw-r--r--src/core/NEON/kernels/NEGEMMMatrixMultiplyKernel.cpp4
6 files changed, 40 insertions, 18 deletions
diff --git a/arm_compute/core/utils/helpers/float_ops.h b/arm_compute/core/utils/helpers/float_ops.h
index 5012c8ef84..9d164150f5 100644
--- a/arm_compute/core/utils/helpers/float_ops.h
+++ b/arm_compute/core/utils/helpers/float_ops.h
@@ -73,25 +73,43 @@ union RawFloat
*
* @param[in] a First number to compare
* @param[in] b Second number to compare
- * @param[in] max_allowed_ulps Number of allowed ULPs
+ * @param[in] max_allowed_ulps (Optional) Number of allowed ULPs
*
* @return True if number is close else false
*/
-bool is_equal_ulps(float a, float b, int max_allowed_ulps = 0)
+inline bool is_equal_ulps(float a, float b, int max_allowed_ulps = 0)
{
RawFloat ra(a);
RawFloat rb(b);
- // Early check for sign
- if(ra.sign() != rb.sign())
- {
- return (a == b);
- }
-
// Check ULP distance
const int ulps = std::abs(ra.i32 - rb.i32);
return ulps <= max_allowed_ulps;
}
+
+/** Checks if the input floating point number is 1.0f checking if the difference is within a range defined with epsilon
+ *
+ * @param[in] a Input floating point number
+ * @param[in] epsilon (Optional) Epsilon used to define the error bounds
+ *
+ * @return True if number is close to 1.0f
+ */
+inline bool is_one(float a, float epsilon = 0.00001f)
+{
+ return std::abs(1.0f - a) <= epsilon;
+}
+
+/** Checks if the input floating point number is 0.0f checking if the difference is within a range defined with epsilon
+ *
+ * @param[in] a Input floating point number
+ * @param[in] epsilon (Optional) Epsilon used to define the error bounds
+ *
+ * @return True if number is close to 0.0f
+ */
+inline bool is_zero(float a, float epsilon = 0.00001f)
+{
+ return std::abs(0.0f - a) <= epsilon;
+}
} // namespace float_ops
} // namespace helpers
} // namespace arm_compute
diff --git a/src/core/CL/kernels/CLGEMMMatrixMultiplyKernel.cpp b/src/core/CL/kernels/CLGEMMMatrixMultiplyKernel.cpp
index 2b004c23db..b3ea309c93 100644
--- a/src/core/CL/kernels/CLGEMMMatrixMultiplyKernel.cpp
+++ b/src/core/CL/kernels/CLGEMMMatrixMultiplyKernel.cpp
@@ -35,6 +35,7 @@
#include "arm_compute/core/Types.h"
#include "arm_compute/core/Utils.h"
#include "arm_compute/core/Window.h"
+#include "arm_compute/core/utils/helpers/float_ops.h"
#include "arm_compute/core/utils/misc/ShapeCalculator.h"
#include <set>
@@ -306,7 +307,7 @@ void CLGEMMMatrixMultiplyKernel::configure(const ICLTensor *input0, const ICLTen
CLBuildOptions build_opts;
// Only define ALPHA when alpha is not 1.0f. This avoids performing unnecessary multiplications.
- if(std::abs(1.0f - alpha) > 0.00001f)
+ if(!(helpers::float_ops::is_one(alpha)))
{
build_opts.add_option("-DALPHA=" + float_to_string_with_full_precision(alpha));
}
diff --git a/src/core/CL/kernels/CLGEMMMatrixMultiplyNativeKernel.cpp b/src/core/CL/kernels/CLGEMMMatrixMultiplyNativeKernel.cpp
index 966bca099a..a3de6e0853 100644
--- a/src/core/CL/kernels/CLGEMMMatrixMultiplyNativeKernel.cpp
+++ b/src/core/CL/kernels/CLGEMMMatrixMultiplyNativeKernel.cpp
@@ -35,6 +35,7 @@
#include "arm_compute/core/Utils.h"
#include "arm_compute/core/Validate.h"
#include "arm_compute/core/Window.h"
+#include "arm_compute/core/utils/helpers/float_ops.h"
#include "arm_compute/core/utils/misc/ShapeCalculator.h"
#include "support/ToolchainSupport.h"
@@ -206,7 +207,7 @@ void CLGEMMMatrixMultiplyNativeKernel::configure(const ICLTensor *input0, const
// Create build options
CLBuildOptions build_opts;
build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input0->info()->data_type()));
- build_opts.add_option_if(std::abs(1.0f - alpha) > 0.00001f, "-DALPHA=" + float_to_string_with_full_precision(alpha));
+ build_opts.add_option_if(!(helpers::float_ops::is_one(alpha)), "-DALPHA=" + float_to_string_with_full_precision(alpha));
build_opts.add_option_if(_reinterpret_input_as_3d, "-DREINTERPRET_INPUT_AS_3D");
build_opts.add_option_if(_reinterpret_output_as_3d, "-DREINTERPRET_OUTPUT_AS_3D");
build_opts.add_option_if(_reinterpret_input_as_3d || _reinterpret_output_as_3d, "-DHEIGHT_GEMM3D=" + support::cpp11::to_string(output->info()->dimension(1)));
diff --git a/src/core/CL/kernels/CLGEMMMatrixMultiplyReshapedKernel.cpp b/src/core/CL/kernels/CLGEMMMatrixMultiplyReshapedKernel.cpp
index 8969124d71..59afa47f6f 100644
--- a/src/core/CL/kernels/CLGEMMMatrixMultiplyReshapedKernel.cpp
+++ b/src/core/CL/kernels/CLGEMMMatrixMultiplyReshapedKernel.cpp
@@ -36,6 +36,7 @@
#include "arm_compute/core/Utils.h"
#include "arm_compute/core/Validate.h"
#include "arm_compute/core/Window.h"
+#include "arm_compute/core/utils/helpers/float_ops.h"
#include "arm_compute/core/utils/misc/ShapeCalculator.h"
#include "support/ToolchainSupport.h"
@@ -200,7 +201,7 @@ void CLGEMMMatrixMultiplyReshapedKernel::configure(const ICLTensor *input0, cons
// Create build options
CLBuildOptions build_opts;
build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input0->info()->data_type()));
- build_opts.add_option_if(std::abs(1.0f - alpha) > 0.00001f, "-DALPHA=" + float_to_string_with_full_precision(alpha));
+ build_opts.add_option_if(!(helpers::float_ops::is_one(alpha)), "-DALPHA=" + float_to_string_with_full_precision(alpha));
build_opts.add_option_if(_reinterpret_output_as_3d, "-DREINTERPRET_OUTPUT_AS_3D");
build_opts.add_option_if(_reinterpret_output_as_3d, "-DHEIGHT_GEMM3D=" + support::cpp11::to_string(output->info()->dimension(1)));
build_opts.add_option_if(_reinterpret_output_as_3d, "-DDEPTH_GEMM3D=" + support::cpp11::to_string(output->info()->dimension(2)));
diff --git a/src/core/CL/kernels/CLGEMMMatrixMultiplyReshapedOnlyRHSKernel.cpp b/src/core/CL/kernels/CLGEMMMatrixMultiplyReshapedOnlyRHSKernel.cpp
index f764b9d36d..99b3d20953 100644
--- a/src/core/CL/kernels/CLGEMMMatrixMultiplyReshapedOnlyRHSKernel.cpp
+++ b/src/core/CL/kernels/CLGEMMMatrixMultiplyReshapedOnlyRHSKernel.cpp
@@ -35,6 +35,7 @@
#include "arm_compute/core/Utils.h"
#include "arm_compute/core/Validate.h"
#include "arm_compute/core/Window.h"
+#include "arm_compute/core/utils/helpers/float_ops.h"
#include "arm_compute/core/utils/misc/ShapeCalculator.h"
#include "support/ToolchainSupport.h"
@@ -73,7 +74,7 @@ Status validate_arguments(const ITensorInfo *input0, const ITensorInfo *input1,
tensor_shape1.set(0, n);
tensor_shape1.set(1, k);
- if(input2 != nullptr && std::abs(0.0f - beta) > 0.00001f)
+ if(input2 != nullptr && !(helpers::float_ops::is_zero(beta)))
{
const int input2_dim0 = static_cast<int>(input2->dimension(0));
const int input2_dim1 = static_cast<int>(input2->dimension(1));
@@ -218,7 +219,7 @@ void CLGEMMMatrixMultiplyReshapedOnlyRHSKernel::configure(const ICLTensor *input
_input0 = input0;
_input1 = input1;
- _input2 = std::abs(0.0f - beta) > 0.00001f ? input2 : nullptr;
+ _input2 = helpers::float_ops::is_zero(beta) ? nullptr : input2;
_output = output;
_reinterpret_input_as_3d = gemm_info.reinterpret_input_as_3d();
_reinterpret_output_as_3d = (gemm_info.depth_output_gemm3d() != 0);
@@ -248,9 +249,9 @@ void CLGEMMMatrixMultiplyReshapedOnlyRHSKernel::configure(const ICLTensor *input
// Create build options
CLBuildOptions build_opts;
build_opts.add_option("-DDATA_TYPE=" + get_cl_type_from_data_type(input0->info()->data_type()));
- build_opts.add_option_if(std::abs(1.0f - alpha) > 0.00001f, "-DALPHA=" + float_to_string_with_full_precision(alpha));
- build_opts.add_option_if(std::abs(0.0f - beta) > 0.00001f && _input2 != nullptr, "-DBETA=" + float_to_string_with_full_precision(beta));
- build_opts.add_option_if(std::abs(1.0f - beta) < 0.00001f, "-DUNIT_BETA");
+ build_opts.add_option_if(!(helpers::float_ops::is_one(alpha)), "-DALPHA=" + float_to_string_with_full_precision(alpha));
+ build_opts.add_option_if(!(helpers::float_ops::is_zero(beta)) && _input2 != nullptr, "-DBETA=" + float_to_string_with_full_precision(beta));
+ build_opts.add_option_if(helpers::float_ops::is_one(beta), "-DUNIT_BETA");
build_opts.add_option_if(_reinterpret_input_as_3d, "-DREINTERPRET_INPUT_AS_3D");
build_opts.add_option_if(_reinterpret_output_as_3d, "-DREINTERPRET_OUTPUT_AS_3D");
build_opts.add_option_if(gemm_info.broadcast_bias(), "-DBROADCAST_BIAS");
diff --git a/src/core/NEON/kernels/NEGEMMMatrixMultiplyKernel.cpp b/src/core/NEON/kernels/NEGEMMMatrixMultiplyKernel.cpp
index a82fae7521..e3508a1234 100644
--- a/src/core/NEON/kernels/NEGEMMMatrixMultiplyKernel.cpp
+++ b/src/core/NEON/kernels/NEGEMMMatrixMultiplyKernel.cpp
@@ -35,7 +35,7 @@
#include "arm_compute/core/Utils.h"
#include "arm_compute/core/Validate.h"
#include "arm_compute/core/Window.h"
-
+#include "arm_compute/core/utils/helpers/float_ops.h"
#include "arm_compute/core/utils/misc/ShapeCalculator.h"
#include <arm_neon.h>
@@ -998,7 +998,7 @@ void NEGEMMMatrixMultiplyKernel::run(const Window &window, const ThreadInfo &inf
ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(INEKernel::window(), window);
- bool multiply_alpha = std::abs(1.0f - _alpha) > 0.00001f;
+ const bool multiply_alpha = !(helpers::float_ops::is_one(_alpha));
// Check if the output tensor is a vector. If so,the kernel runs the vector-matrix multiplication
if((_output->info()->dimension(1) == 1))