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

Georgios Pinitas

7900a9e

2018-11-23 11:44:58 +0000

[diff] [blame]

/*

*

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

#define EQUAL(x, y) ((x) == (y))

27

#define NOTEQUAL(x, y) ((x) != (y))

28

#define GREATER(x, y) ((x) > (y))

29

#define GREATEREQUAL(x, y) ((x) >= (y))

30

#define LESS(x, y) ((x) < (y))

31

#define LESSEQUAL(x, y) ((x) <= (y))

32

33

#define DEFINE_KERNEL_STR(name) compare_##name

34

#define DEFINE_KERNEL(name) DEFINE_KERNEL_STR(name)

35

36

#define DEFINE_KERNEL_QUANTIZED_STR(name) compare_##name##_quantized

37

#define DEFINE_KERNEL_QUANTIZED(name) DEFINE_KERNEL_QUANTIZED_STR(name)

38

39

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

40

/** This function compares two tensors.

41

*

42

* @attention The inputs' data type need to be passed at compile time using -DDATA_TYPE: e.g. -DDATA_TYPE=float

43

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

44

* @attention The comparison operation should be given as a preprocessor argument using -DOP=operation. e.g. -DOP=LESS

45

*

46

* @param[in] in1_ptr Pointer to the source tensor. Supported data types: U8/S16/F16/F32

47

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

48

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

49

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

50

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

51

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

52

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

53

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

54

* @param[in] in2_ptr Pointer to the source tensor. Supported data types: U8/S16/F16/F32

55

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

56

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

57

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

58

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

59

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

60

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

61

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

62

* @param[out] out_ptr Pointer to the destination tensor. Supported data types: U8 (only if both inputs are U8), S16/F16/F32

63

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

64

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

65

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

66

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

67

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

68

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

69

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

70

*/

71

__kernel void DEFINE_KERNEL(OP_NAME)(

72

TENSOR3D_DECLARATION(in1),

73

TENSOR3D_DECLARATION(in2),

74

TENSOR3D_DECLARATION(out))

75

{

76

// Get pixels pointer

77

Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1);

78

Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2);

79

Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out);

80

81

// Load values

82

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

83

in_a = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in1.ptr);

84

VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)

85

in_b = VLOAD(VEC_SIZE)(0, (__global DATA_TYPE *)in2.ptr);

86

87

// Calculate and store result

88

VSTORE(VEC_SIZE)

89

(CONVERT(OP(in_a, in_b), VEC_DATA_TYPE(uchar, VEC_SIZE)), 0, (__global uchar *)out.ptr);

90

}

91

#endif /* defined(DATA_TYPE) && defined(VEC_SIZE) && defined(OP) && defined(OP_NAME) */

92

93

#if defined(OFFSET_IN1) && defined(OFFSET_IN2) && defined(SCALE_IN1) && defined(SCALE_IN2)

94

/** This function compares two quantized tensors.

95

*

96

* @note The quantization offset of the first operand must be passed at compile time using -DOFFSET_IN1, i.e. -DOFFSET_IN1=10

97

* @note The quantization offset of the second operand must be passed at compile time using -DOFFSET_IN2, i.e. -DOFFSET_IN2=10

98

* @note The quantization scale of the first operand must be passed at compile time using -DSCALE_IN1, i.e. -DSCALE_IN1=10

99

* @note The quantization scale of the second operand must be passed at compile time using -DSCALE_IN2, i.e. -DSCALE_IN2=10

100

*

101

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

102

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

103

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

104

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

105

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

106

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

107

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

108

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

109

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

110

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

111

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

112

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

113

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

114

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

115

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

116

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

117

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

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 DEFINE_KERNEL_QUANTIZED(OP_NAME)(

127

TENSOR3D_DECLARATION(in1),

128

TENSOR3D_DECLARATION(in2),

129

TENSOR3D_DECLARATION(out))

130

{

131

// Get pixels pointer

132

Tensor3D in1 = CONVERT_TO_TENSOR3D_STRUCT(in1);

133

Tensor3D in2 = CONVERT_TO_TENSOR3D_STRUCT(in2);

134

Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out);

135

136

int16 in_a = CONVERT(vload16(0, (__global uchar *)in1.ptr), int16);

137

int16 in_b = CONVERT(vload16(0, (__global uchar *)in2.ptr), int16);

138

139

in_a = in_a - (int16)((int)OFFSET_IN1);

140

in_b = in_b - (int16)((int)OFFSET_IN2);

141

142

const float16 in1f32 = convert_float16(in_a) * (float16)((float)SCALE_IN1);

143

const float16 in2f32 = convert_float16(in_b) * (float16)((float)SCALE_IN2);

144

const int16 res = OP(in1f32, in2f32);

145

146

// Store result

147

vstore16(convert_uchar16(res), 0, (__global uchar *)out.ptr);

148

}

149

#endif /* defined(OFFSET_IN1) && defined(OFFSET_IN2) && defined(SCALE_IN1) && defined(SCALE_IN2) */