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

Georgios Pinitas

aaa2718

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

[diff] [blame]

1

/*

Michele Di Giorgio

f6f7876

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

[diff] [blame]

2

Georgios Pinitas

aaa2718

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"

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

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

30

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

31

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

32

*

33

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

34

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

35

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

36

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

37

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

38

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

39

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

40

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

Michele Di Giorgio

f6f7876

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

[diff] [blame]

41

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

Georgios Pinitas

aaa2718

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

[diff] [blame]

42

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

43

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

44

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

45

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

46

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

47

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

48

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

Michele Di Giorgio

f6f7876

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

[diff] [blame]

49

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

Georgios Pinitas

aaa2718

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

[diff] [blame]

50

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

51

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

52

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

53

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

54

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

55

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

56

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

Michele Di Giorgio

f6f7876

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

[diff] [blame]

57

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

Georgios Pinitas

aaa2718

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

[diff] [blame]

58

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

59

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

60

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

61

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

62

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

63

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

64

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

65

*/

66

__kernel void select_same_rank(

67

TENSOR3D_DECLARATION(c),

68

TENSOR3D_DECLARATION(x),

69

TENSOR3D_DECLARATION(y),

70

TENSOR3D_DECLARATION(out))

71

{

72

// Get pixels pointer

73

Tensor3D c_t = CONVERT_TO_TENSOR3D_STRUCT(c);

74

Tensor3D x_t = CONVERT_TO_TENSOR3D_STRUCT(x);

75

Tensor3D y_t = CONVERT_TO_TENSOR3D_STRUCT(y);

76

Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);

77

78

// Load values

79

VEC_DATA_TYPE(SELECT_DATA_TYPE, VEC_SIZE)

80

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

81

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

82

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

83

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

84

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

85

86

// Calculate and store result

87

VSTORE(VEC_SIZE)

88

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

89

}

90

91

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

92

*

93

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

94

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

95

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

96

*

97

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

98

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

99

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

100

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

Michele Di Giorgio

f6f7876

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

[diff] [blame]

101

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

Georgios Pinitas

aaa2718

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

[diff] [blame]

102

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

103

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

104

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

105

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

106

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

107

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

108

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

Michele Di Giorgio

f6f7876

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

[diff] [blame]

109

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

Georgios Pinitas

aaa2718

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

[diff] [blame]

110

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

111

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

112

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

113

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

114

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

115

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

116

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

Michele Di Giorgio

f6f7876

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

[diff] [blame]

117

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

Georgios Pinitas

aaa2718

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

[diff] [blame]

118

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

119

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

120

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

121

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

122

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

123

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

124

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

125

*/

126

__kernel void select_different_rank_2(

127

VECTOR_DECLARATION(c),

128

TENSOR3D_DECLARATION(x),

129

TENSOR3D_DECLARATION(y),

130

TENSOR3D_DECLARATION(out))

131

{

132

const int c_idx = get_global_id(1);

133

134

// Get pixels pointer

135

Vector c_t = CONVERT_TO_VECTOR_STRUCT_NO_STEP(c);

136

Tensor3D x_t = CONVERT_TO_TENSOR3D_STRUCT(x);

137

Tensor3D y_t = CONVERT_TO_TENSOR3D_STRUCT(y);

138

Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);

139

140

// Load values

141

VEC_DATA_TYPE(SELECT_DATA_TYPE, VEC_SIZE)

142

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

143

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

144

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

145

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

146

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

147

148

// Calculate and store result

149

VSTORE(VEC_SIZE)

150

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

151

}

152

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

153

154

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

155

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

156

*

157

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

158

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

159

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

160

*

161

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

162

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

163

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

164

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

Michele Di Giorgio

f6f7876

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

[diff] [blame]

165

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

Georgios Pinitas

aaa2718

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

[diff] [blame]

166

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

167

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

168

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

169

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

170

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

171

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

172

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

Michele Di Giorgio

f6f7876

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

[diff] [blame]

173

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

Georgios Pinitas

aaa2718

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

[diff] [blame]

174

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

175

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

176

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

177

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

178

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

179

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

180

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

Michele Di Giorgio

f6f7876

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

[diff] [blame]

181

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

Georgios Pinitas

aaa2718

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

[diff] [blame]

182

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

183

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

184

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

185

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

186

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

187

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

188

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

189

*/

190

__kernel void select_different_rank_n(

191

VECTOR_DECLARATION(c),

192

TENSOR3D_DECLARATION(x),

193

TENSOR3D_DECLARATION(y),

194

TENSOR3D_DECLARATION(out))

195

{

196

const int c_idx = get_global_id(2) / DEPTH_SIZE;

197

198

// Get pixels pointer

199

Vector c_t = CONVERT_TO_VECTOR_STRUCT_NO_STEP(c);

200

Tensor3D x_t = CONVERT_TO_TENSOR3D_STRUCT(x);

201

Tensor3D y_t = CONVERT_TO_TENSOR3D_STRUCT(y);

202

Tensor3D out_t = CONVERT_TO_TENSOR3D_STRUCT(out);

203

204

// Load values

205

VEC_DATA_TYPE(SELECT_DATA_TYPE, VEC_SIZE)

206

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

207

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

208

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

209

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

210

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

211

212

// Calculate and store result

213

VSTORE(VEC_SIZE)

214

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

215

}

216

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