Blame - src/core/CL/cl_kernels/select.cl - ml/ComputeLibrary

2018-11-21 16:32:15 +0000

[diff] [blame]

1

/*

Michele Di Giorgio

2020-07-06 11:27:21 +0100

[diff] [blame]

2

Georgios Pinitas

2018-11-21 16:32:15 +0000

[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

* SOFTWARE.

*/

#include "helpers.h"

Giorgio Arena

2020-10-12 11:53:51 +0100

[diff] [blame]

26

#if defined(DATA_TYPE) && defined(VEC_SIZE)

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

27

/** This function perform a select operation between two tensors when condition tensor has the same rank.

28

*

29

* @attention The data_type need to be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=uchar

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

30

* @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16

31

*

32

* @param[in] c_ptr Pointer to the source tensor. Supported data types: U8

33

* @param[in] c_stride_x Stride of the source tensor in X dimension (in bytes)

34

* @param[in] c_step_x c_stride_x * number of elements along X processed per workitem(in bytes)

35

* @param[in] c_stride_y Stride of the source tensor in Y dimension (in bytes)

36

* @param[in] c_step_y c_stride_y * number of elements along Y processed per workitem(in bytes)

37

* @param[in] c_stride_z Stride of the source tensor in Z dimension (in bytes)

38

* @param[in] c_step_z c_stride_z * number of elements along Z processed per workitem(in bytes)

39

* @param[in] c_offset_first_element_in_bytes The offset of the first element in the source tensor

Michele Di Giorgio

2020-07-06 11:27:21 +0100

[diff] [blame]

40

* @param[in] x_ptr Pointer to the source tensor. Supported data types: All

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

41

* @param[in] x_stride_x Stride of the source tensor in X dimension (in bytes)

42

* @param[in] x_step_x x_stride_x * number of elements along X processed per workitem(in bytes)

43

* @param[in] x_stride_y Stride of the source tensor in Y dimension (in bytes)

44

* @param[in] x_step_y x_stride_y * number of elements along Y processed per workitem(in bytes)

45

* @param[in] x_stride_z Stride of the source tensor in Z dimension (in bytes)

46

* @param[in] x_step_z x_stride_z * number of elements along Z processed per workitem(in bytes)

47

* @param[in] x_offset_first_element_in_bytes The offset of the first element in the source tensor

Michele Di Giorgio

2020-07-06 11:27:21 +0100

[diff] [blame]

48

* @param[in] y_ptr Pointer to the source tensor. Supported data types: same as @p x_ptr

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

49

* @param[in] y_stride_x Stride of the source tensor in X dimension (in bytes)

50

* @param[in] y_step_x y_stride_x * number of elements along X processed per workitem(in bytes)

51

* @param[in] y_stride_y Stride of the source tensor in Y dimension (in bytes)

52

* @param[in] y_step_y y_stride_y * number of elements along Y processed per workitem(in bytes)

53

* @param[in] y_stride_z Stride of the source tensor in Z dimension (in bytes)

54

* @param[in] y_step_z y_stride_z * number of elements along Z processed per workitem(in bytes)

55

* @param[in] y_offset_first_element_in_bytes The offset of the first element in the source tensor

Michele Di Giorgio

2020-07-06 11:27:21 +0100

[diff] [blame]

56

* @param[out] out_ptr Pointer to the destination tensor. Supported data types: same as @p x_ptr

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

57

* @param[in] out_stride_x Stride of the destination tensor in X dimension (in bytes)

58

* @param[in] out_step_x out_stride_x * number of elements along X processed per workitem(in bytes)

59

* @param[in] out_stride_y Stride of the destination tensor in Y dimension (in bytes)

60

* @param[in] out_step_y out_stride_y * number of elements along Y processed per workitem(in bytes)

61

* @param[in] out_stride_z Stride of the source tensor in Z dimension (in bytes)

62

* @param[in] out_step_z out_stride_z * number of elements along Z processed per workitem(in bytes)

63

* @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination tensor

64

*/

65

__kernel void select_same_rank(

66

TENSOR3D_DECLARATION(c),

67

TENSOR3D_DECLARATION(x),

68

TENSOR3D_DECLARATION(y),

69

TENSOR3D_DECLARATION(out))

70

{

71

// Get pixels pointer

72

Tensor3D c_t = CONVERT_TO_TENSOR3D_STRUCT(c);

73

Tensor3D x_t = CONVERT_TO_TENSOR3D_STRUCT(x);

74

Tensor3D y_t = CONVERT_TO_TENSOR3D_STRUCT(y);

75

Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);

76

77

// Load values

Giorgio Arena

2020-10-12 11:53:51 +0100

[diff] [blame]

78

SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE)

79

in_c = CONVERT((VLOAD(VEC_SIZE)(0, (__global uchar *)c_t.ptr)), SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE));

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

