aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--src/core/CL/cl_kernels/normalization_layer.cl2
-rw-r--r--src/core/CL/cl_kernels/pixelwise_mul_float.cl2
-rw-r--r--tests/AssetsLibrary.h6
-rw-r--r--tests/benchmark_new/CL/NormalizationLayer.cpp12
-rw-r--r--tests/benchmark_new/CL/SYSTEM/AlexNet.cpp1
-rw-r--r--tests/benchmark_new/NEON/NormalizationLayer.cpp15
-rw-r--r--tests/benchmark_new/NEON/SYSTEM/AlexNet.cpp12
-rw-r--r--tests/datasets_new/AlexNetNormalizationLayerDataset.h58
-rw-r--r--tests/datasets_new/GoogLeNetNormalizationLayerDataset.h (renamed from tests/datasets_new/NormalizationLayerDataset.h)22
-rw-r--r--tests/validation/CL/PixelWiseMultiplication.cpp19
10 files changed, 119 insertions, 30 deletions
diff --git a/src/core/CL/cl_kernels/normalization_layer.cl b/src/core/CL/cl_kernels/normalization_layer.cl
index 076b0d8909..2305ae0d15 100644
--- a/src/core/CL/cl_kernels/normalization_layer.cl
+++ b/src/core/CL/cl_kernels/normalization_layer.cl
@@ -146,7 +146,7 @@ __kernel void normalization_layer_in_map_1D(TENSOR3D_DECLARATION(input),
acc_vec += vload4(0, (__global DATA_TYPE *)tensor3D_offset(&squared_in, i - current_pos, 0, 0));
}
- const float4 normalized = pow((float4)kappa + coeff * (float4)acc_vec, beta);
+ const float4 normalized = pow((float4)kappa + coeff * CONVERT(acc_vec, float4), beta);
const float4 normalized_pixel = CONVERT(vload4(0, (__global DATA_TYPE *)in.ptr), float4) / normalized;
diff --git a/src/core/CL/cl_kernels/pixelwise_mul_float.cl b/src/core/CL/cl_kernels/pixelwise_mul_float.cl
index 98127e0311..f4f36a02af 100644
--- a/src/core/CL/cl_kernels/pixelwise_mul_float.cl
+++ b/src/core/CL/cl_kernels/pixelwise_mul_float.cl
@@ -84,7 +84,7 @@ __kernel void pixelwise_mul_float(
// Perform multiplication
#ifdef DATA_TYPE_FLOAT
VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
- res = CONVERT(in1_data * in2_data * scale, VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
+ res = CONVERT(in1_data * in2_data * (DATA_TYPE_RES)scale, VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
#else /* DATA_TYPE_FLOAT */
VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
res = CONVERT_OP_FLOAT(CONVERT_OP_FLOAT((convert_float16(in1_data * in2_data) * scale), VEC_DATA_TYPE(DATA_TYPE_RES, 16), ROUND), VEC_DATA_TYPE(DATA_TYPE_OUT, 16), ROUND);
diff --git a/tests/AssetsLibrary.h b/tests/AssetsLibrary.h
index 58738f871d..3dd30e7629 100644
--- a/tests/AssetsLibrary.h
+++ b/tests/AssetsLibrary.h
@@ -473,6 +473,12 @@ void AssetsLibrary::fill_tensor_uniform(T &&tensor, std::random_device::result_t
break;
}
case DataType::F16:
+ {
+ // It doesn't make sense to check [-inf, inf], so hard code it to a big number
+ std::uniform_real_distribution<float> distribution_f16(-100.f, 100.f);
+ fill(tensor, distribution_f16, seed_offset);
+ break;
+ }
case DataType::F32:
{
// It doesn't make sense to check [-inf, inf], so hard code it to a big number
diff --git a/tests/benchmark_new/CL/NormalizationLayer.cpp b/tests/benchmark_new/CL/NormalizationLayer.cpp
index 6d5142c7d7..7ba78a5e99 100644
--- a/tests/benchmark_new/CL/NormalizationLayer.cpp
+++ b/tests/benchmark_new/CL/NormalizationLayer.cpp
@@ -30,25 +30,31 @@
#include "framework/datasets/Datasets.h"
#include "tests/CL/CLAccessor.h"
#include "tests/TypePrinter.h"
-#include "tests/datasets_new/NormalizationLayerDataset.h"
+#include "tests/datasets_new/AlexNetNormalizationLayerDataset.h"
+#include "tests/datasets_new/GoogLeNetNormalizationLayerDataset.h"
#include "tests/fixtures_new/NormalizationLayerFixture.h"
namespace arm_compute
{
namespace test
{
+namespace
+{
+const auto normalization_layer_data_types = framework::dataset::make("DataType", { DataType::F16, DataType::F32 });
+} // namespace
+
using CLNormalizationLayerFixture = NormalizationLayerFixture<CLTensor, CLNormalizationLayer, CLAccessor>;
TEST_SUITE(CL)
REGISTER_FIXTURE_DATA_TEST_CASE(AlexNetNormalizationLayer, CLNormalizationLayerFixture, framework::DatasetMode::ALL,
framework::dataset::combine(framework::dataset::combine(datasets::AlexNetNormalizationLayerDataset(),
- framework::dataset::make("DataType", { DataType::F32 })),
+ normalization_layer_data_types),
framework::dataset::make("Batches", { 1, 4, 8 })));
REGISTER_FIXTURE_DATA_TEST_CASE(GoogLeNetNormalizationLayer, CLNormalizationLayerFixture, framework::DatasetMode::ALL,
framework::dataset::combine(framework::dataset::combine(datasets::GoogLeNetNormalizationLayerDataset(),
- framework::dataset::make("DataType", DataType::F32)),
+ normalization_layer_data_types),
framework::dataset::make("Batches", { 1, 4, 8 })));
TEST_SUITE_END()
diff --git a/tests/benchmark_new/CL/SYSTEM/AlexNet.cpp b/tests/benchmark_new/CL/SYSTEM/AlexNet.cpp
index 85cd701648..e6d91d9706 100644
--- a/tests/benchmark_new/CL/SYSTEM/AlexNet.cpp
+++ b/tests/benchmark_new/CL/SYSTEM/AlexNet.cpp
@@ -36,7 +36,6 @@
#include "framework/datasets/Datasets.h"
#include "tests/CL/CLAccessor.h"
#include "tests/TypePrinter.h"
-#include "tests/datasets_new/ActivationLayerDataset.h"
#include "tests/fixtures_new/AlexNetFixture.h"
namespace arm_compute
diff --git a/tests/benchmark_new/NEON/NormalizationLayer.cpp b/tests/benchmark_new/NEON/NormalizationLayer.cpp
index c3d4969a10..71dd9c354c 100644
--- a/tests/benchmark_new/NEON/NormalizationLayer.cpp
+++ b/tests/benchmark_new/NEON/NormalizationLayer.cpp
@@ -30,25 +30,34 @@
#include "framework/datasets/Datasets.h"
#include "tests/NEON/Accessor.h"
#include "tests/TypePrinter.h"
-#include "tests/datasets_new/NormalizationLayerDataset.h"
+#include "tests/datasets_new/AlexNetNormalizationLayerDataset.h"
+#include "tests/datasets_new/GoogLeNetNormalizationLayerDataset.h"
#include "tests/fixtures_new/NormalizationLayerFixture.h"
namespace arm_compute
{
namespace test
{
+namespace
+{
+#ifdef ARM_COMPUTE_ENABLE_FP16
+const auto normalization_layer_data_types = framework::dataset::make("DataType", { DataType::F16, DataType::F32, DataType::QS8 });
+#else /* ARM_COMPUTE_ENABLE_FP16 */
+const auto normalization_layer_data_types = framework::dataset::make("DataType", { DataType::F32, DataType::QS8 });
+#endif /* ARM_COMPUTE_ENABLE_FP16 */
+} // namespace
using NENormalizationLayerFixture = NormalizationLayerFixture<Tensor, NENormalizationLayer, Accessor>;
TEST_SUITE(NEON)
REGISTER_FIXTURE_DATA_TEST_CASE(AlexNetNormalizationLayer, NENormalizationLayerFixture, framework::DatasetMode::ALL,
framework::dataset::combine(framework::dataset::combine(datasets::AlexNetNormalizationLayerDataset(),
- framework::dataset::make("DataType", { DataType::F32, DataType::QS8 })),
+ normalization_layer_data_types),
framework::dataset::make("Batches", { 1, 4, 8 })));
REGISTER_FIXTURE_DATA_TEST_CASE(GoogLeNetNormalizationLayer, NENormalizationLayerFixture, framework::DatasetMode::ALL,
framework::dataset::combine(framework::dataset::combine(datasets::GoogLeNetNormalizationLayerDataset(),
- framework::dataset::make("DataType", DataType::F32)),
+ normalization_layer_data_types),
framework::dataset::make("Batches", { 1, 4, 8 })));
TEST_SUITE_END()
diff --git a/tests/benchmark_new/NEON/SYSTEM/AlexNet.cpp b/tests/benchmark_new/NEON/SYSTEM/AlexNet.cpp
index e97919a21f..282d3e67bd 100644
--- a/tests/benchmark_new/NEON/SYSTEM/AlexNet.cpp
+++ b/tests/benchmark_new/NEON/SYSTEM/AlexNet.cpp
@@ -36,13 +36,21 @@
#include "framework/datasets/Datasets.h"
#include "tests/NEON/Accessor.h"
#include "tests/TypePrinter.h"
-#include "tests/datasets_new/ActivationLayerDataset.h"
#include "tests/fixtures_new/AlexNetFixture.h"
namespace arm_compute
{
namespace test
{
+namespace
+{
+#ifdef ARM_COMPUTE_ENABLE_FP16
+const auto alex_net_data_types = framework::dataset::make("DataType", { DataType::F16, DataType::F32, DataType::QS8 });
+#else /* ARM_COMPUTE_ENABLE_FP16 */
+const auto alex_net_data_types = framework::dataset::make("DataType", { DataType::F32, DataType::QS8 });
+#endif /* ARM_COMPUTE_ENABLE_FP16 */
+} // namespace
+
using NEAlexNetFixture = AlexNetFixture<ITensor,
Tensor,
SubTensor,
@@ -58,7 +66,7 @@ TEST_SUITE(NEON)
TEST_SUITE(SYSTEM_TEST)
REGISTER_FIXTURE_DATA_TEST_CASE(AlexNet, NEAlexNetFixture, framework::DatasetMode::ALL,
- framework::dataset::combine(framework::dataset::make("DataType", { DataType::F32, DataType::QS8 }),
+ framework::dataset::combine(alex_net_data_types,
framework::dataset::make("Batches", { 1, 4, 8 })));
TEST_SUITE_END()
diff --git a/tests/datasets_new/AlexNetNormalizationLayerDataset.h b/tests/datasets_new/AlexNetNormalizationLayerDataset.h
new file mode 100644
index 0000000000..5714c5cf27
--- /dev/null
+++ b/tests/datasets_new/AlexNetNormalizationLayerDataset.h
@@ -0,0 +1,58 @@
+/*
+ * Copyright (c) 2017 ARM Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+#ifndef ARM_COMPUTE_TEST_ALEXNET_NORMALIZATION_LAYER_DATASET
+#define ARM_COMPUTE_TEST_ALEXNET_NORMALIZATION_LAYER_DATASET
+
+#include "framework/datasets/Datasets.h"
+
+#include "tests/TypePrinter.h"
+
+#include "arm_compute/core/TensorShape.h"
+#include "arm_compute/core/Types.h"
+
+namespace arm_compute
+{
+namespace test
+{
+namespace datasets
+{
+class AlexNetNormalizationLayerDataset final : public
+ framework::dataset::CartesianProductDataset<framework::dataset::InitializerListDataset<TensorShape>, framework::dataset::SingletonDataset<NormalizationLayerInfo>>
+{
+public:
+ AlexNetNormalizationLayerDataset()
+ : CartesianProductDataset
+ {
+ framework::dataset::make("Shape", { TensorShape(55U, 55U, 96U), TensorShape(27U, 27U, 256U) }),
+ framework::dataset::make("Info", NormalizationLayerInfo(NormType::CROSS_MAP, 5, 0.0001f, 0.75f))
+ }
+ {
+ }
+ AlexNetNormalizationLayerDataset(AlexNetNormalizationLayerDataset &&) = default;
+ ~AlexNetNormalizationLayerDataset() = default;
+};
+} // namespace datasets
+} // namespace test
+} // namespace arm_compute
+#endif /* ARM_COMPUTE_TEST_ALEXNET_NORMALIZATION_LAYER_DATASET */
diff --git a/tests/datasets_new/NormalizationLayerDataset.h b/tests/datasets_new/GoogLeNetNormalizationLayerDataset.h
index 73e215be48..2c20a6b9af 100644
--- a/tests/datasets_new/NormalizationLayerDataset.h
+++ b/tests/datasets_new/GoogLeNetNormalizationLayerDataset.h
@@ -21,8 +21,8 @@
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*/
-#ifndef ARM_COMPUTE_TEST_NORMALIZATION_LAYER_DATASET
-#define ARM_COMPUTE_TEST_NORMALIZATION_LAYER_DATASET
+#ifndef ARM_COMPUTE_TEST_GOOGLENET_NORMALIZATION_LAYER_DATASET
+#define ARM_COMPUTE_TEST_GOOGLENET_NORMALIZATION_LAYER_DATASET
#include "framework/datasets/Datasets.h"
@@ -37,22 +37,6 @@ namespace test
{
namespace datasets
{
-class AlexNetNormalizationLayerDataset final : public
- framework::dataset::CartesianProductDataset<framework::dataset::InitializerListDataset<TensorShape>, framework::dataset::SingletonDataset<NormalizationLayerInfo>>
-{
-public:
- AlexNetNormalizationLayerDataset()
- : CartesianProductDataset
- {
- framework::dataset::make("Shape", { TensorShape(55U, 55U, 96U), TensorShape(27U, 27U, 256U) }),
- framework::dataset::make("Info", NormalizationLayerInfo(NormType::CROSS_MAP, 5, 0.0001f, 0.75f))
- }
- {
- }
- AlexNetNormalizationLayerDataset(AlexNetNormalizationLayerDataset &&) = default;
- ~AlexNetNormalizationLayerDataset() = default;
-};
-
class GoogLeNetNormalizationLayerDataset final : public
framework::dataset::CartesianProductDataset<framework::dataset::InitializerListDataset<TensorShape>, framework::dataset::SingletonDataset<NormalizationLayerInfo>>
{
@@ -74,4 +58,4 @@ public:
} // namespace datasets
} // namespace test
} // namespace arm_compute
-#endif /* ARM_COMPUTE_TEST_NORMALIZATION_LAYER_DATASET */
+#endif /* ARM_COMPUTE_TEST_GOOGLENET_NORMALIZATION_LAYER_DATASET */
diff --git a/tests/validation/CL/PixelWiseMultiplication.cpp b/tests/validation/CL/PixelWiseMultiplication.cpp
index a95c823aa5..375c77dedf 100644
--- a/tests/validation/CL/PixelWiseMultiplication.cpp
+++ b/tests/validation/CL/PixelWiseMultiplication.cpp
@@ -38,6 +38,7 @@ using namespace arm_compute::test::validation;
namespace
{
const float tolerance_f32 = 1.f; /**< Tolerance value for comparing reference's output against implementation's output for float input */
+const float tolerance_f16 = 1.f; /**< Tolerance value for comparing reference's output against implementation's output for float input */
/** Compute CL pixel-wise multiplication function.
*
@@ -88,6 +89,24 @@ CLTensor compute_pixel_wise_multiplication(const TensorShape &shape, DataType dt
BOOST_AUTO_TEST_SUITE(CL)
BOOST_AUTO_TEST_SUITE(PixelWiseMultiplication)
+BOOST_AUTO_TEST_SUITE(Float16)
+BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
+BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * DataType::F16 *ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP,
+ shape, dt, convert_policy, rounding_policy)
+{
+ constexpr float scale = 1.f / 255.f;
+
+ // Compute function
+ CLTensor dst = compute_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy);
+
+ // Compute reference
+ RawTensor ref_dst = Reference::compute_reference_pixel_wise_multiplication(shape, dt, dt, dt, scale, convert_policy, rounding_policy);
+
+ // Validate output
+ validate(CLAccessor(dst), ref_dst, tolerance_f16);
+}
+BOOST_AUTO_TEST_SUITE_END()
+
BOOST_AUTO_TEST_SUITE(Float)
BOOST_TEST_DECORATOR(*boost::unit_test::label("precommit"))
BOOST_DATA_TEST_CASE(RunSmall, SmallShapes() * DataType::F32 *ConvertPolicies() * RoundingPolicy::TO_NEAREST_UP,
generated by cgit v1.2.1 (git 2.23.0) at 2024-05-20 11:40:46 +0000