X-Git-Url: http://review.tizen.org/git/?a=blobdiff_plain;f=documentation%2Fgaussian__pyramid_8cl_source.xhtml;h=a51db9478b6671893c0c3a918bddb11da64c7aad;hb=HEAD;hp=a1c5870a6e294a2ca9b296873523008e685adc9d;hpb=871448ee8eff790c4ccc3250008dd71170cc78b2;p=platform%2Fupstream%2Farmcl.git diff --git a/documentation/gaussian__pyramid_8cl_source.xhtml b/documentation/gaussian__pyramid_8cl_source.xhtml index a1c5870..a51db94 100644 --- a/documentation/gaussian__pyramid_8cl_source.xhtml +++ b/documentation/gaussian__pyramid_8cl_source.xhtml @@ -6,7 +6,7 @@ -ARM Compute Library: src/core/CL/cl_kernels/gaussian_pyramid.cl Source File +Compute Library: src/core/CL/cl_kernels/gaussian_pyramid.cl Source File @@ -39,8 +39,8 @@ -
ARM Compute Library -  17.03.1 +
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) 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 
43 __kernel void gaussian1x5_sub_x(
44  IMAGE_DECLARATION(src),
45  IMAGE_DECLARATION(dst))
46 {
47  Image src = CONVERT_TO_IMAGE_STRUCT(src);
48  Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
49 
50  // Load values for the convolution (20 bytes needed)
51  uchar16 temp0 = vload16(0, src.ptr);
52  uchar4 temp1 = vload4(0, src.ptr + 16);
53 
54  // Convert to USHORT8
55  ushort8 l2_data = convert_ushort8((uchar8)(temp0.s02468ACE));
56  ushort8 l1_data = convert_ushort8((uchar8)(temp0.s13579BDF));
57  ushort8 m_data = convert_ushort8((uchar8)(temp0.s2468, temp0.sACE, temp1.s0));
58  ushort8 r1_data = convert_ushort8((uchar8)(temp0.s3579, temp0.sBDF, temp1.s1));
59  ushort8 r2_data = convert_ushort8((uchar8)(temp0.s468A, temp0.sCE, temp1.s02));
60 
61  // Compute convolution along the X direction
62  ushort8 pixels = l2_data + r2_data;
63  pixels += l1_data * (ushort8)4;
64  pixels += m_data * (ushort8)6;
65  pixels += r1_data * (ushort8)4;
66 
67  // Store result
68  vstore8(pixels, 0, (__global ushort *)dst.ptr);
69 }
70 
88 __kernel void gaussian5x1_sub_y(
89  IMAGE_DECLARATION(src),
90  IMAGE_DECLARATION(dst))
91 {
92  Image src = CONVERT_TO_IMAGE_STRUCT(src);
93  Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
94 
95  // Load values
96  ushort8 u2_data = vload8(0, (__global ushort *)offset(&src, 0, -2));
97  ushort8 u1_data = vload8(0, (__global ushort *)offset(&src, 0, -1));
98  ushort8 m_data = vload8(0, (__global ushort *)src.ptr);
99  ushort8 d1_data = vload8(0, (__global ushort *)offset(&src, 0, 1));
100  ushort8 d2_data = vload8(0, (__global ushort *)offset(&src, 0, 2));
101 
102  // Compute convolution along the Y direction
103  ushort8 pixels = u2_data + d2_data;
104  pixels += u1_data * (ushort8)4;
105  pixels += m_data * (ushort8)6;
106  pixels += d1_data * (ushort8)4;
107 
108  // Scale result
109  pixels >>= (ushort8)8;
110 
111  // Store result
112  vstore8(convert_uchar8_sat(pixels), 0, dst.ptr);
113 }
IMAGE_DECLARATION
#define IMAGE_DECLARATION(name)
Definition: helpers.h:49
-
offset
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
Definition: helpers.h:201
+Go to the documentation of this file.
1 /*
2  * Copyright (c) 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 
43 __kernel void gaussian1x5_sub_x(
44  IMAGE_DECLARATION(src),
45  IMAGE_DECLARATION(dst))
46 {
47  Image src = CONVERT_TO_IMAGE_STRUCT(src);
48  Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
49 
50  // Load values for the convolution (20 bytes needed)
51  uchar16 temp0 = vload16(0, src.ptr);
52  uchar4 temp1 = vload4(0, src.ptr + 16);
53 
54  // Convert to USHORT8
55  ushort8 l2_data = convert_ushort8((uchar8)(temp0.s02468ACE));
56  ushort8 l1_data = convert_ushort8((uchar8)(temp0.s13579BDF));
57  ushort8 m_data = convert_ushort8((uchar8)(temp0.s2468, temp0.sACE, temp1.s0));
58  ushort8 r1_data = convert_ushort8((uchar8)(temp0.s3579, temp0.sBDF, temp1.s1));
59  ushort8 r2_data = convert_ushort8((uchar8)(temp0.s468A, temp0.sCE, temp1.s02));
60 
61  // Compute convolution along the X direction
62  ushort8 pixels = l2_data + r2_data;
63  pixels += l1_data * (ushort8)4;
64  pixels += m_data * (ushort8)6;
65  pixels += r1_data * (ushort8)4;
66 
67  // Store result
68  vstore8(pixels, 0, (__global ushort *)dst.ptr);
69 }
70 
88 __kernel void gaussian5x1_sub_y(
89  IMAGE_DECLARATION(src),
90  IMAGE_DECLARATION(dst))
91 {
92  Image src = CONVERT_TO_IMAGE_STRUCT(src);
93  Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
94 
95  // Load values
96  ushort8 u2_data = vload8(0, (__global ushort *)offset(&src, 0, 0));
97  ushort8 u1_data = vload8(0, (__global ushort *)offset(&src, 0, 1));
98  ushort8 m_data = vload8(0, (__global ushort *)offset(&src, 0, 2));
99  ushort8 d1_data = vload8(0, (__global ushort *)offset(&src, 0, 3));
100  ushort8 d2_data = vload8(0, (__global ushort *)offset(&src, 0, 4));
101 
102  // Compute convolution along the Y direction
103  ushort8 pixels = u2_data + d2_data;
104  pixels += u1_data * (ushort8)4;
105  pixels += m_data * (ushort8)6;
106  pixels += d1_data * (ushort8)4;
107 
108  // Scale result
109  pixels >>= (ushort8)8;
110 
111  // Store result
112  vstore8(convert_uchar8_sat(pixels), 0, dst.ptr);
113 }
IMAGE_DECLARATION
#define IMAGE_DECLARATION(name)
Definition: helpers.h:68
+
offset
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
Definition: helpers.h:303
-
CONVERT_TO_IMAGE_STRUCT
#define CONVERT_TO_IMAGE_STRUCT(name)
Definition: helpers.h:73
+
CONVERT_TO_IMAGE_STRUCT
#define CONVERT_TO_IMAGE_STRUCT(name)
Definition: helpers.h:104
gaussian5x1_sub_y
__kernel void gaussian5x1_sub_y(__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_offset_first_element_in_bytes)
Computes the Gaussian Filter 5x1 + sub-sampling along the Y direction.
Definition: gaussian_pyramid.cl:88
gaussian1x5_sub_x
__kernel void gaussian1x5_sub_x(__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_ptr, uint dst_stride_x, uint dst_step_x, uint dst_stride_y, uint dst_step_y, uint dst_offset_first_element_in_bytes)
Computes the Gaussian Filter 1x5 + sub-sampling along the X direction.
Definition: gaussian_pyramid.cl:43
-
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
+
arm_compute::test::validation::dst
CLTensor dst
Definition: Convolution.cpp:137
+
Image
Structure to hold Image information.
Definition: helpers.h:142
+
Image::ptr
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:144
+
arm_compute::test::validation::src
convolution configure & src
Definition: Convolution.cpp:147