diff options
-rw-r--r-- | src/core/CL/cl_kernels/l2_normalize.cl | 20 | ||||
-rw-r--r-- | src/core/CL/kernels/CLL2NormalizeLayerKernel.cpp | 14 | ||||
-rw-r--r-- | tests/benchmark/CL/L2NormalizeLayer.cpp | 6 | ||||
-rw-r--r-- | tests/benchmark/NEON/L2NormalizeLayer.cpp | 6 |
4 files changed, 26 insertions, 20 deletions
diff --git a/src/core/CL/cl_kernels/l2_normalize.cl b/src/core/CL/cl_kernels/l2_normalize.cl index 5f66efbcc4..70b8b363fa 100644 --- a/src/core/CL/cl_kernels/l2_normalize.cl +++ b/src/core/CL/cl_kernels/l2_normalize.cl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2016-2018 ARM Limited. + * Copyright (c) 2016-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -31,26 +31,32 @@ * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] src_step_y src_stride_y * number of elements along X processed per workitem(in bytes) * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor * @param[in] sum_ptr Pointer to the source tensor. Supported data types: F16/F32 * @param[in] sum_stride_x Stride of the source tensor in X dimension (in bytes) * @param[in] sum_step_x sum_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] sum_stride_y Stride of the source tensor in Y dimension (in bytes) + * @param[in] sum_step_y sum_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] sum_offset_first_element_in_bytes The offset of the first element in the source tensor * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes) * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) + * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor * @param[in] epsilon Epsilon value */ __kernel void l2_normalize_x( - VECTOR_DECLARATION(src), - VECTOR_DECLARATION(sum), - VECTOR_DECLARATION(dst), + IMAGE_DECLARATION(src), + IMAGE_DECLARATION(sum), + IMAGE_DECLARATION(dst), DATA_TYPE epsilon) { - Vector src = CONVERT_TO_VECTOR_STRUCT(src); - Vector sum = CONVERT_TO_VECTOR_STRUCT(sum); - Vector dst = CONVERT_TO_VECTOR_STRUCT(dst); + Image src = CONVERT_TO_IMAGE_STRUCT(src); + Image sum = CONVERT_TO_IMAGE_STRUCT(sum); + Image dst = CONVERT_TO_IMAGE_STRUCT(dst); VEC_DATA_TYPE(DATA_TYPE, 16) in = vload16(0, (__global DATA_TYPE *)src.ptr); diff --git a/src/core/CL/kernels/CLL2NormalizeLayerKernel.cpp b/src/core/CL/kernels/CLL2NormalizeLayerKernel.cpp index e33dab0f09..cb2e29449c 100644 --- a/src/core/CL/kernels/CLL2NormalizeLayerKernel.cpp +++ b/src/core/CL/kernels/CLL2NormalizeLayerKernel.cpp @@ -117,7 +117,7 @@ void CLL2NormalizeLayerKernel::configure(const ICLTensor *input, const ICLTensor { case 0: kernel_name = "x"; - idx = num_arguments_per_1D_tensor() * 3; + idx = num_arguments_per_2D_tensor() * 3; break; case 1: kernel_name = "y"; @@ -169,17 +169,17 @@ void CLL2NormalizeLayerKernel::run(const Window &window, cl::CommandQueue &queue case 0: { window_sum.set(Window::DimX, Window::Dimension(0, 0, 0)); - Window in_slice = window.first_slice_window_1D(); - Window sum_slice = window_sum.first_slice_window_1D(); + Window in_slice = window.first_slice_window_2D(); + Window sum_slice = window_sum.first_slice_window_2D(); do { unsigned int idx = 0; - add_1D_tensor_argument(idx, _input, in_slice); - add_1D_tensor_argument(idx, _sum, sum_slice); - add_1D_tensor_argument(idx, _output, in_slice); + add_2D_tensor_argument(idx, _input, in_slice); + add_2D_tensor_argument(idx, _sum, sum_slice); + add_2D_tensor_argument(idx, _output, in_slice); enqueue(queue, *this, in_slice); } - while(window.slide_window_slice_1D(in_slice) && window.slide_window_slice_1D(sum_slice)); + while(window.slide_window_slice_2D(in_slice) && window.slide_window_slice_2D(sum_slice)); } break; case 1: diff --git a/tests/benchmark/CL/L2NormalizeLayer.cpp b/tests/benchmark/CL/L2NormalizeLayer.cpp index c88792c1ed..23bf612069 100644 --- a/tests/benchmark/CL/L2NormalizeLayer.cpp +++ b/tests/benchmark/CL/L2NormalizeLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -49,10 +49,10 @@ using CLL2NormalizeLayerFixture = L2NormalizeLayerFixture<CLTensor, CLL2Normaliz TEST_SUITE(CL) REGISTER_FIXTURE_DATA_TEST_CASE(L2NormalizeLayer, CLL2NormalizeLayerFixture, framework::DatasetMode::ALL, - framework::dataset::combine(framework::dataset::combine(datasets::SmallShapes(), data_types), framework::dataset::make("Axis", { 0 }))); + framework::dataset::combine(framework::dataset::combine(datasets::SmallShapes(), data_types), framework::dataset::make("Axis", { 0, 1, 2 }))); TEST_SUITE(NIGHTLY) REGISTER_FIXTURE_DATA_TEST_CASE(L2NormalizeLayer, CLL2NormalizeLayerFixture, framework::DatasetMode::NIGHTLY, - framework::dataset::combine(framework::dataset::combine(datasets::LargeShapes(), data_types), framework::dataset::make("Axis", { 0 }))); + framework::dataset::combine(framework::dataset::combine(datasets::LargeShapes(), data_types), framework::dataset::make("Axis", { 0, 1, 2 }))); TEST_SUITE_END() TEST_SUITE_END() } // namespace benchmark diff --git a/tests/benchmark/NEON/L2NormalizeLayer.cpp b/tests/benchmark/NEON/L2NormalizeLayer.cpp index 14a9784312..164c64b2aa 100644 --- a/tests/benchmark/NEON/L2NormalizeLayer.cpp +++ b/tests/benchmark/NEON/L2NormalizeLayer.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2017 ARM Limited. + * Copyright (c) 2017-2019 ARM Limited. * * SPDX-License-Identifier: MIT * @@ -49,10 +49,10 @@ using NEL2NormalizeLayerFixture = L2NormalizeLayerFixture<Tensor, NEL2NormalizeL TEST_SUITE(NEON) REGISTER_FIXTURE_DATA_TEST_CASE(L2NormalizeLayer, NEL2NormalizeLayerFixture, framework::DatasetMode::ALL, - framework::dataset::combine(framework::dataset::combine(datasets::SmallShapes(), data_types), framework::dataset::make("Axis", { 0 }))); + framework::dataset::combine(framework::dataset::combine(datasets::SmallShapes(), data_types), framework::dataset::make("Axis", { 0, 1, 2 }))); TEST_SUITE(NIGHTLY) REGISTER_FIXTURE_DATA_TEST_CASE(L2NormalizeLayer, NEL2NormalizeLayerFixture, framework::DatasetMode::NIGHTLY, - framework::dataset::combine(framework::dataset::combine(datasets::LargeShapes(), data_types), framework::dataset::make("Axis", { 0 }))); + framework::dataset::combine(framework::dataset::combine(datasets::LargeShapes(), data_types), framework::dataset::make("Axis", { 0, 1, 2 }))); TEST_SUITE_END() TEST_SUITE_END() } // namespace benchmark |