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