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

#if defined(DATA_TYPE) && defined(POOLED_DIM_X) && defined(POOLED_DIM_Y) && defined(MAX_DIM_X) && defined(MAX_DIM_Y) && defined(MAX_DIM_Z) && defined(SPATIAL_SCALE) // Check for compile time constants

31

32

/** Performs a roi align on a single output pixel.

33

*

34

* @param[in] input Pointer to input Tensor3D struct.

35

* @param[in] region_start_x Start x index projected onto the input tensor.

36

* @param[in] region_end_x End x index projected onto the input tensor.

37

* @param[in] region_start_y Start y index projected onto the input tensor.

38

* @param[in] region_end_y End y index projected onto the input tensor.

39

* @param[in] pz z index of the input tensor.

40

*

41

* @return An average pooled value from the region specified in the input tensor.

42

*/

43

inline DATA_TYPE roi_align_1x1(const Tensor3D *input, float region_start_x,

float bin_size_x,

float grid_size_x,

float region_end_x,

float region_start_y,

float bin_size_y,

float grid_size_y,

float region_end_y,

int pz)

{

// Iterate through the pooling region

54

float sum = 0;

55

for(int iy = 0; iy < grid_size_y; ++iy)

56

{

57

for(int ix = 0; ix < grid_size_x; ++ix)

58

{

59

// Align the window in the middle of every bin

60

const float y = region_start_y + (iy + 0.5f) * bin_size_y / (float)grid_size_y;

61

const float x = region_start_x + (ix + 0.5f) * bin_size_x / (float)grid_size_x;

62

63

// Interpolation in the unit square

64

const int y_low = (int)y;

65

const int x_low = (int)x;

66

const int y_high = y_low + 1;

67

const int x_high = x_low + 1;

68

69

const float ly = y - y_low;

70

const float lx = x - x_low;

71

const float hy = 1.f - ly;

72

const float hx = 1.f - lx;

73

74

const float w1 = hy * hx;

75

const float w2 = hy * lx;

76

const float w3 = ly * hx;

77

const float w4 = ly * lx;

78

79

const DATA_TYPE data1 = *(__global DATA_TYPE *)tensor3D_offset(input, x_low, y_low, pz);

80

const DATA_TYPE data2 = *(__global DATA_TYPE *)tensor3D_offset(input, x_high, y_low, pz);

81

const DATA_TYPE data3 = *(__global DATA_TYPE *)tensor3D_offset(input, x_low, y_high, pz);

82

const DATA_TYPE data4 = *(__global DATA_TYPE *)tensor3D_offset(input, x_high, y_high, pz);

83

sum += w1 * data1 + w2 * data2 + w3 * data3 + w4 * data4;

}

}

return (DATA_TYPE)(sum / (grid_size_x * grid_size_y));

88

}

89

90

/** Performs a roi align function.

91

*

92

* @note Datatype must be passed using -DDATA_TYPE e.g. -DDATA_TYPE=float. Supported data types are F16, F32;

93

* @note Datasize must be passed using -DDATA_SIZE e.g. -DDATA_SIZE=32;

94

* @note Input dimensions must be passed using -DMAX_DIM_X, -DMAX_DIM_Y and -DMAX_DIM_Z;

95

* @note Pooled region dimensions must be passed using -DPOOLED_DIM_X and -DPOOLED_DIM_Y;

96

* @note Spatial scale must be passed using -DSPATIAL_SCALE;

97

* @note Sampling ratio (i.e., the number of samples in each bin) may be passed using -DSAMPLING_RATIO. If not defined each roi

98

* will have a default sampling ratio of roi_dims/pooling_dims

99

*

100

* @param[in] input_ptr Pointer to the source image. Supported data types: F16, F32

101

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

102

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

103

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

104

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

105

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

106

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

107

* @param[in] input_offset_first_element_in_bytes The offset of the first element in the pooled region of the source image as specifed by ROI

108

* @param[in] rois_ptr Pointer to the rois array. Layout: {x, y, width, height, batch_indx}

109

* @param[in] rois_stride_x Stride of the rois array in X dimension (in bytes)

110

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

111

* @param[in] rois_offset_first_element_in_bytes The offset of the first element in the rois array

112

* @param[out] output_ptr Pointer to the destination image. Supported data types: F16, F32

113

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

114

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

115

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

116

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

117

* @param[in] output_stride_z Stride of the destination tensor in Z dimension (in bytes)

118

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

119

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

120

* @param[in] input_stride_w Stride of the source image in W dimension (in bytes)

121

* @param[in] output_stride_w Stride of the destination image in W dimension (in bytes)

122

*/

123

__kernel void roi_align_layer(

124

TENSOR3D_DECLARATION(input),

125

VECTOR_DECLARATION(rois),

126

TENSOR3D_DECLARATION(output),

127

unsigned int input_stride_w, unsigned int output_stride_w)

128

{

129

// Get pixels pointer

130

Tensor3D input = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(input);

131

Vector rois = CONVERT_TO_VECTOR_STRUCT_NO_STEP(rois);

132

Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(output);

133

134

const int px = get_global_id(0);

135

const int py = get_global_id(1);

136

const int pw = get_global_id(2);

137

138

// Load roi parameters

139

// roi is laid out as follows:

140

// { x, y, width, height, batch_index }

141

const ushort4 roi = vload4(0, (__global ushort *)vector_offset(&rois, pw));

142

const ushort roi_batch = *((__global ushort *)vector_offset(&rois, pw) + 4);

143

const float2 roi_anchor = convert_float2(roi.s01) * convert_float(SPATIAL_SCALE);

144

const float2 roi_dims = fmax(convert_float2(roi.s23) * convert_float(SPATIAL_SCALE), 1.f);

145

146

// Calculate pooled region start and end

147

const float2 spatial_indx = (float2)(px, py);

148

const float2 pooled_dims = (float2)(POOLED_DIM_X, POOLED_DIM_Y);

149

const float2 max_spatial_dims = (float2)(MAX_DIM_X, MAX_DIM_Y);

150

151

const float2 bin_size = roi_dims / pooled_dims;

152

float2 region_start = spatial_indx * bin_size + roi_anchor;

153

float2 region_end = (spatial_indx + 1) * bin_size + roi_anchor;

154

155

region_start = clamp(region_start, 0, max_spatial_dims);

156

region_end = clamp(region_end, 0, max_spatial_dims);

157

158

#if defined(SAMPLING_RATIO)

159

const float2 roi_bin_grid = SAMPLING_RATIO;

160

#else // !defined(SAMPLING_RATIO)

161

// Note that we subtract EPS_GRID before ceiling. This is to avoid situations where 1.000001 gets ceiled to 2.

162

const float2 roi_bin_grid = ceil(roi_dims / pooled_dims - EPS_GRID);

163

#endif // defined(SAMPLING_RATIO)

164

165

// Move input and output pointer across the fourth dimension

166

input.ptr += roi_batch * input_stride_w;

167

output.ptr += pw * output_stride_w;

168

for(int pz = 0; pz < MAX_DIM_Z; ++pz)

169

{

170

*(__global DATA_TYPE *)tensor3D_offset(&output, px, py, pz) = (__global DATA_TYPE)roi_align_1x1(&input,

region_start.x,

bin_size.x,

roi_bin_grid.x,

region_end.x,

region_start.y,

bin_size.y,

roi_bin_grid.y,

region_end.y, pz);

}

}

#endif // Check for compile time constants