80

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

81

in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_t.ptr);

82

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

83

in_y = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)y_t.ptr);

84

85

// Calculate and store result

86

VSTORE(VEC_SIZE)

Giorgio Arena

2020-10-12 11:53:51 +0100

[diff] [blame]

87

(select(in_y, in_x, in_c > (SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

88

}

89

90

/** This function perform a select operation between two tensors when condition tensor has a different rank.

91

*

92

* @attention The data_type need to be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=uchar

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

93

* @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16

94

*

95

* @param[in] c_ptr Pointer to the source tensor. Supported data types: U8

96

* @param[in] c_stride_x Stride of the source tensor in X dimension (in bytes)

97

* @param[in] c_step_x c_stride_x * number of elements along X processed per workitem(in bytes)

98

* @param[in] c_offset_first_element_in_bytes The offset of the first element in the source tensor

Michele Di Giorgio

2020-07-06 11:27:21 +0100

[diff] [blame]

99

* @param[in] x_ptr Pointer to the source tensor. Supported data types: All

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

100

* @param[in] x_stride_x Stride of the source tensor in X dimension (in bytes)

101

* @param[in] x_step_x x_stride_x * number of elements along X processed per workitem(in bytes)

102

* @param[in] x_stride_y Stride of the source tensor in Y dimension (in bytes)

103

* @param[in] x_step_y x_stride_y * number of elements along Y processed per workitem(in bytes)

104

* @param[in] x_stride_z Stride of the source tensor in Z dimension (in bytes)

105

* @param[in] x_step_z x_stride_z * number of elements along Z processed per workitem(in bytes)

106

* @param[in] x_offset_first_element_in_bytes The offset of the first element in the source tensor

Michele Di Giorgio

2020-07-06 11:27:21 +0100

[diff] [blame]

107

* @param[in] y_ptr Pointer to the source tensor. Supported data types: same as @p x_ptr

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

108

* @param[in] y_stride_x Stride of the source tensor in X dimension (in bytes)

109

* @param[in] y_step_x y_stride_x * number of elements along X processed per workitem(in bytes)

110

* @param[in] y_stride_y Stride of the source tensor in Y dimension (in bytes)

111

* @param[in] y_step_y y_stride_y * number of elements along Y processed per workitem(in bytes)

112

* @param[in] y_stride_z Stride of the source tensor in Z dimension (in bytes)

113

* @param[in] y_step_z y_stride_z * number of elements along Z processed per workitem(in bytes)

114

* @param[in] y_offset_first_element_in_bytes The offset of the first element in the source tensor

Michele Di Giorgio

2020-07-06 11:27:21 +0100

[diff] [blame]

115

* @param[out] out_ptr Pointer to the destination tensor. Supported data types: same as @p x_ptr

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

116

* @param[in] out_stride_x Stride of the destination tensor in X dimension (in bytes)

117

* @param[in] out_step_x out_stride_x * number of elements along X processed per workitem(in bytes)

118

* @param[in] out_stride_y Stride of the destination tensor in Y dimension (in bytes)

119

* @param[in] out_step_y out_stride_y * number of elements along Y processed per workitem(in bytes)

120

* @param[in] out_stride_z Stride of the source tensor in Z dimension (in bytes)

121

* @param[in] out_step_z out_stride_z * number of elements along Z processed per workitem(in bytes)

122

* @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination tensor

123

*/

124

__kernel void select_different_rank_2(

125

VECTOR_DECLARATION(c),

126

TENSOR3D_DECLARATION(x),

127

TENSOR3D_DECLARATION(y),

128

TENSOR3D_DECLARATION(out))

129

{

130

const int c_idx = get_global_id(1);

131

132

// Get pixels pointer

133

Vector c_t = CONVERT_TO_VECTOR_STRUCT_NO_STEP(c);

134

Tensor3D x_t = CONVERT_TO_TENSOR3D_STRUCT(x);

135

Tensor3D y_t = CONVERT_TO_TENSOR3D_STRUCT(y);

136

Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);

137

138

// Load values

Giorgio Arena

2020-10-12 11:53:51 +0100

[diff] [blame]

139

SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE)

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

140

in_c = *((__global uchar *)(c_t.ptr + c_idx * c_t.stride_x));

141

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

142

in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_t.ptr);

143

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

144

in_y = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)y_t.ptr);

145

146

// Calculate and store result

147

VSTORE(VEC_SIZE)

Giorgio Arena

2020-10-12 11:53:51 +0100

[diff] [blame]

148

