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

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

1

/*

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

2

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +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

* SOFTWARE.

*/

#include "helpers.h"

Manuel Bottini

2020-12-02 13:22:14 +0000

[diff] [blame]

26

#if defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER)

27

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

28

/** This function computes the bitwise OR of two input images.

29

*

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

30

* @note The following variables must be passed at compile time:

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

*

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

34

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

35

* @param[in] in1_stride_x Stride of the source image in X dimension (in bytes)

36

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

37

* @param[in] in1_stride_y Stride of the source image in Y dimension (in bytes)

38

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

39

* @param[in] in1_offset_first_element_in_bytes The offset of the first element in the source image

40

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

41

* @param[in] in2_stride_x Stride of the source image in X dimension (in bytes)

42

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

43

* @param[in] in2_stride_y Stride of the source image in Y dimension (in bytes)

44

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

45

* @param[in] in2_offset_first_element_in_bytes The offset of the first element in the source image

46

* @param[out] out_ptr Pointer to the destination image. Supported data types: U8

47

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

48

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

49

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

50

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

51

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

52

*/

53

__kernel void bitwise_or(

54

IMAGE_DECLARATION(in1),

55

IMAGE_DECLARATION(in2),

56

IMAGE_DECLARATION(out))

57

{

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

58

uint x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0);

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

59

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

60

// Get pixels pointer

61

__global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x_offs + get_global_id(1) * in1_step_y;

62

__global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + x_offs + get_global_id(1) * in2_step_y;

63

__global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x_offs + get_global_id(1) * out_step_y;

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

64

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

65

// Load data

66

VEC_DATA_TYPE(uchar, VEC_SIZE)

67

in_a = VLOAD(VEC_SIZE)(0, (__global uchar *)in1_addr);

68

VEC_DATA_TYPE(uchar, VEC_SIZE)

69

in_b = VLOAD(VEC_SIZE)(0, (__global uchar *)in2_addr);

70

71

VEC_DATA_TYPE(uchar, VEC_SIZE)

72

data0 = in_a | in_b;

73

74

// Boundary-aware store

75

STORE_VECTOR_SELECT(data, uchar, (__global uchar *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0);

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

76

}

77

78

/** This function computes the bitwise AND of two input images.

79

*

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

80

* @note The following variables must be passed at compile time:

81

* -# -DVEC_SIZE : The number of elements processed in X dimension

82

* -# -DVEC_SIZE_LEFTOVER: Leftover size in the X dimension; x_dimension % VEC_SIZE

83

*

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

84

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

85

* @param[in] in1_stride_x Stride of the source image in X dimension (in bytes)

86

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

87

* @param[in] in1_stride_y Stride of the source image in Y dimension (in bytes)

88

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

89

* @param[in] in1_offset_first_element_in_bytes The offset of the first element in the source image

90

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

91

* @param[in] in2_stride_x Stride of the source image in X dimension (in bytes)

92

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

93

* @param[in] in2_stride_y Stride of the source image in Y dimension (in bytes)

94

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

95

* @param[in] in2_offset_first_element_in_bytes The offset of the first element in the source image

96

* @param[out] out_ptr Pointer to the destination image. Supported data types: U8

97

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

98

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

99

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

100

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

101

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

102

*/

103

__kernel void bitwise_and(

104

IMAGE_DECLARATION(in1),

105

IMAGE_DECLARATION(in2),

106

IMAGE_DECLARATION(out))

107

{

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

108

uint x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0);

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

109

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

110

// Get pixels pointer

111

__global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x_offs + get_global_id(1) * in1_step_y;

112

__global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + x_offs + get_global_id(1) * in2_step_y;

113

__global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x_offs + get_global_id(1) * out_step_y;

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

114

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

115

// Load data

116

VEC_DATA_TYPE(uchar, VEC_SIZE)

117

in_a = VLOAD(VEC_SIZE)(0, (__global uchar *)in1_addr);

118

VEC_DATA_TYPE(uchar, VEC_SIZE)

119

in_b = VLOAD(VEC_SIZE)(0, (__global uchar *)in2_addr);

120

121

VEC_DATA_TYPE(uchar, VEC_SIZE)

122

data0 = in_a & in_b;

123

124

// Boundary-aware store

125

STORE_VECTOR_SELECT(data, uchar, (__global uchar *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0);

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

126

}

127

128

