2 * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
4 * Licensed under the Apache License, Version 2.0 (the "License");
5 * you may not use this file except in compliance with the License.
6 * You may obtain a copy of the License at
8 * http://www.apache.org/licenses/LICENSE-2.0
10 * Unless required by applicable law or agreed to in writing, software
11 * distributed under the License is distributed on an "AS IS" BASIS,
12 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 * See the License for the specific language governing permissions and
14 * limitations under the License.
18 * Copyright (c) 2016-2018 ARM Limited.
20 * SPDX-License-Identifier: MIT
22 * Permission is hereby granted, free of charge, to any person obtaining a copy
23 * of this software and associated documentation files (the "Software"), to
24 * deal in the Software without restriction, including without limitation the
25 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
26 * sell copies of the Software, and to permit persons to whom the Software is
27 * furnished to do so, subject to the following conditions:
29 * The above copyright notice and this permission notice shall be included in all
30 * copies or substantial portions of the Software.
32 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
33 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
34 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
35 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
36 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
37 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
47 #if defined(OP_CODE) && defined(DATA_TYPE)
48 /** returns truth value of the two input tensors for BINARY LOGICAL OP.
49 * where BINARY LOGICAL OP can be AND, OR.
51 * @attention Data type can be passed using the -DDATA_TYPE compile flag, e.g. -DDATA_TYPE=uchar
52 * @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size.
54 * @attention Operation type(code) specifying which operation to perform should be passed as
55 * preprocessor argument using -DOP_CODE = number. e.g. -DOP_CODE=1
57 * @param[in] input1_ptr Pointer to the source tensor.
58 * Supported data types: QASYMM8
59 * @param[in] input1_stride_x Stride of the source tensor in X dimension
61 * @param[in] input1_step_x input1_stride_x * number of elements along X
62 * processed per workitem(in bytes)
63 * @param[in] input1_stride_y Stride of the source tensor in Y dimension
65 * @param[in] input1_step_y input1_stride_y * number of elements along Y
66 * processed per workitem(in bytes)
67 * @param[in] input1_stride_z Stride of the source tensor in Z dimension
69 * @param[in] input1_step_z input1_stride_z * number of elements along Z
70 * processed per workitem(in bytes)
71 * @param[in] input1_offset_first_element_in_bytes The offset of the first element in the source
73 * @param[in] input2_ptr Pointer to the source tensor.
74 * Supported data types: QASYMM8
75 * @param[in] input2_stride_x Stride of the source tensor in X dimension
77 * @param[in] input2_step_x input2_stride_x * number of elements along X
78 * processed per workitem(in bytes)
79 * @param[in] input2_stride_y Stride of the source tensor in Y dimension
81 * @param[in] input2_step_y input2_stride_y * number of elements along Y
82 * processed per workitem(in bytes)
83 * @param[in] input2_stride_z Stride of the source tensor in Z dimension
85 * @param[in] input2_step_z input2_stride_z * number of elements along Z
86 * processed per workitem(in bytes)
87 * @param[in] input2_offset_first_element_in_bytes The offset of the first element in the source
89 * @param[out] output_ptr Pointer to the destination tensor.
90 * Supported data types: QASYMM8
91 * @param[in] output_stride_x Stride of the destination tensor in X dimension
93 * @param[in] output_step_x output_stride_x * number of elements along X
94 * processed per workitem(in bytes)
95 * @param[in] output_stride_y Stride of the destination tensor in Y dimension
97 * @param[in] output_step_y output_stride_y * number of elements along Y
98 * processed per workitem(in bytes)
99 * @param[in] output_stride_z Stride of the destination tensor in Z dimension
101 * @param[in] output_step_z output_stride_z * number of elements along Z
102 * processed per workitem(in bytes)
104 __kernel void binary_logical_op(TENSOR3D_DECLARATION(input1), TENSOR3D_DECLARATION(input2),
105 TENSOR3D_DECLARATION(output))
107 Tensor3D input1 = CONVERT_TO_TENSOR3D_STRUCT(input1);
108 Tensor3D input2 = CONVERT_TO_TENSOR3D_STRUCT(input2);
109 Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
111 #if OP_CODE == 1 // LOGICAL AND
113 (CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input1.ptr) &&
114 VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input2.ptr),
115 VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)),
116 0, (__global DATA_TYPE *)output.ptr);
118 #elif OP_CODE == 2 // LOGICAL OR
120 (CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input1.ptr) ||
121 VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)input2.ptr),
122 VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)),
123 0, (__global DATA_TYPE *)output.ptr);
125 #else // OP NOT SUPPORTED
130 #endif // if defined(OP_CODE) && defined(DATA_TYPE)