X-Git-Url: http://review.tizen.org/git/?a=blobdiff_plain;f=documentation%2Farithmetic__op_8cl_source.xhtml;h=c10ccc6c9b59fd2e2d421e0bb9b0bdb2ed37e56f;hb=b3a371bc429d2ba45e56baaf239d8200c2662a74;hp=7039e2a20a5e06837cd8dfeed562f9d786fbc694;hpb=46d5927c3461ec270de8b0fc087ac5dc2431488f;p=platform%2Fupstream%2Farmcl.git diff --git a/documentation/arithmetic__op_8cl_source.xhtml b/documentation/arithmetic__op_8cl_source.xhtml index 7039e2a..c10ccc6 100644 --- a/documentation/arithmetic__op_8cl_source.xhtml +++ b/documentation/arithmetic__op_8cl_source.xhtml @@ -6,7 +6,7 @@ -ARM Compute Library: src/core/CL/cl_kernels/arithmetic_op.cl Source File +Compute Library: src/core/CL/cl_kernels/arithmetic_op.cl Source File @@ -39,8 +39,8 @@ -
ARM Compute Library -  17.05 +
Compute Library +  18.05
@@ -55,6 +55,7 @@ var searchBox = new SearchBox("searchBox", "search",false,'Search');
-Go to the documentation of this file.
1 /*
2  * Copyright (c) 2016, 2017 ARM Limited.
3  *
4  * SPDX-License-Identifier: MIT
5  *
6  * Permission is hereby granted, free of charge, to any person obtaining a copy
7  * of this software and associated documentation files (the "Software"), to
8  * deal in the Software without restriction, including without limitation the
9  * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10  * sell copies of the Software, and to permit persons to whom the Software is
11  * furnished to do so, subject to the following conditions:
12  *
13  * The above copyright notice and this permission notice shall be included in all
14  * copies or substantial portions of the Software.
15  *
16  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22  * SOFTWARE.
23  */
24 #include "helpers.h"
25 
26 #ifdef SATURATE
27 #define ADD(x, y) add_sat((x), (y))
28 #define SUB(x, y) sub_sat((x), (y))
29 #else
30 #define ADD(x, y) (x) + (y)
31 #define SUB(x, y) (x) - (y)
32 #endif
33 
59 __kernel void arithmetic_add(
60  IMAGE_DECLARATION(in1),
61  IMAGE_DECLARATION(in2),
62  IMAGE_DECLARATION(out))
63 {
64  // Get pixels pointer
65  Image in1 = CONVERT_TO_IMAGE_STRUCT(in1);
66  Image in2 = CONVERT_TO_IMAGE_STRUCT(in2);
67  Image out = CONVERT_TO_IMAGE_STRUCT(out);
68 
69  // Load values
70  VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
71  in_a = CONVERT(vload16(0, (__global DATA_TYPE_IN1 *)in1.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
72  VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
73  in_b = CONVERT(vload16(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
74 
75  // Calculate and store result
76  vstore16(ADD(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr);
77 }
78 
104 __kernel void arithmetic_sub(
105  IMAGE_DECLARATION(in1),
106  IMAGE_DECLARATION(in2),
107  IMAGE_DECLARATION(out))
108 {
109  // Get pixels pointer
110  Image in1 = CONVERT_TO_IMAGE_STRUCT(in1);
111  Image in2 = CONVERT_TO_IMAGE_STRUCT(in2);
112  Image out = CONVERT_TO_IMAGE_STRUCT(out);
113 
114  // Load values
115  VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
116  in_a = CONVERT(vload16(0, (__global DATA_TYPE_IN1 *)in1.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
117  VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
118  in_b = CONVERT(vload16(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
119 
120  // Calculate and store result
121  vstore16(SUB(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr);
122 }
IMAGE_DECLARATION
#define IMAGE_DECLARATION(name)
Definition: helpers.h:49
+Go to the documentation of this file.
1 /*
2  * Copyright (c) 2016, 2017 ARM Limited.
3  *
4  * SPDX-License-Identifier: MIT
5  *
6  * Permission is hereby granted, free of charge, to any person obtaining a copy
7  * of this software and associated documentation files (the "Software"), to
8  * deal in the Software without restriction, including without limitation the
9  * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10  * sell copies of the Software, and to permit persons to whom the Software is
11  * furnished to do so, subject to the following conditions:
12  *
13  * The above copyright notice and this permission notice shall be included in all
14  * copies or substantial portions of the Software.
15  *
16  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22  * SOFTWARE.
23  */
24 #include "helpers.h"
25 
26 #if defined(FIXED_POINT_POSITION)
27 #include "fixed_point.h"
28 #endif /* FIXED_POINT_POSITION */
29 
30 #ifdef SATURATE
31 #define ADD(x, y) add_sat((x), (y))
32 #define SUB(x, y) sub_sat((x), (y))
33 #else /* SATURATE */
34 #define ADD(x, y) (x) + (y)
35 #define SUB(x, y) (x) - (y)
36 #endif /* SATURATE */
37 
69 __kernel void arithmetic_add(
70  TENSOR3D_DECLARATION(in1),
71  TENSOR3D_DECLARATION(in2),
72  TENSOR3D_DECLARATION(out))
73 {
74  // Get pixels pointer
75  Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1);
76  Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2);
77  Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out);
78 
79  // Load values
80  VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
81  in_a = CONVERT(vload16(0, (__global DATA_TYPE_IN1 *)in1.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
82  VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
83  in_b = CONVERT(vload16(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
84 
85  // Calculate and store result
86  vstore16(ADD(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr);
87 }
88 
120 __kernel void arithmetic_sub(
121  TENSOR3D_DECLARATION(in1),
122  TENSOR3D_DECLARATION(in2),
123  TENSOR3D_DECLARATION(out))
124 {
125  // Get pixels pointer
126  Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1);
127  Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2);
128  Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out);
129 
130  // Load values
131  VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
132  in_a = CONVERT(vload16(0, (__global DATA_TYPE_IN1 *)in1.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
133  VEC_DATA_TYPE(DATA_TYPE_OUT, 16)
134  in_b = CONVERT(vload16(0, (__global DATA_TYPE_IN2 *)in2.ptr), VEC_DATA_TYPE(DATA_TYPE_OUT, 16));
135 
136  // Calculate and store result
137  vstore16(SUB(in_a, in_b), 0, (__global DATA_TYPE_OUT *)out.ptr);
138 }
CONVERT
#define CONVERT(x, type)
Definition: fixed_point.h:98
+
CONVERT_TO_TENSOR3D_STRUCT
#define CONVERT_TO_TENSOR3D_STRUCT(name)
Definition: helpers.h:119
+
Tensor3D
Structure to hold 3D tensor information.
Definition: helpers.h:151
+
arithmetic_sub
__kernel void arithmetic_sub(__global uchar *in1_ptr, uint in1_stride_x, uint in1_step_x, uint in1_stride_y, uint in1_step_y, uint in1_stride_z, uint in1_step_z, uint in1_offset_first_element_in_bytes, __global uchar *in2_ptr, uint in2_stride_x, uint in2_step_x, uint in2_stride_y, uint in2_step_y, uint in2_stride_z, uint in2_step_z, uint in2_offset_first_element_in_bytes, __global uchar *out_ptr, uint out_stride_x, uint out_step_x, uint out_stride_y, uint out_step_y, uint out_stride_z, uint out_step_z, uint out_offset_first_element_in_bytes)
This function subtracts one tensors from another.
Definition: arithmetic_op.cl:120
-
CONVERT_TO_IMAGE_STRUCT
#define CONVERT_TO_IMAGE_STRUCT(name)
Definition: helpers.h:73
-
VEC_DATA_TYPE
#define VEC_DATA_TYPE(type, size)
Definition: helpers.h:32
-
ADD
#define ADD(x, y)
Definition: arithmetic_op.cl:30
-
arithmetic_add
__kernel void arithmetic_add(__global uchar *in1_ptr, uint in1_stride_x, uint in1_step_x, uint in1_stride_y, uint in1_step_y, uint in1_offset_first_element_in_bytes, __global uchar *in2_ptr, uint in2_stride_x, uint in2_step_x, uint in2_stride_y, uint in2_step_y, uint in2_offset_first_element_in_bytes, __global uchar *out_ptr, uint out_stride_x, uint out_step_x, uint out_stride_y, uint out_step_y, uint out_offset_first_element_in_bytes)
This function add two images.
Definition: arithmetic_op.cl:59
-
Image
Structure to hold Image information.
Definition: helpers.h:95
-
Image::ptr
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:97
-
arithmetic_sub
__kernel void arithmetic_sub(__global uchar *in1_ptr, uint in1_stride_x, uint in1_step_x, uint in1_stride_y, uint in1_step_y, uint in1_offset_first_element_in_bytes, __global uchar *in2_ptr, uint in2_stride_x, uint in2_step_x, uint in2_stride_y, uint in2_step_y, uint in2_offset_first_element_in_bytes, __global uchar *out_ptr, uint out_stride_x, uint out_step_x, uint out_stride_y, uint out_step_y, uint out_offset_first_element_in_bytes)
This function subtracts one image from another.
Definition: arithmetic_op.cl:104
-
SUB
#define SUB(x, y)
Definition: arithmetic_op.cl:31
+
ADD
#define ADD(x, y)
Definition: arithmetic_op.cl:34
+
fixed_point.h
+
arithmetic_add
__kernel void arithmetic_add(__global uchar *in1_ptr, uint in1_stride_x, uint in1_step_x, uint in1_stride_y, uint in1_step_y, uint in1_stride_z, uint in1_step_z, uint in1_offset_first_element_in_bytes, __global uchar *in2_ptr, uint in2_stride_x, uint in2_step_x, uint in2_stride_y, uint in2_step_y, uint in2_stride_z, uint in2_step_z, uint in2_offset_first_element_in_bytes, __global uchar *out_ptr, uint out_stride_x, uint out_step_x, uint out_stride_y, uint out_step_y, uint out_stride_z, uint out_step_z, uint out_offset_first_element_in_bytes)
This function adds two tensors.
Definition: arithmetic_op.cl:69
+
TENSOR3D_DECLARATION
#define TENSOR3D_DECLARATION(name)
Definition: helpers.h:76
+
SUB
#define SUB(x, y)
Definition: arithmetic_op.cl:35
+
VEC_DATA_TYPE
#define VEC_DATA_TYPE(type, size)
Definition: fixed_point.h:93
DATA_TYPE_OUT
#define DATA_TYPE_OUT
Definition: convolution3x3.cl:31
-
CONVERT
#define CONVERT(x, type)
Definition: helpers.h:35
+
Tensor3D::ptr
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:153