24.02.1
copy_tensor.cl
Go to the documentation of this file.
1
/*
2
* Copyright (c) 2018-2021 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(DATA_TYPE) && defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER)
27
/** Performs a copy of input tensor to the output tensor.
28
*
29
* @note The following variables must be passed at compile time:
30
* -# -DDATA_TYPE : Input and output datatypes.
31
* -# -DVEC_SIZE : The number of elements processed in X dimension
32
* -# -DVEC_SIZE_LEFTOVER: Leftover size in the X dimension; x_dimension % VEC_SIZE
33
*
34
* @param[in] in_ptr Pointer to the source tensor. Supported data types: All
35
* @param[in] in_stride_x Stride of the source tensor in X dimension (in bytes)
36
* @param[in] in_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
37
* @param[in] in_stride_y Stride of the source tensor in Y dimension (in bytes)
38
* @param[in] in_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
39
* @param[in] in_stride_z Stride of the source tensor in Z dimension (in bytes)
40
* @param[in] in_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
41
* @param[in] in_offset_first_element_in_bytes The offset of the first element in the source tensor
42
* @param[out] out_ptr Pointer to the destination tensor. Supported data types: same as @p in_ptr
43
* @param[in] out_stride_x Stride of the destination tensor in X dimension (in bytes)
44
* @param[in] out_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
45
* @param[in] out_stride_y Stride of the destination tensor in Y dimension (in bytes)
46
* @param[in] out_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
47
* @param[in] out_stride_z Stride of the source tensor in Z dimension (in bytes)
48
* @param[in] out_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
49
* @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination tensor
50
*/
51
__kernel
void
copy_tensor
(
52
TENSOR3D_DECLARATION
(in),
53
TENSOR3D_DECLARATION
(out))
54
{
55
Tensor3D
in =
CONVERT_TO_TENSOR3D_STRUCT
(in);
56
Tensor3D
out =
CONVERT_TO_TENSOR3D_STRUCT
(out);
57
58
// Boundary-aware access:
59
// If the there's left-over in width (VEC_SIZE_LEFTOVER > 0):
60
// Shift all accesses other than the first to avoid accessing out of bounds
61
const
int
shift = max((
int
)(get_global_id(0) *
VEC_SIZE
) - (
int
)VEC_SIZE_LEFTOVER, 0) %
VEC_SIZE
;
62
in.
ptr
-= shift * in.
stride_x
;
63
out.
ptr
-= shift * out.
stride_x
;
64
65
// Load data
66
VEC_DATA_TYPE
(DATA_TYPE,
VEC_SIZE
)
67
data0 =
VLOAD
(
VEC_SIZE
)(0, (__global DATA_TYPE *)in.
ptr
);
68
69
// Boundary-aware store
70
STORE_VECTOR_SELECT
(data, DATA_TYPE, (__global DATA_TYPE *)out.
ptr
,
VEC_SIZE
, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0);
71
}
72
#endif // defined(DATA_TYPE) && defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER)
VEC_SIZE
#define VEC_SIZE
Definition:
qlstm_layer_normalization.cl:54
CONVERT_TO_TENSOR3D_STRUCT
#define CONVERT_TO_TENSOR3D_STRUCT(name)
Definition:
helpers.h:893
VEC_DATA_TYPE
#define VEC_DATA_TYPE(type, size)
Definition:
helpers.h:764
STORE_VECTOR_SELECT
#define STORE_VECTOR_SELECT(basename, data_type, ptr, vec_size, leftover, cond)
Definition:
load_store_utility.h:594
Tensor3D
Structure to hold 3D tensor information.
Definition:
helpers.h:932
VLOAD
#define VLOAD(size)
Definition:
helpers.h:204
TENSOR3D_DECLARATION
#define TENSOR3D_DECLARATION(name)
Definition:
helpers.h:852
Tensor3D::stride_x
int stride_x
Stride of the image in X dimension (in bytes)
Definition:
helpers.h:936
Tensor3D::ptr
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition:
helpers.h:934
arm_compute::test::copy_tensor
SimpleTensor< T1 > copy_tensor(const SimpleTensor< T2 > &tensor)
Definition:
SimpleTensor.h:232
helpers.h
src
core
CL
cl_kernels
common
copy_tensor.cl
Generated on Mon Mar 18 2024 11:31:56 for Compute Library by
1.8.17