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

Anthony Barbier

6ff3b19

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

[diff] [blame]

1

/*

Michele Di Giorgio

d9eaf61

2020-07-08 11:12:57 +0100

[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"

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

27

*

28

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

29

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

30

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

31

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

32

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

33

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

34

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

35

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

36

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

37

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

38

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

39

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

40

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

41

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

42

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

43

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

44

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

45

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

46

*/

47

__kernel void bitwise_or(

48

IMAGE_DECLARATION(in1),

49

IMAGE_DECLARATION(in2),

50

IMAGE_DECLARATION(out))

51

{

52

Image in1 = CONVERT_TO_IMAGE_STRUCT(in1);

53

Image in2 = CONVERT_TO_IMAGE_STRUCT(in2);

54

Image out = CONVERT_TO_IMAGE_STRUCT(out);

55

56

uchar16 in_a = vload16(0, in1.ptr);

57

uchar16 in_b = vload16(0, in2.ptr);

58

59

vstore16(in_a | in_b, 0, out.ptr);

60

}

61

62

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

63

*

64

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

65

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

66

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

67

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

68

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

69

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

70

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

71

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

72

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

73

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

74

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

75

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

76

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

77

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

78

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

79

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

80

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

81

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

82

*/

83

__kernel void bitwise_and(

84

IMAGE_DECLARATION(in1),

85

IMAGE_DECLARATION(in2),

86

IMAGE_DECLARATION(out))

87

{

88

Image in1 = CONVERT_TO_IMAGE_STRUCT(in1);

89

Image in2 = CONVERT_TO_IMAGE_STRUCT(in2);

90

Image out = CONVERT_TO_IMAGE_STRUCT(out);

91

92

uchar16 in_a = vload16(0, in1.ptr);

93

uchar16 in_b = vload16(0, in2.ptr);

94

95

vstore16(in_a & in_b, 0, out.ptr);

96

}

97

98

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

99

*

100

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

101

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

102

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

103

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

104

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

105

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

106

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

107

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

108

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

109

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

110

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

111

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

112

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

113

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

114

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

115

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

116

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

117

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

118

*/

119

__kernel void bitwise_xor(

120

IMAGE_DECLARATION(in1),

121

IMAGE_DECLARATION(in2),

122

IMAGE_DECLARATION(out))

123

{

124

Image in1 = CONVERT_TO_IMAGE_STRUCT(in1);

125

Image in2 = CONVERT_TO_IMAGE_STRUCT(in2);

126

Image out = CONVERT_TO_IMAGE_STRUCT(out);

127

128

uchar16 in_a = vload16(0, in1.ptr);

129

uchar16 in_b = vload16(0, in2.ptr);

130

131

vstore16(in_a ^ in_b, 0, out.ptr);

132

}

133

134

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

135

*

136

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

137

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

138

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

139

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

140

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

141

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

142

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

143

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

144

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

145

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

146

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

147

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

148

*/

149

__kernel void bitwise_not(

150

IMAGE_DECLARATION(in),

151

IMAGE_DECLARATION(out))

152

{

153

Image in = CONVERT_TO_IMAGE_STRUCT(in);

154

Image out = CONVERT_TO_IMAGE_STRUCT(out);

155

156

uchar16 in_data = vload16(0, in.ptr);

157

158

vstore16(~in_data, 0, out.ptr);

159

}