(select(in_y, in_x, in_c > (SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

149

}

150

#endif /* defined(DATA_TYPE) && defined(SELECT_DATA_TYPE) && defined(VEC_SIZE) */

151

152

#if defined(DATA_TYPE) && defined(SELECT_DATA_TYPE) && defined(VEC_SIZE) && defined(DEPTH_SIZE)

153

/** This function perform a select operation between two tensors when condition tensor has a different rank.

154

*

155

* @attention The data_type need to be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=uchar

156

* @attention The select operation data_type need to be passed at compile time using -DSELECT_DATA_TYPE: e.g. -DSELECT_DATA_TYPE=uchar

157

* @attention Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16

158

*

159

* @param[in] c_ptr Pointer to the source tensor. Supported data types: U8

160

* @param[in] c_stride_x Stride of the source tensor in X dimension (in bytes)

161

* @param[in] c_step_x c_stride_x * number of elements along X processed per workitem(in bytes)

162

* @param[in] c_offset_first_element_in_bytes The offset of the first element in the source tensor

Michele Di Giorgio

2020-07-06 11:27:21 +0100

[diff] [blame]

163

* @param[in] x_ptr Pointer to the source tensor. Supported data types: All

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

164

* @param[in] x_stride_x Stride of the source tensor in X dimension (in bytes)

165

* @param[in] x_step_x x_stride_x * number of elements along X processed per workitem(in bytes)

166

* @param[in] x_stride_y Stride of the source tensor in Y dimension (in bytes)

167

* @param[in] x_step_y x_stride_y * number of elements along Y processed per workitem(in bytes)

168

* @param[in] x_stride_z Stride of the source tensor in Z dimension (in bytes)

169

* @param[in] x_step_z x_stride_z * number of elements along Z processed per workitem(in bytes)

170

* @param[in] x_offset_first_element_in_bytes The offset of the first element in the source tensor

Michele Di Giorgio

2020-07-06 11:27:21 +0100

[diff] [blame]

171

* @param[in] y_ptr Pointer to the source tensor. Supported data types: same as @p x_ptr

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

172

* @param[in] y_stride_x Stride of the source tensor in X dimension (in bytes)

173

* @param[in] y_step_x y_stride_x * number of elements along X processed per workitem(in bytes)

174

* @param[in] y_stride_y Stride of the source tensor in Y dimension (in bytes)

175

* @param[in] y_step_y y_stride_y * number of elements along Y processed per workitem(in bytes)

176

* @param[in] y_stride_z Stride of the source tensor in Z dimension (in bytes)

177

* @param[in] y_step_z y_stride_z * number of elements along Z processed per workitem(in bytes)

178

* @param[in] y_offset_first_element_in_bytes The offset of the first element in the source tensor

Michele Di Giorgio

2020-07-06 11:27:21 +0100

[diff] [blame]

179

* @param[out] out_ptr Pointer to the destination tensor. Supported data types: same as @p x_ptr

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

180

* @param[in] out_stride_x Stride of the destination tensor in X dimension (in bytes)

181

* @param[in] out_step_x out_stride_x * number of elements along X processed per workitem(in bytes)

182

* @param[in] out_stride_y Stride of the destination tensor in Y dimension (in bytes)

183

* @param[in] out_step_y out_stride_y * number of elements along Y processed per workitem(in bytes)

184

* @param[in] out_stride_z Stride of the source tensor in Z dimension (in bytes)

185

* @param[in] out_step_z out_stride_z * number of elements along Z processed per workitem(in bytes)

186

* @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination tensor

187

*/

188

__kernel void select_different_rank_n(

189

VECTOR_DECLARATION(c),

190

TENSOR3D_DECLARATION(x),

191

TENSOR3D_DECLARATION(y),

192

TENSOR3D_DECLARATION(out))

193

{

194

const int c_idx = get_global_id(2) / DEPTH_SIZE;

195

196

// Get pixels pointer

197

Vector c_t = CONVERT_TO_VECTOR_STRUCT_NO_STEP(c);

198

Tensor3D x_t = CONVERT_TO_TENSOR3D_STRUCT(x);

199

Tensor3D y_t = CONVERT_TO_TENSOR3D_STRUCT(y);

200

Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);

201

202

// Load values

Giorgio Arena

2020-10-12 11:53:51 +0100

[diff] [blame]

203

SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE)

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

204

in_c = *((__global uchar *)(c_t.ptr + c_idx * c_t.stride_x));

205

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

206

in_x = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)x_t.ptr);

207

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

208

in_y = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)y_t.ptr);

209

210

// Calculate and store result

211

VSTORE(VEC_SIZE)

Giorgio Arena

2020-10-12 11:53:51 +0100

[diff] [blame]

212

(select(in_y, in_x, in_c > (SELECT_DATA_TYPE(DATA_TYPE, VEC_SIZE))0), 0, (__global DATA_TYPE *)out_t.ptr);

Georgios Pinitas

2018-11-21 16:32:15 +0000

[diff] [blame]

213

}

Giorgio Arena