This patch adds CL Kernels for SquaredDifference op.
Signed-off-by: prasannar <prasanna.r@samsung.com>
--- /dev/null
+/*
+ * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
+ * Copyright (c) 2016-2018 ARM Limited.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#ifndef __ARM_COMPUTE_CLSQUARED_DIFFERENCE_KERNEL_H__
+#define __ARM_COMPUTE_CLSQUARED_DIFFERENCE_KERNEL_H__
+
+#include "arm_compute/core/CL/ICLKernel.h"
+#include "arm_compute/core/Types.h"
+
+namespace arm_compute
+{
+class ICLTensor;
+
+/** OpenCL kernel to return squared difference value of two tensors (x-y)^2*/
+class CLSquaredDifferenceKernel : public ICLKernel
+{
+public:
+ /** Default constructor */
+ CLSquaredDifferenceKernel();
+ /** Prevent instances of this class from being copied (As this class contains pointers). */
+ CLSquaredDifferenceKernel(const CLSquaredDifferenceKernel &) = delete;
+ /** Prevent instances of this class from being copied (As this class contains pointers). */
+ CLSquaredDifferenceKernel &operator=(const CLSquaredDifferenceKernel &) = delete;
+ /** Allow instances of this class to be moved */
+ CLSquaredDifferenceKernel(CLSquaredDifferenceKernel &&) = default;
+ /** Allow instances of this class to be moved */
+ CLSquaredDifferenceKernel &operator=(CLSquaredDifferenceKernel &&) = default;
+ /** Initialize the kernel's input, output.
+ *
+ * @param[in] input1 Source tensor1.
+ * @param[in] input2 Source tensor2.
+ * @param[out] output Output tensor.
+ */
+ void configure(const ICLTensor *input1, const ICLTensor *input2, ICLTensor *output);
+
+ // Inherited methods overridden:
+ void run(const Window &window, cl::CommandQueue &queue) override;
+
+private:
+ const ICLTensor *_input1;
+ const ICLTensor *_input2;
+ ICLTensor *_output;
+};
+} // namespace arm_compute
+#endif /*__ARM_COMPUTE_CLSQUARED_DIFFERENCE_KERNEL_H__ */
--- /dev/null
+/*
+ * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
+ * Copyright (c) 2016-2018 ARM Limited.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#ifndef __ARM_COMPUTE_CLSQUARED_DIFFERENCE_H__
+#define __ARM_COMPUTE_CLSQUARED_DIFFERENCE_H__
+
+#include "arm_compute/core/Types.h"
+#include "arm_compute/runtime/CL/ICLSimpleFunction.h"
+
+#include <cstdint>
+
+namespace arm_compute
+{
+class ICLTensor;
+
+class CLSquaredDifference : public ICLSimpleFunction
+{
+public:
+ /** Initialise the function's source and destination.
+ *
+ * @param[in] input1 Source tensor1. Data types supported:
+ * U8/S8/QS8/QASYMM8/U16/S16/QS16/F16/U32/S32/F32.
+ * @param[in] input2 Source tensor2. Data types supported:
+ * U8/S8/QS8/QASYMM8/U16/S16/QS16/F16/U32/S32/F32.
+ * @param[out] output Output tensor. Data types supported: Same as @p input.
+ *
+ */
+ void configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output);
+};
+} // namespace arm_compute
+#endif /*__ARM_COMPUTE_CLSQUARED_DIFFERENCESQUARED_DIFFERENCE*/
{"softmax_layer_max_shift_exp_sum_quantized_parallel", "softmax_layer_quantized.cl"},
{"softmax_layer_max_shift_exp_sum_serial", "softmax_layer.cl"},
{"softmax_layer_max_shift_exp_sum_parallel", "softmax_layer.cl"},
+ {"squared_difference", "squared_difference.cl"},
{"strided_slice", "strided_slice.cl"},
{"suppress_non_maximum", "canny.cl"},
{"tablelookup_U8", "tablelookup.cl"},
#include "./cl_kernels/reduction_mean.clembed"
},
{
+ "squared_difference.cl",
+#include "./cl_kernels/squared_difference.clembed"
+ },
+ {
"strided_slice.cl",
#include "./cl_kernels/strided_slice.clembed"
},
--- /dev/null
+/*
+ * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
+ * Copyright (c) 2016-2018 ARM Limited.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#include "helpers.h"
+
+/** Returns true value of squared_difference of two tensors.
+ *
+ * @attention Data type can be passed using the -DDATA_TYPE_IN compile flag, e.g. -DDATA_TYPE_IN=float
+ * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
+ * @note Can only take floating point data types.
+ *
+ * @param[in] input1_ptr Pointer to the source image. Supported data types: F16/F32
+ * @param[in] input1_stride_x Stride of the source image in X dimension (in bytes)
+ * @param[in] input1_step_x input1_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] input1_stride_y Stride of the source image in Y dimension (in bytes)
+ * @param[in] input1_step_y input1_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in] input1_stride_z Stride of the source tensor in Z dimension (in bytes)
+ * @param[in] input1_step_z input1_stride_z * number of elements along Z processed per workitem(in bytes)
+ * @param[in] input1_offset_first_element_in_bytes The offset of the first element in the source image
+ *
+ * @param[in] input2_ptr Pointer to the source image. Supported data types: F16/F32
+ * @param[in] input2_stride_x Stride of the source image in X dimension (in bytes)
+ * @param[in] input2_step_x input2_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] input2_stride_y Stride of the source image in Y dimension (in bytes)
+ * @param[in] input2_step_y input2_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in] input2_stride_z Stride of the source tensor in Z dimension (in bytes)
+ * @param[in] input2_step_z input2_stride_z * number of elements along Z processed per workitem(in bytes)
+ * @param[in] input2_offset_first_element_in_bytes The offset of the first element in the source image
+ *
+ * @param[out] output_ptr Pointer to the destination image.
+ * @param[in] output_stride_x Stride of the destination image in X dimension (in bytes)
+ * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
+ * @param[in] output_stride_y Stride of the destination image in Y dimension (in bytes)
+ * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
+ * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
+ * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
+ * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination image
+ */
+__kernel void squared_difference(
+ TENSOR3D_DECLARATION(input1),
+ TENSOR3D_DECLARATION(input2),
+ TENSOR3D_DECLARATION(output))
+{
+ Tensor3D input1 = CONVERT_TO_TENSOR3D_STRUCT(input1);
+ Tensor3D input2 = CONVERT_TO_TENSOR3D_STRUCT(input2);
+ Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
+
+ VSTORE(VEC_SIZE)
+ (pow(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input1.ptr) - VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input2.ptr), 2),
+ 0, (__global DATA_TYPE *)output.ptr);
+}
--- /dev/null
+/*
+ * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
+ * Copyright (c) 2016-2018 ARM Limited.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#include "arm_compute/core/CL/kernels/CLSquaredDifferenceKernel.h"
+
+#include "arm_compute/core/AccessWindowStatic.h"
+#include "arm_compute/core/CL/CLHelpers.h"
+#include "arm_compute/core/CL/CLKernelLibrary.h"
+#include "arm_compute/core/CL/CLKernelLibraryEx.h"
+#include "arm_compute/core/CL/ICLTensor.h"
+#include "arm_compute/core/Error.h"
+#include "arm_compute/core/Helpers.h"
+#include "arm_compute/core/TensorInfo.h"
+#include "arm_compute/core/Validate.h"
+#include "arm_compute/core/Window.h"
+
+#include <algorithm>
+
+using namespace arm_compute;
+
+CLSquaredDifferenceKernel::CLSquaredDifferenceKernel()
+ : _input1(nullptr), _input2(nullptr), _output(nullptr)
+{
+}
+
+void CLSquaredDifferenceKernel::configure(const ICLTensor *input1, const ICLTensor *input2,
+ ICLTensor *output)
+{
+ ARM_COMPUTE_ERROR_ON_MISMATCHING_DIMENSIONS(input1->info()->tensor_shape(),
+ input2->info()->tensor_shape());
+ ARM_COMPUTE_ERROR_ON_MISMATCHING_DIMENSIONS(input1->info()->tensor_shape(),
+ output->info()->tensor_shape());
+ ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input1, input2);
+ ARM_COMPUTE_ERROR_ON_MISMATCHING_DATA_TYPES(input1, output);
+
+ _input1 = input1;
+ _input2 = input2;
+ _output = output;
+
+ constexpr unsigned int num_elems_processed_per_iteration = 16;
+
+ // Create kernel
+ std::set<std::string> build_opts;
+ build_opts.emplace(("-DDATA_TYPE=" + get_cl_type_from_data_type(input1->info()->data_type())));
+ build_opts.emplace(
+ ("-DVEC_SIZE=" + support::cpp11::to_string(num_elems_processed_per_iteration)));
+ _kernel = static_cast<cl::Kernel>(
+ CLKernelLibraryEx::get().create_kernel("squared_difference", build_opts));
+
+ // Configure window
+ Window win = calculate_max_window(*input1->info(), Steps(num_elems_processed_per_iteration));
+
+ AccessWindowHorizontal input1_access(input1->info(), 0, num_elems_processed_per_iteration);
+ AccessWindowHorizontal input2_access(input2->info(), 0, num_elems_processed_per_iteration);
+ AccessWindowHorizontal output_access(output->info(), 0, num_elems_processed_per_iteration);
+
+ ValidRegion valid_region =
+ intersect_valid_regions(input1->info()->valid_region(), input2->info()->valid_region());
+
+ update_window_and_padding(win, input1_access, input2_access, output_access);
+
+ output_access.set_valid_region(win, valid_region);
+
+ ICLKernel::configure(win);
+}
+
+void CLSquaredDifferenceKernel::run(const Window &window, cl::CommandQueue &queue)
+{
+ ARM_COMPUTE_ERROR_ON_UNCONFIGURED_KERNEL(this);
+ ARM_COMPUTE_ERROR_ON_INVALID_SUBWINDOW(ICLKernel::window(), window);
+
+ Window collapsed = window.collapse_if_possible(ICLKernel::window(), Window::DimZ);
+ Window slice = collapsed.first_slice_window_3D();
+
+ do
+ {
+ unsigned int idx = 0;
+ add_3D_tensor_argument(idx, _input1, slice);
+ add_3D_tensor_argument(idx, _input2, slice);
+ add_3D_tensor_argument(idx, _output, slice);
+ enqueue(queue, *this, slice);
+ } while (collapsed.slide_window_slice_3D(slice));
+}
--- /dev/null
+/*
+ * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
+ * Copyright (c) 2016-2018 ARM Limited.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#include "arm_compute/runtime/CL/functions/CLSquaredDifference.h"
+
+#include "arm_compute/core/CL/kernels/CLSquaredDifferenceKernel.h"
+
+#include <utility>
+
+using namespace arm_compute;
+
+void CLSquaredDifference::configure(ICLTensor *input1, ICLTensor *input2, ICLTensor *output)
+{
+ auto k = arm_compute::support::cpp14::make_unique<CLSquaredDifferenceKernel>();
+ k->configure(input1, input2, output);
+ _kernel = std::move(k);
+}
#include <arm_compute/runtime/CL/functions/CLCopy.h>
#include <arm_compute/runtime/CL/functions/CLNormalizationLayer.h>
#include <arm_compute/runtime/CL/functions/CLExp.h>
+#include <arm_compute/runtime/CL/functions/CLSquaredDifference.h>
#include <arm_compute/runtime/CL/functions/CLNeg.h>
#include <arm_compute/runtime/SubTensor.h>
#include "internal/layers/PadLayer.h"
#include "internal/layers/SimpleSpaceToDepth.h"
#include "internal/layers/SimpleEmbeddingLookup.h"
-#include "internal/layers/SquaredDifferenceOperation.h"
#include "internal/layers/SimpleDepthToSpace.h"
#include "internal/layers/HashtableLookupLayer.h"
#include "internal/layers/SimpleSpaceToBatchND.h"
auto lhs_alloc = ctx.at(::internal::tflite::operand::Index{param.lhs_index});
auto rhs_alloc = ctx.at(::internal::tflite::operand::Index{param.rhs_index});
- auto fn = nnfw::make_unique<SquaredDifferenceOperation>();
-
- // TODO Decide ConvertPolicy (WARP? SATURATE?) according to NN API specification
- fn->configure(lhs_alloc, rhs_alloc, ofm_alloc);
+ if (::internal::arm_compute::isGpuMode())
+ {
+ auto fn = nnfw::make_unique<::arm_compute::CLSquaredDifference>();
- builder.append("SquaredDifference", std::move(fn));
+ fn->configure(CAST_CL(lhs_alloc), CAST_CL(rhs_alloc), CAST_CL(ofm_alloc));
+ builder.append("SquaredDifference", std::move(fn));
+ }
+ else
+ {
+ // TODO Enable NEON Support
+ throw std::runtime_error("Not supported, yet");
+ }
};
+++ /dev/null
-/*
- * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-#include "SquaredDifferenceOperation.h"
-#include "internal/arm_compute.h"
-
-void SquaredDifferenceOperation::configure(::arm_compute::ITensor *input1,
- ::arm_compute::ITensor *input2,
- ::arm_compute::ITensor *output)
-{
- _input1 = input1;
- _input2 = input2;
- _output = output;
-
- if (::internal::arm_compute::isGpuMode())
- {
- _cl_intermediate.allocator()->init(*input1->info());
-
- _cl_sub.configure(CAST_CL(input1), CAST_CL(input2), &_cl_intermediate,
- ::arm_compute::ConvertPolicy::SATURATE);
-
- const ::arm_compute::ActivationLayerInfo act_info{
- ::arm_compute::ActivationLayerInfo::ActivationFunction::SQUARE};
-
- _cl_act.configure(&_cl_intermediate, CAST_CL(output), act_info);
-
- _cl_intermediate.allocator()->allocate();
- }
- else
- {
- _neon_intermediate.allocator()->init(*input1->info());
-
- _neon_sub.configure(input1, input2, &_neon_intermediate,
- ::arm_compute::ConvertPolicy::SATURATE);
-
- const ::arm_compute::ActivationLayerInfo act_info{
- ::arm_compute::ActivationLayerInfo::ActivationFunction::SQUARE};
-
- _neon_act.configure(&_neon_intermediate, output, act_info);
-
- _neon_intermediate.allocator()->allocate();
- }
-}
-
-void SquaredDifferenceOperation::run(void)
-{
- if (::internal::arm_compute::isGpuMode())
- {
- _cl_sub.run();
- _cl_act.run();
- }
- else
- {
- _neon_sub.run();
- _neon_act.run();
- }
-}
+++ /dev/null
-/*
- * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-#ifndef __SQUARED_DIFFERENCE_OPERATION_H__
-#define __SQUARED_DIFFERENCE_OPERATION_H__
-
-#include <arm_compute/runtime/Tensor.h>
-#include <arm_compute/runtime/CL/CLTensor.h>
-
-#include <arm_compute/runtime/CL/functions/CLArithmeticSubtraction.h>
-#include <arm_compute/runtime/CL/functions/CLActivationLayer.h>
-#include <arm_compute/runtime/NEON/functions/NEArithmeticSubtraction.h>
-#include <arm_compute/runtime/NEON/functions/NEActivationLayer.h>
-
-class SquaredDifferenceOperation : public ::arm_compute::IFunction
-{
-public:
- void configure(::arm_compute::ITensor *input1, ::arm_compute::ITensor *input2,
- ::arm_compute::ITensor *output);
-
-public:
- void run(void) override;
-
-private:
- ::arm_compute::ITensor *_input1;
- ::arm_compute::ITensor *_input2;
-
- ::arm_compute::CLTensor _cl_intermediate;
- ::arm_compute::Tensor _neon_intermediate;
-
- ::arm_compute::ITensor *_output;
-
-private:
- ::arm_compute::CLArithmeticSubtraction _cl_sub;
- ::arm_compute::CLActivationLayer _cl_act;
-
- ::arm_compute::NEArithmeticSubtraction _neon_sub;
- ::arm_compute::NEActivationLayer _neon_act;
-};
-#endif // __SQUARED_DIFFERENCE_OPERATION_H__
projects / platform / core / ml / nnfw.git / commitdiff