/** This function computes the bitwise XOR of two input images.

129

*

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

130

* @note The following variables must be passed at compile time:

131

* -# -DVEC_SIZE : The number of elements processed in X dimension

132

* -# -DVEC_SIZE_LEFTOVER: Leftover size in the X dimension; x_dimension % VEC_SIZE

133

*

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

134

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

135

* @param[in] in1_stride_x Stride of the source image in X dimension (in bytes)

136

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

137

* @param[in] in1_stride_y Stride of the source image in Y dimension (in bytes)

138

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

139

* @param[in] in1_offset_first_element_in_bytes The offset of the first element in the source image

140

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

141

* @param[in] in2_stride_x Stride of the source image in X dimension (in bytes)

142

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

143

* @param[in] in2_stride_y Stride of the source image in Y dimension (in bytes)

144

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

145

* @param[in] in2_offset_first_element_in_bytes The offset of the first element in the source image

146

* @param[out] out_ptr Pointer to the destination image. Supported data types: U8

147

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

148

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

149

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

150

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

151

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

152

*/

153

__kernel void bitwise_xor(

154

IMAGE_DECLARATION(in1),

155

IMAGE_DECLARATION(in2),

156

IMAGE_DECLARATION(out))

157

{

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

158

uint x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0);

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

159

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

160

// Get pixels pointer

161

__global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x_offs + get_global_id(1) * in1_step_y;

162

__global uchar *in2_addr = in2_ptr + in2_offset_first_element_in_bytes + x_offs + get_global_id(1) * in2_step_y;

163

__global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x_offs + get_global_id(1) * out_step_y;

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

164

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

165

// Load data

166

VEC_DATA_TYPE(uchar, VEC_SIZE)

167

in_a = VLOAD(VEC_SIZE)(0, (__global uchar *)in1_addr);

168

VEC_DATA_TYPE(uchar, VEC_SIZE)

169

in_b = VLOAD(VEC_SIZE)(0, (__global uchar *)in2_addr);

170

171

VEC_DATA_TYPE(uchar, VEC_SIZE)

172

data0 = in_a ^ in_b;

173

174

// Boundary-aware store

175

STORE_VECTOR_SELECT(data, uchar, (__global uchar *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0);

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

176

}

177

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

178

/** This function computes the bitwise NOT of an images.

179

*

180

* @note The following variables must be passed at compile time:

181

* -# -DVEC_SIZE : The number of elements processed in X dimension

182

* -# -DVEC_SIZE_LEFTOVER: Leftover size in the X dimension; x_dimension % VEC_SIZE

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

183

*

184

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

185

* @param[in] in_stride_x Stride of the source image in X dimension (in bytes)

186

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

187

* @param[in] in_stride_y Stride of the source image in Y dimension (in bytes)

188

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

189

* @param[in] in_offset_first_element_in_bytes The offset of the first element in the source image

190

* @param[out] out_ptr Pointer to the destination image. Supported data types: U8

191

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

192

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

193

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

194

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

195

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

196

*/

197

__kernel void bitwise_not(

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

198

IMAGE_DECLARATION(in1),

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

199

IMAGE_DECLARATION(out))

200

{

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

201

uint x_offs = max((int)(get_global_id(0) * VEC_SIZE - (VEC_SIZE - VEC_SIZE_LEFTOVER) % VEC_SIZE), 0);

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

202

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

203

// Get pixels pointer

204

__global uchar *in1_addr = in1_ptr + in1_offset_first_element_in_bytes + x_offs + get_global_id(1) * in1_step_y;

205

__global uchar *out_addr = out_ptr + out_offset_first_element_in_bytes + x_offs + get_global_id(1) * out_step_y;

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

206

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

207

// Load data

208

VEC_DATA_TYPE(uchar, VEC_SIZE)

209

in_a = VLOAD(VEC_SIZE)(0, (__global uchar *)in1_addr);

210

211

VEC_DATA_TYPE(uchar, VEC_SIZE)

212

data0 = ~in_a;

213

214

// Boundary-aware store

215

STORE_VECTOR_SELECT(data, uchar, (__global uchar *)out_addr, VEC_SIZE, VEC_SIZE_LEFTOVER, VEC_SIZE_LEFTOVER != 0 && get_global_id(0) == 0);

Anthony Barbier

6ff3b19

2017-09-04 18:44:23 +0100

[diff] [blame]

216

}

Manuel Bottini

63bb7ca

2020-12-02 13:22:14 +0000

[diff] [blame]

217

218

#endif // defined(VEC_SIZE) && defined(VEC_SIZE_LEFTOVER)