Gitiles
Code Review Sign In
review.mlplatform.org / ml / ComputeLibrary / 7333e1f10f5da9dc67b511d326121a843771a107 / . / src / core / CL / cl_kernels / l2_normalize.cl
blob: 14b37e325766f7e916d12f6bdb325ef685abd5cf [file] [log] [blame]
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]1/*
Michele Di Giorgiod9eaf612020-07-08 11:12:57 +0100[diff] [blame]2 * Copyright (c) 2016-2019 Arm Limited.
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]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
Michalis Spyrou5538d342018-11-14 08:10:13 +0000[diff] [blame]26/** This kernel performs l2 normalization on x-axis
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]27 *
28 * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
29 * @note The data size must be passed at compile time using -DDATA_SIZE e.g. -DDATA_SIZE=32
30 *
Vidhya Sudhan Loganathan7485d5a2018-07-04 09:34:00 +0100[diff] [blame]31 * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]32 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes)
33 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
Usama Arifae0001e2019-03-26 13:44:01 +0000[diff] [blame]34 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes)
35 * @param[in] src_step_y src_stride_y * number of elements along X processed per workitem(in bytes)
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]36 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor
Vidhya Sudhan Loganathan7485d5a2018-07-04 09:34:00 +0100[diff] [blame]37 * @param[in] sum_ptr Pointer to the source tensor. Supported data types: F16/F32
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]38 * @param[in] sum_stride_x Stride of the source tensor in X dimension (in bytes)
39 * @param[in] sum_step_x sum_stride_x * number of elements along X processed per workitem(in bytes)
Usama Arifae0001e2019-03-26 13:44:01 +0000[diff] [blame]40 * @param[in] sum_stride_y Stride of the source tensor in Y dimension (in bytes)
41 * @param[in] sum_step_y sum_stride_y * number of elements along Y processed per workitem(in bytes)
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]42 * @param[in] sum_offset_first_element_in_bytes The offset of the first element in the source tensor
43 * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr
44 * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes)
45 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
Usama Arifae0001e2019-03-26 13:44:01 +0000[diff] [blame]46 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes)
47 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]48 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
49 * @param[in] epsilon Epsilon value
50 */
Michalis Spyrou5538d342018-11-14 08:10:13 +0000[diff] [blame]51__kernel void l2_normalize_x(
Usama Arifae0001e2019-03-26 13:44:01 +0000[diff] [blame]52 IMAGE_DECLARATION(src),
53 IMAGE_DECLARATION(sum),
54 IMAGE_DECLARATION(dst),
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]55 DATA_TYPE epsilon)
56{
Usama Arifae0001e2019-03-26 13:44:01 +0000[diff] [blame]57 Image src = CONVERT_TO_IMAGE_STRUCT(src);
58 Image sum = CONVERT_TO_IMAGE_STRUCT(sum);
59 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]60
61 VEC_DATA_TYPE(DATA_TYPE, 16)
62 in = vload16(0, (__global DATA_TYPE *)src.ptr);
63 VEC_DATA_TYPE(DATA_TYPE, 16)
Michalis Spyrou8aaf93e2018-10-11 17:33:32 +0100[diff] [blame]64 normalize_value = (VEC_DATA_TYPE(DATA_TYPE, 16))rsqrt(fmax(((__global DATA_TYPE *)sum.ptr)[0], epsilon));
65
66 vstore16(in * normalize_value, 0, (__global DATA_TYPE *)dst.ptr);
67}
68
Michalis Spyrou5538d342018-11-14 08:10:13 +0000[diff] [blame]69/** This kernel performs l2 normalization on y-axis.
Michalis Spyrou8aaf93e2018-10-11 17:33:32 +0100[diff] [blame]70 *
71 * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
72 * @note The data size must be passed at compile time using -DDATA_SIZE e.g. -DDATA_SIZE=32
73 *
74 * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32
75 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes)
76 * @param[in] src_step_x src_stride_x * number of elements along Y processed per workitem(in bytes)
77 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes)
78 * @param[in] src_step_y src_stride_y * number of elements along X processed per workitem(in bytes)
79 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor
80 * @param[in] sum_ptr Pointer to the source tensor. Supported data types: F16/F32
81 * @param[in] sum_stride_x Stride of the source tensor in X dimension (in bytes)
82 * @param[in] sum_step_x sum_stride_x * number of elements along X processed per workitem(in bytes)
83 * @param[in] sum_stride_y Stride of the source tensor in Y dimension (in bytes)
84 * @param[in] sum_step_y sum_stride_y * number of elements along Y processed per workitem(in bytes)
85 * @param[in] sum_offset_first_element_in_bytes The offset of the first element in the source tensor
86 * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr
87 * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes)
88 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
89 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes)
90 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)
91 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
92 * @param[in] epsilon Epsilon value
93 */
Michalis Spyrou5538d342018-11-14 08:10:13 +0000[diff] [blame]94__kernel void l2_normalize_y(
Michalis Spyrou8aaf93e2018-10-11 17:33:32 +0100[diff] [blame]95 IMAGE_DECLARATION(src),
96 IMAGE_DECLARATION(sum),
97 IMAGE_DECLARATION(dst),
98 DATA_TYPE epsilon)
99{
100 Image src = CONVERT_TO_IMAGE_STRUCT(src);
101 Image sum = CONVERT_TO_IMAGE_STRUCT(sum);
102 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
103
104 VEC_DATA_TYPE(DATA_TYPE, 16)
105 in = vload16(0, (__global DATA_TYPE *)src.ptr);
106 VEC_DATA_TYPE(DATA_TYPE, 16)
107 sums = vload16(0, (__global DATA_TYPE *)sum.ptr);
108
109 VEC_DATA_TYPE(DATA_TYPE, 16)
110 normalize_value = (VEC_DATA_TYPE(DATA_TYPE, 16))rsqrt(fmax(sums, epsilon));
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]111
112 vstore16(in * normalize_value, 0, (__global DATA_TYPE *)dst.ptr);
Michalis Spyrou5538d342018-11-14 08:10:13 +0000[diff] [blame]113}
114/** This kernel performs l2 normalization on z-axis.
115 *
116 * @note The data type must be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float
117 * @note The data size must be passed at compile time using -DDATA_SIZE e.g. -DDATA_SIZE=32
118 *
119 * @param[in] src_ptr Pointer to the source tensor. Supported data types: F16/F32
120 * @param[in] src_stride_x Stride of the source tensor in X dimension (in bytes)
121 * @param[in] src_step_x src_stride_x * number of elements along Y processed per workitem(in bytes)
122 * @param[in] src_stride_y Stride of the source tensor in Y dimension (in bytes)
123 * @param[in] src_step_y src_stride_y * number of elements along X processed per workitem(in bytes)
124 * @param[in] src_stride_z Stride of the source tensor in Z dimension (in bytes)
125 * @param[in] src_step_z src_stride_z * number of elements along Z processed per workitem(in bytes)
126 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source tensor
127 * @param[in] sum_ptr Pointer to the source tensor. Supported data types: F16/F32
128 * @param[in] sum_stride_x Stride of the source tensor in X dimension (in bytes)
129 * @param[in] sum_step_x sum_stride_x * number of elements along X processed per workitem(in bytes)
130 * @param[in] sum_stride_y Stride of the source tensor in Y dimension (in bytes)
131 * @param[in] sum_step_y sum_stride_y * number of elements along Y processed per workitem(in bytes)
132 * @param[in] sum_stride_z Stride of the source tensor in Z dimension (in bytes)
133 * @param[in] sum_step_z sum_stride_z * number of elements along Z processed per workitem(in bytes)
134 * @param[in] sum_offset_first_element_in_bytes The offset of the first element in the source tensor
135 * @param[out] dst_ptr Pointer to the destination tensor. Supported data types: same as @p src_ptr
136 * @param[in] dst_stride_x Stride of the destination tensor in X dimension (in bytes)
137 * @param[in] dst_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
138 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes)
139 * @param[in] dst_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes)
140 * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes)
141 * @param[in] dst_step_z dst_stride_z * number of elements along Y processed per workitem(in bytes)
142 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor
143 * @param[in] epsilon Epsilon value
144 */
145__kernel void l2_normalize_z(
146 TENSOR3D_DECLARATION(src),
147 TENSOR3D_DECLARATION(sum),
148 TENSOR3D_DECLARATION(dst),
149 DATA_TYPE epsilon)
150{
151 Tensor3D src = CONVERT_TO_TENSOR3D_STRUCT(src);
152 Tensor3D sum = CONVERT_TO_TENSOR3D_STRUCT(sum);
153 Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);
154
155 VEC_DATA_TYPE(DATA_TYPE, 16)
156 in = vload16(0, (__global DATA_TYPE *)src.ptr);
157 VEC_DATA_TYPE(DATA_TYPE, 16)
158 sums = vload16(0, (__global DATA_TYPE *)sum.ptr);
159
160 VEC_DATA_TYPE(DATA_TYPE, 16)
161 normalize_value = (VEC_DATA_TYPE(DATA_TYPE, 16))rsqrt(fmax(sums, epsilon));
162
163 vstore16(in * normalize_value, 0, (__global DATA_TYPE *)dst.ptr);
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100[diff] [blame]164}