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

2017-11-03 17:33:15 +0700

[diff] [blame]

1

/*

Michele Di Giorgio

d9eaf61

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

[diff] [blame]

2

Chunosov

2017-11-03 17:33:15 +0700

[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

22

* SOFTWARE.

23

*/

24

#include "helpers_asymm.h"

25

Sheri Zhang

2020-02-20 11:23:08 +0000

[diff] [blame]

26

#undef CONVERT_SAT_STR

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

27

#undef CONVERT_SAT

28

Michele Di Giorgio

2019-10-23 10:53:10 +0100

[diff] [blame]

29

#if defined(DATA_TYPE) && defined(STRIDE_X) && defined(WEIGHTS_DEPTH) && defined(OUTPUT_MULTIPLIER) && defined(OUTPUT_SHIFT)

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

30

Sheri Zhang

2020-02-20 11:23:08 +0000

[diff] [blame]

31

#define CONVERT_SAT_STR(x, type) (convert_##type##8_sat((x)))

32

#define CONVERT_SAT(x, type) CONVERT_SAT_STR(x, type)

33

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

34

#if defined(DATA_LAYOUT_NHWC)

35

Georgios Pinitas

aa95ddc

2020-07-21 22:45:13 +0100

[diff] [blame]

#if KERNEL_SIZE == 9

#if STRIDE_X == 1

#define CONVOLUTION1x9(acc, src_ptr, weights_ptr) CONVOLUTION1x9_STRIDE1(acc, src_ptr, weights_ptr)

40

#elif STRIDE_X == 2

41

#define CONVOLUTION1x9(acc, src_ptr, weights_ptr) CONVOLUTION1x9_STRIDE2(acc, src_ptr, weights_ptr)

42

#else /* STRIDE_X not equals 1 or 2 */

43

#error "STRIDE_X larger than 2 is not supported"

44

#endif /* STRIDE_X */

45

46

#define CONVOLUTION1x9_STRIDE1(acc, src_ptr, weights_ptr) \

47

({ \

48

int8 weights_values0 = 0; \

49

int weights_value1 = 0; \

50

weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \

51

weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \

52

weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \

53

weights_values0.s3 = convert_int(*(weights_ptr + 3 * weights_stride_y)); \

54

weights_values0.s4 = convert_int(*(weights_ptr + 4 * weights_stride_y)); \

55

weights_values0.s5 = convert_int(*(weights_ptr + 5 * weights_stride_y)); \

56

weights_values0.s6 = convert_int(*(weights_ptr + 6 * weights_stride_y)); \

57

weights_values0.s7 = convert_int(*(weights_ptr + 7 * weights_stride_y)); \

58

weights_value1 = convert_int(*(weights_ptr + 8 * weights_stride_y)); \

\

int8 src0 = 0; \

int8 src1 = 0; \

src0.s0 = convert_int(*(src_ptr + 0 * weights_stride_y)); \

63

src0.s1 = convert_int(*(src_ptr + 1 * weights_stride_y)); \

64

src0.s2 = convert_int(*(src_ptr + 2 * weights_stride_y)); \

65

src0.s3 = convert_int(*(src_ptr + 3 * weights_stride_y)); \

66

src0.s4 = convert_int(*(src_ptr + 4 * weights_stride_y)); \

67

src0.s5 = convert_int(*(src_ptr + 5 * weights_stride_y)); \

68

src0.s6 = convert_int(*(src_ptr + 6 * weights_stride_y)); \

69

src0.s7 = convert_int(*(src_ptr + 7 * weights_stride_y)); \

70

src1.s0 = convert_int(*(src_ptr + 8 * weights_stride_y)); \

71

src1.s1 = convert_int(*(src_ptr + 9 * weights_stride_y)); \

72

src1.s2 = convert_int(*(src_ptr + 10 * weights_stride_y)); \

73

src1.s3 = convert_int(*(src_ptr + 11 * weights_stride_y)); \

74

src1.s4 = convert_int(*(src_ptr + 12 * weights_stride_y)); \

75

src1.s5 = convert_int(*(src_ptr + 13 * weights_stride_y)); \

76

src1.s6 = convert_int(*(src_ptr + 14 * weights_stride_y)); \

77

src1.s7 = convert_int(*(src_ptr + 15 * weights_stride_y)); \

78

\

79

acc += src0 * (int8)weights_values0.s0; \

80

acc += (int8)(src0.s1234, src0.s567, src1.s0) * (int8)weights_values0.s1; \

81

acc += (int8)(src0.s234, src0.s567, src1.s01) * (int8)weights_values0.s2; \

82

acc += (int8)(src0.s345, src0.s67, src1.s012) * (int8)weights_values0.s3; \

83

acc += (int8)(src0.s4567, src1.s0123) * (int8)weights_values0.s4; \

84

acc += (int8)(src0.s567, src1.s0123, src1.s4) * (int8)weights_values0.s5; \

85

acc += (int8)(src0.s67, src1.s012, src1.s345) * (int8)weights_values0.s6; \

86

acc += (int8)(src0.s7, src1.s0123, src1.s456) * (int8)weights_values0.s7; \

87

acc += src1 * (int8)weights_value1; \

88

})

89

90

#define CONVOLUTION1x9_STRIDE2(acc, src_ptr, weights_ptr) \

91

({ \

92

int8 weights_values0 = 0; \

93

int weights_value1 = 0; \

94

weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \

95

weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \

96

weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \

97

weights_values0.s3 = convert_int(*(weights_ptr + 3 * weights_stride_y)); \

98

weights_values0.s4 = convert_int(*(weights_ptr + 4 * weights_stride_y)); \

99

weights_values0.s5 = convert_int(*(weights_ptr + 5 * weights_stride_y)); \

100

weights_values0.s6 = convert_int(*(weights_ptr + 6 * weights_stride_y)); \

101

weights_values0.s7 = convert_int(*(weights_ptr + 7 * weights_stride_y)); \

102

weights_value1 = convert_int(*(weights_ptr + 8 * weights_stride_y)); \

\

int16 src0 = 0; \

int8 src1 = 0; \

src0.s0 = convert_int(*(src_ptr + 0 * weights_stride_y)); \

107

src0.s1 = convert_int(*(src_ptr + 1 * weights_stride_y)); \

108

src0.s2 = convert_int(*(src_ptr + 2 * weights_stride_y)); \

109

src0.s3 = convert_int(*(src_ptr + 3 * weights_stride_y)); \

110

src0.s4 = convert_int(*(src_ptr + 4 * weights_stride_y)); \

111

src0.s5 = convert_int(*(src_ptr + 5 * weights_stride_y)); \

112

src0.s6 = convert_int(*(src_ptr + 6 * weights_stride_y)); \

113

src0.s7 = convert_int(*(src_ptr + 7 * weights_stride_y)); \

114

src0.s8 = convert_int(*(src_ptr + 8 * weights_stride_y)); \

115

src0.s9 = convert_int(*(src_ptr + 9 * weights_stride_y)); \

116

src0.sA = convert_int(*(src_ptr + 10 * weights_stride_y)); \

117

src0.sB = convert_int(*(src_ptr + 11 * weights_stride_y)); \

118

src0.sC = convert_int(*(src_ptr + 12 * weights_stride_y)); \

119

src0.sD = convert_int(*(src_ptr + 13 * weights_stride_y)); \

120

src0.sE = convert_int(*(src_ptr + 14 * weights_stride_y)); \

121

src0.sF = convert_int(*(src_ptr + 15 * weights_stride_y)); \

122

src1.s0 = convert_int(*(src_ptr + 16 * weights_stride_y)); \

123

src1.s1 = convert_int(*(src_ptr + 17 * weights_stride_y)); \

124

src1.s2 = convert_int(*(src_ptr + 18 * weights_stride_y)); \

125

src1.s3 = convert_int(*(src_ptr + 19 * weights_stride_y)); \

126

src1.s4 = convert_int(*(src_ptr + 20 * weights_stride_y)); \

127

src1.s5 = convert_int(*(src_ptr + 21 * weights_stride_y)); \

128

src1.s6 = convert_int(*(src_ptr + 22 * weights_stride_y)); \

129

src1.s7 = convert_int(*(src_ptr + 23 * weights_stride_y)); \

130

\

131

acc += src0.s02468ACE * (int8)weights_values0.s0; \

132

acc += (int8)(src0.s1357, src0.s9BDF) * (int8)weights_values0.s1; \

133

acc += (int8)(src0.s2468, src0.sACE, src1.s0) * (int8)weights_values0.s2; \

134

acc += (int8)(src0.s3579, src0.sBDF, src1.s1) * (int8)weights_values0.s3; \

135

acc += (int8)(src0.s468A, src0.sCE, src1.s02) * (int8)weights_values0.s4; \

136

acc += (int8)(src0.s579, src0.sBDF, src1.s13) * (int8)weights_values0.s5; \

137

acc += (int8)(src0.s68A, src0.sCE, src1.s024) * (int8)weights_values0.s6; \

138

acc += (int8)(src0.s79B, src0.sDF, src1.s135) * (int8)weights_values0.s7; \

139

acc += (int8)(src0.s8AC, src0.sE, src1.s0246) * (int8)weights_value1; \

140

})

141

142

#elif KERNEL_SIZE == 5

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

143

144

#if STRIDE_X == 1

145

#define CONVOLUTION1x5(acc, src_ptr, weights_ptr) CONVOLUTION1x5_STRIDE1(acc, src_ptr, weights_ptr)

146

#elif STRIDE_X == 2

147

#define CONVOLUTION1x5(acc, src_ptr, weights_ptr) CONVOLUTION1x5_STRIDE2(acc, src_ptr, weights_ptr)

148

#else /* STRIDE_X not equals 1 or 2 */

149

#error "STRIDE_X larger than 2 is not supported"

150

#endif /* STRIDE_X */

151

152

#define CONVOLUTION1x5_STRIDE1(acc, src_ptr, weights_ptr) \

153

({ \

154

int4 weights_values0 = 0; \

155

int weights_value1 = 0; \

156

weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \

157

weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \

158

weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \

159

weights_values0.s3 = convert_int(*(weights_ptr + 3 * weights_stride_y)); \

160

weights_value1 = convert_int(*(weights_ptr + 4 * weights_stride_y)); \

\

int8 src0 = 0; \

int4 src1 = 0; \

src0.s0 = convert_int(*(src_ptr + 0 * weights_stride_y)); \

165

src0.s1 = convert_int(*(src_ptr + 1 * weights_stride_y)); \

166

src0.s2 = convert_int(*(src_ptr + 2 * weights_stride_y)); \

167

src0.s3 = convert_int(*(src_ptr + 3 * weights_stride_y)); \

168

src0.s4 = convert_int(*(src_ptr + 4 * weights_stride_y)); \

169

src0.s5 = convert_int(*(src_ptr + 5 * weights_stride_y)); \

170

src0.s6 = convert_int(*(src_ptr + 6 * weights_stride_y)); \

171

src0.s7 = convert_int(*(src_ptr + 7 * weights_stride_y)); \

172

src1.s0 = convert_int(*(src_ptr + 8 * weights_stride_y)); \

173

src1.s1 = convert_int(*(src_ptr + 9 * weights_stride_y)); \

174

src1.s2 = convert_int(*(src_ptr + 10 * weights_stride_y)); \

175

src1.s3 = convert_int(*(src_ptr + 11 * weights_stride_y)); \

176

\

177

acc += (src0 + input_offset) * ((int8)weights_values0.s0 + weight_offset); \

178

acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \

179

acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \

180

acc += ((int8)(src0.s345, src0.s67, src1.s012) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \

181

acc += ((int8)(src0.s45, src0.s67, src1.s0123) + input_offset) * ((int8)weights_value1 + weight_offset); \

182

})

183

184

#define CONVOLUTION1x5_STRIDE2(acc, src_ptr, weights_ptr) \

185

({ \

186

int4 weights_values0 = 0; \

187

int weights_value1 = 0; \

188

weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \

189

weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \

190

weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \

191

weights_values0.s3 = convert_int(*(weights_ptr + 3 * weights_stride_y)); \

192

weights_value1 = convert_int(*(weights_ptr + 4 * weights_stride_y)); \

\

int16 src0 = 0; \

int4 src1 = 0; \

src0.s0 = convert_int(*(src_ptr + 0 * weights_stride_y)); \

197

src0.s1 = convert_int(*(src_ptr + 1 * weights_stride_y)); \

198

src0.s2 = convert_int(*(src_ptr + 2 * weights_stride_y)); \

199

src0.s3 = convert_int(*(src_ptr + 3 * weights_stride_y)); \

200

src0.s4 = convert_int(*(src_ptr + 4 * weights_stride_y)); \

201

src0.s5 = convert_int(*(src_ptr + 5 * weights_stride_y)); \

202

src0.s6 = convert_int(*(src_ptr + 6 * weights_stride_y)); \

203

src0.s7 = convert_int(*(src_ptr + 7 * weights_stride_y)); \

204

src0.s8 = convert_int(*(src_ptr + 8 * weights_stride_y)); \

205

src0.s9 = convert_int(*(src_ptr + 9 * weights_stride_y)); \

206

src0.sa = convert_int(*(src_ptr + 10 * weights_stride_y)); \

207

src0.sb = convert_int(*(src_ptr + 11 * weights_stride_y)); \

208

src0.sc = convert_int(*(src_ptr + 12 * weights_stride_y)); \

209

src0.sd = convert_int(*(src_ptr + 13 * weights_stride_y)); \

210

src0.se = convert_int(*(src_ptr + 14 * weights_stride_y)); \

211

src0.sf = convert_int(*(src_ptr + 15 * weights_stride_y)); \

212

src1.s0 = convert_int(*(src_ptr + 16 * weights_stride_y)); \

213

src1.s1 = convert_int(*(src_ptr + 17 * weights_stride_y)); \

214

src1.s2 = convert_int(*(src_ptr + 18 * weights_stride_y)); \

215

src1.s3 = convert_int(*(src_ptr + 19 * weights_stride_y)); \

216

\

217

acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \

218

acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \

219

acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \

220

acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \

221

acc += ((int8)(src0.s468a, src0.sCE, src1.s02) + input_offset) * ((int8)weights_value1 + weight_offset); \

222

})

223

224

#elif KERNEL_SIZE == 3

225

226

#if STRIDE_X == 1

227

#define CONVOLUTION1x3(acc, src_ptr, weights_ptr) CONVOLUTION1x3_STRIDE1(acc, src_ptr, weights_ptr)

228

#elif STRIDE_X == 2

229

#define CONVOLUTION1x3(acc, src_ptr, weights_ptr) CONVOLUTION1x3_STRIDE2(acc, src_ptr, weights_ptr)

230

#else /* STRIDE_X not equals 1 or 2 */

231

#error "STRIDE_X larger than 2 is not supported"

232

#endif /* STRIDE_X */

233

234

#define CONVOLUTION1x3_STRIDE1(acc, src_ptr, weights_ptr) \

235

({ \

236

int3 weights_values0 = 0; \

237

weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \

238

weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \

239

weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \

\

int8 src0 = 0; \

int2 src1 = 0; \

src0.s0 = convert_int(*(src_ptr + 0 * weights_stride_y)); \

244

src0.s1 = convert_int(*(src_ptr + 1 * weights_stride_y)); \

245

src0.s2 = convert_int(*(src_ptr + 2 * weights_stride_y)); \

246

src0.s3 = convert_int(*(src_ptr + 3 * weights_stride_y)); \

247

src0.s4 = convert_int(*(src_ptr + 4 * weights_stride_y)); \

248

src0.s5 = convert_int(*(src_ptr + 5 * weights_stride_y)); \

249

src0.s6 = convert_int(*(src_ptr + 6 * weights_stride_y)); \

250

src0.s7 = convert_int(*(src_ptr + 7 * weights_stride_y)); \

251

src1.s0 = convert_int(*(src_ptr + 8 * weights_stride_y)); \

252

src1.s1 = convert_int(*(src_ptr + 9 * weights_stride_y)); \

253

\

254

acc += (src0 + input_offset) * ((int8)weights_values0.s0 + weight_offset); \

255

acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \

256

acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \

257

})

258

259

#define CONVOLUTION1x3_STRIDE2(acc, src_ptr, weights_ptr) \

260

({ \

261

int3 weights_values0 = 0; \

262

weights_values0.s0 = convert_int(*(weights_ptr + 0 * weights_stride_y)); \

263

weights_values0.s1 = convert_int(*(weights_ptr + 1 * weights_stride_y)); \

264

weights_values0.s2 = convert_int(*(weights_ptr + 2 * weights_stride_y)); \

\

int16 src0 = 0; \

int src1 = 0; \

src0.s0 = convert_int(*(src_ptr + 0 * src_stride_y)); \

269

src0.s1 = convert_int(*(src_ptr + 1 * src_stride_y)); \

270

src0.s2 = convert_int(*(src_ptr + 2 * src_stride_y)); \

271

src0.s3 = convert_int(*(src_ptr + 3 * src_stride_y)); \

272

src0.s4 = convert_int(*(src_ptr + 4 * src_stride_y)); \

273

src0.s5 = convert_int(*(src_ptr + 5 * src_stride_y)); \

274

src0.s6 = convert_int(*(src_ptr + 6 * src_stride_y)); \

275

src0.s7 = convert_int(*(src_ptr + 7 * src_stride_y)); \

276

src0.s8 = convert_int(*(src_ptr + 8 * src_stride_y)); \

277

src0.s9 = convert_int(*(src_ptr + 9 * src_stride_y)); \

278

src0.sa = convert_int(*(src_ptr + 10 * src_stride_y)); \

279

src0.sb = convert_int(*(src_ptr + 11 * src_stride_y)); \

280

src0.sc = convert_int(*(src_ptr + 12 * src_stride_y)); \

281

src0.sd = convert_int(*(src_ptr + 13 * src_stride_y)); \

282

src0.se = convert_int(*(src_ptr + 14 * src_stride_y)); \

283

src0.sf = convert_int(*(src_ptr + 15 * src_stride_y)); \

284

src1 = convert_int(*(src_ptr + 16 * src_stride_y)); \

285

acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \

286

acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \

287

acc += ((int8)(src0.s2468, src0.sACE, src1) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \

288

})

289

290

#elif KERNEL_SIZE == 1

291

292

#if STRIDE_X == 3

293

#define INPUT_VALUE extract_input_stride3

294

#elif STRIDE_X == 2

295

#define INPUT_VALUE extract_input_stride2

296

#elif STRIDE_X == 1

297

#define INPUT_VALUE extract_input_stride1

298

299

#else /* STRIDE_X not equals 1, 2 or 3 */

300

#error "Only support strides 1, 2 and 3"

301

#endif /* STRIDE_X */

302

303

#endif // KERNEL_SIZE == 1

304

305

/** Extracts a 1D horizontal vector from the input tensor with stride as 1.

306

*

307

* @param[in] input_value Pointer to the first value.

308

*

309

* @return extracted input values.

310

*/

311

inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride1(__global const DATA_TYPE *input_value, const uchar stride_y)

312

{

313

VEC_DATA_TYPE(DATA_TYPE, 8)

314

vals;

315

vals.s0 = *(input_value + 0 * stride_y);

316

vals.s1 = *(input_value + 1 * stride_y);

317

vals.s2 = *(input_value + 2 * stride_y);

318

vals.s3 = *(input_value + 3 * stride_y);

319

vals.s4 = *(input_value + 4 * stride_y);

320

vals.s5 = *(input_value + 5 * stride_y);

321

vals.s6 = *(input_value + 6 * stride_y);

322

vals.s7 = *(input_value + 7 * stride_y);

return vals;

}

/** Extracts a 1D horizontal vector from the input tensor with stride as 2.

328

*

329

* @param[in] input_value Pointer to the first value.

330

*

331

* @return extracted input values.

332

*/

333

inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride2(__global const DATA_TYPE *input_value, const uchar stride_y)

334

{

335

VEC_DATA_TYPE(DATA_TYPE, 8)

336

vals;

337

vals.s0 = *(input_value + 0 * stride_y);

338

vals.s1 = *(input_value + 2 * stride_y);

339

vals.s2 = *(input_value + 4 * stride_y);

340

vals.s3 = *(input_value + 6 * stride_y);

341

vals.s4 = *(input_value + 8 * stride_y);

342

vals.s5 = *(input_value + 10 * stride_y);

343

vals.s6 = *(input_value + 12 * stride_y);

344

vals.s7 = *(input_value + 14 * stride_y);

return vals;

}

/** Extracts a 1D horizontal vector from the input tensor with stride as 3 and 8-bit data size.

350

*

351

* @param[in] input_value Pointer to the first value.

352

*

353

* @return extracted input values.

354

*/

355

inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3(__global const DATA_TYPE *input_value, const uchar stride_y)

356

{

357

VEC_DATA_TYPE(DATA_TYPE, 8)

358

vals;

359

vals.s0 = *(input_value + 0 * stride_y);

360

vals.s1 = *(input_value + 3 * stride_y);

361

vals.s2 = *(input_value + 6 * stride_y);

362

vals.s3 = *(input_value + 9 * stride_y);

363

vals.s4 = *(input_value + 12 * stride_y);

364

vals.s5 = *(input_value + 15 * stride_y);

365

vals.s6 = *(input_value + 18 * stride_y);

366

vals.s7 = *(input_value + 21 * stride_y);

return vals;

}

/** This kernel performs a direct convolution to convolve the low three dimensions.

372

*

373

* @note The convolution stride x must be passed at compile time using -DSTRIDE_X e.g. -DSTRIDE_X=1

374

* @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH

375

* @note If biases are used then -DHAS_BIAS has to be passed at compile time

376

* @note The output quantization multiplier must be passed at compile time using -DOUTPUT_MULTIPLIER e.g. -DOUTPUT_MULTIPLIER=1234

377

* @note The output quantization shift must be passed at compile time using -DOUTPUT_SHIFT e.g. -DOUTPUT_SHIFT=4

378

*

379

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

380

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

381

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

382

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

383

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

384

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

385

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

386

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

387

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

388

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

389

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

390

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

391

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

392

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

393

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

394

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

395

* @param[in] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr

396

* @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes)

397

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

398

* @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes)

399

* @param[in] weights_step_y weights_stride_y * number of elements along y processed per workitem(in bytes)

400

* @param[in] weights_stride_z Stride of the weights tensor in Z dimension (in bytes)

401

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

402

* @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor

403

* @param[in] biases_ptr Pointer to the biases tensor. Supported data types: S32

404

* @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes)

405

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

406

* @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor

407

* @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension

408

* @param[in] input_offset Input offset quantization parameter

409

* @param[in] weight_offset Weights offset quantization parameter

410

* @param[in] output_offset Output offset quantization parameter

411

*/

412

__kernel void direct_convolution_quantized(

413

TENSOR3D_DECLARATION(src),

414

TENSOR3D_DECLARATION(dst),

415

TENSOR3D_DECLARATION(weights),

416

#ifdef HAS_BIAS

417

VECTOR_DECLARATION(biases),

418

#endif /* defined(HAS_BIAS) */

419

unsigned int weights_stride_w,

int input_offset,

int weight_offset,

int output_offset)

{

Image src = CONVERT_TO_IMAGE_STRUCT(src);

425

Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights);

426

Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);

int8 values0 = 0;

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

430

const int id0 = get_global_id(0);

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

431

const int y_coord = (get_global_id(2) * STRIDE_Y) - PAD_TOP;

432

433

__global DATA_TYPE *weights_addr = (__global DATA_TYPE *)tensor3D_offset(&weights, 0, 0, 0);

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

434

__global DATA_TYPE *src_addr = (__global DATA_TYPE *)offset(&src, 0, 0) - src_stride_x * id0 + y_coord * (int)src_stride_z;

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

435

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

436

weights_addr += id0 * weights_stride_w;

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

437

438

for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d)

439

{

Georgios Pinitas

aa95ddc

2020-07-21 22:45:13 +0100

[diff] [blame]

#if KERNEL_SIZE == 9

if(y_coord < 0)

{

const int start_z = -y_coord;

444

for(int i = start_z; i < 9; ++i)

445

{

446

CONVOLUTION1x9(values0, (src_addr + i * (int)src_stride_z), (weights_addr + i * (int)weights_stride_z));

447

}

448

}

449

else if(y_coord > (SRC_HEIGHT - 9))

450

{

451

// Avoid loading rows beyond the input height

452

const int end_z = SRC_HEIGHT - y_coord;

453

for(int i = 0; i < end_z; ++i)

454

{

455

CONVOLUTION1x9(values0, (src_addr + i * (int)src_stride_z), (weights_addr + i * (int)weights_stride_z));

}

}

else

{

CONVOLUTION1x9(values0, src_addr, weights_addr);

461

CONVOLUTION1x9(values0, (src_addr + 1 * (int)src_stride_z), (weights_addr + 1 * (int)weights_stride_z));

462

CONVOLUTION1x9(values0, (src_addr + 2 * (int)src_stride_z), (weights_addr + 2 * (int)weights_stride_z));

463

CONVOLUTION1x9(values0, (src_addr + 3 * (int)src_stride_z), (weights_addr + 3 * (int)weights_stride_z));

464

CONVOLUTION1x9(values0, (src_addr + 4 * (int)src_stride_z), (weights_addr + 4 * (int)weights_stride_z));

465

CONVOLUTION1x9(values0, (src_addr + 5 * (int)src_stride_z), (weights_addr + 5 * (int)weights_stride_z));

466

CONVOLUTION1x9(values0, (src_addr + 6 * (int)src_stride_z), (weights_addr + 6 * (int)weights_stride_z));

467

CONVOLUTION1x9(values0, (src_addr + 7 * (int)src_stride_z), (weights_addr + 7 * (int)weights_stride_z));

468

CONVOLUTION1x9(values0, (src_addr + 8 * (int)src_stride_z), (weights_addr + 8 * (int)weights_stride_z));

469

}

470

#elif KERNEL_SIZE == 5

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

471

#if(PAD_TOP == 1) || (PAD_BOTTM == 1)

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

472

if(y_coord < 0) // special case Z = -1 doesn't exists

473

{

474

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

475

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

476

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z));

477

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z));

478

}

479

else if(get_global_id(2) == (DST_HEIGHT - 1))

480

{

481

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z));

482

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

483

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

484

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z));

}

else

{

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z));

489

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

490

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

491

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z));

492

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z));

493

}

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

494

#elif(PAD_TOP == 2) || (PAD_BOTTM == 2)

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

495

if(y_coord < -1)

496

{

497

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

498

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z));

499

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z));

500

}

501

else if(y_coord == -1)

502

{

503

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

504

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

505

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z));

506

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z));

507

}

508

else if(y_coord == (SRC_HEIGHT - 3))

509

{

510

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z));

511

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

512

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

513

}

514

else if(y_coord >= (SRC_HEIGHT - 4))

515

{

516

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z));

517

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

518

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

519

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z));

}

else

{

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z));

524

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

525

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

526

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z));

527

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z));

528

}

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

529

#else /* PAD_TOP == 2 || || PAD_BOTTM == 2 */

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

530

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z));

531

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

532

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

533

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_z));

534

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_z));

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

535

#endif /* PAD_TOP == 1 || || PAD_BOTTM == 1 */

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

536

#elif KERNEL_SIZE == 3

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

537

#if(PAD_TOP > 0) || (PAD_BOTTOM > 0)

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

538

if(y_coord < 0) // special case Z = -1 doesn't exists

539

{

540

//skip first row and load the two next ones

541

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

542

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

543

}

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

544

else if(y_coord == (SRC_HEIGHT - PAD_BOTTOM - 1))

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

545

{

546

// special case when computing the last row of the output we must read the last three rows from the input buffer (including padding) but the

547

// Z axis has no padding at all.

548

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z));

549

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

}

else

{

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z));

554

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

555

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

556

}

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

557

#else // PAD_TOP > 0 || PAD_BOTTOM > 0

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

558

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_z));

559

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_z));

560

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_z), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_z));

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

561

#endif // PAD_TOP > 0 || PAD_BOTTOM > 0

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

562

#elif KERNEL_SIZE == 1

563

int weight = convert_int(*(__global DATA_TYPE *)weights_addr);

564

int8 input_value = convert_int8(INPUT_VALUE((__global DATA_TYPE *)src_addr, src_stride_y));

565

values0 += (input_value + input_offset) * ((int8)weight + weight_offset);

566

#endif /* (KERNEL_SIZE == 1) || (KERNEL_SIZE == 3) || (KERNEL_SIZE == 5) */

567

568

src_addr += src_stride_x;

569

weights_addr += weights_stride_x;

}

#ifdef HAS_BIAS

Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases);

Giorgio Arena

2020-03-02 09:49:29 +0000

[diff] [blame]

574

__global int *bias_addr = ((__global int *)(vector_offset(&biases, id0)));

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

575

values0 += (int8)(*bias_addr);

576

#endif /* defined(HAS_BIAS) */

577

578

#if OUTPUT_SHIFT < 0

579

values0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_GREATER_THAN_ONE(values0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8);

580

#else // OUTPUT_SHIFT < 0

581

values0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(values0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8);

582

#endif // OUTPUT_SHIFT < 0

583

values0 = values0 + output_offset;

584

585

VEC_DATA_TYPE(DATA_TYPE, 8)

586

values = CONVERT_SAT(values0, DATA_TYPE);

587

*(dst.ptr + 0 * dst_stride_y) = values.s0;

588

*(dst.ptr + 1 * dst_stride_y) = values.s1;

589

*(dst.ptr + 2 * dst_stride_y) = values.s2;

590

*(dst.ptr + 3 * dst_stride_y) = values.s3;

591

*(dst.ptr + 4 * dst_stride_y) = values.s4;

592

*(dst.ptr + 5 * dst_stride_y) = values.s5;

593

*(dst.ptr + 6 * dst_stride_y) = values.s6;

594

*(dst.ptr + 7 * dst_stride_y) = values.s7;

595

}

596

597

#else // defined(DATA_LAYOUT_NHWC)

598

Sang-Hoon Park

2019-10-15 09:29:13 +0100

[diff] [blame]

#if KERNEL_SIZE == 9

#if STRIDE_X == 1

#define CONVOLUTION1x9(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x9_STRIDE1(acc, src_row_ptr, weights_row_ptr)

603

#elif STRIDE_X == 2

604

#define CONVOLUTION1x9(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x9_STRIDE2(acc, src_row_ptr, weights_row_ptr)

605

#else /* STRIDE_X not equals 1 or 2 */

606

#error "STRIDE_X larger than 2 is not supported"

607

#endif /* STRIDE_X */

608

609

#define CONVOLUTION1x9_STRIDE1(acc, src_row_ptr, weights_row_ptr) \

610

({ \

611

int8 weights_values0 = convert_int8(vload8(0, weights_row_ptr)); \

612

int weights_value1 = convert_int(*(weights_row_ptr + 8)); \

613

int16 src0 = convert_int16(vload16(0, src_row_ptr)); \

614

acc += (src0.lo + input_offset) * ((int8)weights_values0.s0 + weight_offset); \

615

acc += ((int8)(src0.s1234, src0.s5678) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \

616

acc += ((int8)(src0.s2345, src0.s6789) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \

617

acc += ((int8)(src0.s3456, src0.s789A) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \

618

acc += ((int8)(src0.s4567, src0.s89AB) + input_offset) * ((int8)weights_values0.s4 + weight_offset); \

619

acc += ((int8)(src0.s5678, src0.s9ABC) + input_offset) * ((int8)weights_values0.s5 + weight_offset); \

620

acc += ((int8)(src0.s6789, src0.sABCD) + input_offset) * ((int8)weights_values0.s6 + weight_offset); \

621

acc += ((int8)(src0.s789A, src0.sBCDE) + input_offset) * ((int8)weights_values0.s7 + weight_offset); \

622

acc += ((int8)(src0.s89AB, src0.sCDEF) + input_offset) * ((int8)weights_value1 + weight_offset); \

623

})

624

625

#define CONVOLUTION1x9_STRIDE2(acc, src_row_ptr, weights_row_ptr) \

626

({ \

627

int8 weights_values0 = convert_int8(vload8(0, weights_row_ptr)); \

628

int weights_value1 = convert_int(*(weights_row_ptr + 8)); \

629

int16 src0 = convert_int16(vload16(0, src_row_ptr)); \

630

int8 src1 = convert_int8(vload8(0, src_row_ptr + 16)); \

631

acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \

632

acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \

633

acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \

634

acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \

635

acc += ((int8)(src0.s468A, src0.sCE, src1.s02) + input_offset) * ((int8)weights_values0.s4 + weight_offset); \

636

acc += ((int8)(src0.s579B, src0.sDF, src1.s13) + input_offset) * ((int8)weights_values0.s5 + weight_offset); \

637

acc += ((int8)(src0.s68AC, src0.sE, src1.s024) + input_offset) * ((int8)weights_values0.s6 + weight_offset); \

638

acc += ((int8)(src0.s79BD, src0.sF, src1.s135) + input_offset) * ((int8)weights_values0.s7 + weight_offset); \

639

acc += ((int8)(src0.s8ACE, src1.s0246) + input_offset) * ((int8)weights_value1 + weight_offset); \

640

})

641

642

#elif KERNEL_SIZE == 5

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

643

644

#if STRIDE_X == 1

645

#define CONVOLUTION1x5(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x5_STRIDE1(acc, src_row_ptr, weights_row_ptr)

646

#elif STRIDE_X == 2

647

#define CONVOLUTION1x5(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x5_STRIDE2(acc, src_row_ptr, weights_row_ptr)

648

#else /* STRIDE_X not equals 1 or 2 */

649

#error "STRIDE_X larger than 2 is not supported"

650

#endif /* STRIDE_X */

651

652

#define CONVOLUTION1x5_STRIDE1(acc, src_row_ptr, weights_row_ptr) \

653

({ \

654

int4 weights_values0 = convert_int4(vload4(0, weights_row_ptr)); \

655

int weights_value1 = convert_int(*(weights_row_ptr + 4)); \

656

int8 src0 = convert_int8(vload8(0, src_row_ptr)); \

657

int4 src1 = convert_int4(vload4(0, src_row_ptr + 8)); \

658

acc += (src0 + input_offset) * ((int8)weights_values0.s0 + weight_offset); \

659

acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \

660

acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \

661

acc += ((int8)(src0.s345, src0.s67, src1.s012) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \

662

acc += ((int8)(src0.s45, src0.s67, src1.s0123) + input_offset) * ((int8)weights_value1 + weight_offset); \

663

})

664

665

#define CONVOLUTION1x5_STRIDE2(acc, src_row_ptr, weights_row_ptr) \

666

({ \

667

int4 weights_values0 = convert_int4(vload4(0, weights_row_ptr)); \

668

int weights_value1 = convert_int(*(weights_row_ptr + 4)); \

669

int16 src0 = convert_int16(vload16(0, src_row_ptr)); \

670

int4 src1 = convert_int4(vload4(0, src_row_ptr + 16)); \

671

acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \

672

acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \

673

acc += ((int8)(src0.s2468, src0.sACE, src1.s0) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \

674

acc += ((int8)(src0.s3579, src0.sBDF, src1.s1) + input_offset) * ((int8)weights_values0.s3 + weight_offset); \

675

acc += ((int8)(src0.s468a, src0.sCE, src1.s02) + input_offset) * ((int8)weights_value1 + weight_offset); \

676

})

677

678

#elif KERNEL_SIZE == 3

679

680

#if STRIDE_X == 1

681

#define CONVOLUTION1x3(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x3_STRIDE1(acc, src_row_ptr, weights_row_ptr)

682

#elif STRIDE_X == 2

683

#define CONVOLUTION1x3(acc, src_row_ptr, weights_row_ptr) CONVOLUTION1x3_STRIDE2(acc, src_row_ptr, weights_row_ptr)

684

#else /* STRIDE_X not equals 1 or 2 */

685

#error "STRIDE_X larger than 2 is not supported"

686

#endif /* STRIDE_X */

687

688

#define CONVOLUTION1x3_STRIDE1(acc, src_row_ptr, weights_row_ptr) \

689

({ \

690

int3 weights_values0 = convert_int3(vload3(0, weights_row_ptr)); \

691

int8 src0 = convert_int8(vload8(0, src_row_ptr)); \

692

int2 src1 = convert_int2(vload2(0, src_row_ptr + 8)); \

693

acc += (src0 + input_offset) * ((int8)weights_values0.s0 + weight_offset); \

694

acc += ((int8)(src0.s1234, src0.s567, src1.s0) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \

695

acc += ((int8)(src0.s234, src0.s567, src1.s01) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \

696

})

697

698

#define CONVOLUTION1x3_STRIDE2(acc, src_row_ptr, weights_row_ptr) \

699

({ \

700

int3 weights_values0 = convert_int3(vload3(0, weights_row_ptr)); \

701

int16 src0 = convert_int16(vload16(0, src_row_ptr)); \

702

int src1 = convert_int(*(src_row_ptr + 16)); \

703

acc += (src0.even + input_offset) * ((int8)weights_values0.s0 + weight_offset); \

704

acc += ((int8)(src0.s1357, src0.s9BDF) + input_offset) * ((int8)weights_values0.s1 + weight_offset); \

705

acc += ((int8)(src0.s2468, src0.sACE, src1) + input_offset) * ((int8)weights_values0.s2 + weight_offset); \

706

})

707

708

#elif KERNEL_SIZE == 1

709

710

#if STRIDE_X == 3

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

711

#define INPUT_VALUE extract_input_stride3

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

712

#elif STRIDE_X == 2

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

713

#define INPUT_VALUE extract_input_stride2

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

714

#elif STRIDE_X == 1

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

715

#define INPUT_VALUE extract_input_stride1

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

716

717

#else /* STRIDE_X not equals 1, 2 or 3 */

718

#error "Only support strides 1, 2 and 3"

719

#endif /* STRIDE_X */

720

721

/** Extracts a 1D horizontal vector from the input tensor with stride as 1.

722

*

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

723

* @param[in] input_value Pointer to the first value.

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

724

*

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

725

* @return extracted input values.

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

726

*/

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

727

inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride1(__global const DATA_TYPE *input_value)

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

728

{

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

729

return vload8(0, input_value);

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

730

}

731

732

/** Extracts a 1D horizontal vector from the input tensor with stride as 2.

733

*

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

734

* @param[in] input_value Pointer to the first value.

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

735

*

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

736

* @return extracted input values.

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

737

*/

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

738

inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride2(__global const DATA_TYPE *input_value)

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

739

{

Sheri Zhang

2020-02-20 11:23:08 +0000

[diff] [blame]

740

VEC_DATA_TYPE(DATA_TYPE, 16)

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

741

temp = vload16(0, input_value);

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

742

return temp.s02468ace;

743

}

744

745

/** Extracts a 1D horizontal vector from the input tensor with stride as 3 and 8-bit data size.

746

*

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

747

* @param[in] input_value Pointer to the first value.

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

748

*

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

749

* @return extracted input values.

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

750

*/

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

751

inline VEC_DATA_TYPE(DATA_TYPE, 8) extract_input_stride3(__global const DATA_TYPE *input_value)

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

752

{

Sheri Zhang

2020-02-20 11:23:08 +0000

[diff] [blame]

753

VEC_DATA_TYPE(DATA_TYPE, 16)

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

754

temp1 = vload16(0, input_value);

Sheri Zhang

2020-02-20 11:23:08 +0000

[diff] [blame]

755

VEC_DATA_TYPE(DATA_TYPE, 16)

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

756

temp2 = vload16(0, input_value + 12);

Sheri Zhang

2020-02-20 11:23:08 +0000

[diff] [blame]

757

return (VEC_DATA_TYPE(DATA_TYPE, 8))(temp1.s0369, temp2.s0369);

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

758

}

759

Sang-Hoon Park

2019-10-15 09:29:13 +0100

[diff] [blame]

760

#else /* KERNEL_SIZE not equals 1, 3 , 5, 9 */

761

#error "Only kernel sizes 1, 3, 5 and 9 are supported"

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

762

#endif /* KERNEL_SIZE */

763

764

/** This kernel performs a direct convolution to convolve the low three dimensions.

765

*

766

* @note The convolution stride x must be passed at compile time using -DSTRIDE_X e.g. -DSTRIDE_X=1

767

* @note The third dimensions of the weights tensors must be passed at compile time using -DWEIGHTS_DEPTH

768

* @note If biases are used then -DHAS_BIAS has to be passed at compile time

Michele Di Giorgio

2019-10-23 10:53:10 +0100

[diff] [blame]

769

* @note The output quantization multiplier must be passed at compile time using -DOUTPUT_MULTIPLIER e.g. -DOUTPUT_MULTIPLIER=1234

770

* @note The output quantization shift must be passed at compile time using -DOUTPUT_SHIFT e.g. -DOUTPUT_SHIFT=4

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

771

*

Sheri Zhang

2020-02-20 11:23:08 +0000

[diff] [blame]

772

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

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

773

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

774

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

775

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

776

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

777

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

778

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

779

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

780

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

781

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

782

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

783

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

784

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

785

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

786

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

787

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

Sheri Zhang

2020-02-20 11:23:08 +0000

[diff] [blame]

788

* @param[in] weights_ptr Pointer to the weights tensor. Supported data types: same as @p src_ptr

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

789

* @param[in] weights_stride_x Stride of the weights tensor in X dimension (in bytes)

790

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

791

* @param[in] weights_stride_y Stride of the weights tensor in Y dimension (in bytes)

792

* @param[in] weights_step_y weights_stride_y * number of elements along y processed per workitem(in bytes)

793

* @param[in] weights_stride_z Stride of the weights tensor in Z dimension (in bytes)

794

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

795

* @param[in] weights_offset_first_element_in_bytes The offset of the first element in the weights tensor

Georgios Pinitas

540d008

2017-11-17 10:55:00 +0000

[diff] [blame]

796

* @param[in] biases_ptr Pointer to the biases tensor. Supported data types: S32

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

797

* @param[in] biases_stride_x Stride of the biases tensor in X dimension (in bytes)

798

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

799

* @param[in] biases_offset_first_element_in_bytes The offset of the first element in the biases tensor

800

* @param[in] weights_stride_w Stride of the weights tensor in the 4th dimension

801

* @param[in] input_offset Input offset quantization parameter

802

* @param[in] weight_offset Weights offset quantization parameter

803

* @param[in] output_offset Output offset quantization parameter

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

804

*/

Sang-Hoon Park

2019-10-15 09:29:13 +0100

[diff] [blame]

805

__kernel void direct_convolution_quantized(

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

806

TENSOR3D_DECLARATION(src),

807

TENSOR3D_DECLARATION(dst),

808

TENSOR3D_DECLARATION(weights),

809

#ifdef HAS_BIAS

810

VECTOR_DECLARATION(biases),

811

#endif /* defined(HAS_BIAS) */

812

unsigned int weights_stride_w,

813

int input_offset,

814

int weight_offset,

Michele Di Giorgio

2019-10-23 10:53:10 +0100

[diff] [blame]

815

int output_offset)

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

816

{

817

Image src = CONVERT_TO_IMAGE_STRUCT(src);

818

Tensor3D weights = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(weights);

819

Tensor3D dst = CONVERT_TO_TENSOR3D_STRUCT(dst);

820

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

821

int8 values0 = 0;

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

822

Sheri Zhang

2020-02-20 11:23:08 +0000

[diff] [blame]

823

__global DATA_TYPE *weights_addr = (__global DATA_TYPE *)tensor3D_offset(&weights, 0, 0, 0);

824

__global DATA_TYPE *src_addr = (__global DATA_TYPE *)offset(&src, 0, 0);

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

825

826

const int kernel_index = get_global_id(2);

827

weights_addr += kernel_index * weights_stride_w;

828

829

for(volatile int d = 0; d < WEIGHTS_DEPTH; ++d)

830

{

Sang-Hoon Park

2019-10-15 09:29:13 +0100

[diff] [blame]

831

#if KERNEL_SIZE == 9

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

832

CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_y));

833

CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y));

834

CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_y));

835

CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_y));

836

CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_y));

837

CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 5 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 5 * weights_stride_y));

838

CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 6 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 6 * weights_stride_y));

839

CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 7 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 7 * weights_stride_y));

840

CONVOLUTION1x9(values0, (__global DATA_TYPE *)(src_addr + 8 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 8 * weights_stride_y));

Sang-Hoon Park

2019-10-15 09:29:13 +0100

[diff] [blame]

841

#elif KERNEL_SIZE == 5

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

842

CONVOLUTION1x5(values0, (__global DATA_TYPE *)src_addr, (__global DATA_TYPE *)weights_addr);

843

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y));

844

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_y));

845

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 3 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 3 * weights_stride_y));

846

CONVOLUTION1x5(values0, (__global DATA_TYPE *)(src_addr + 4 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 4 * weights_stride_y));

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

847

#elif KERNEL_SIZE == 3

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

848

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 0 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 0 * weights_stride_y));

849

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 1 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 1 * weights_stride_y));

850

CONVOLUTION1x3(values0, (__global DATA_TYPE *)(src_addr + 2 * src_stride_y), (__global DATA_TYPE *)(weights_addr + 2 * weights_stride_y));

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

851

#elif KERNEL_SIZE == 1

Sheri Zhang

2020-02-20 11:23:08 +0000

[diff] [blame]

852

int weight = convert_int(*(__global DATA_TYPE *)weights_addr);

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

853

int8 input_value = convert_int8(INPUT_VALUE((__global DATA_TYPE *)src_addr));

854

values0 += (input_value + input_offset) * ((int8)weight + weight_offset);

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

855

#endif /* (KERNEL_SIZE == 1) || (KERNEL_SIZE == 3) || (KERNEL_SIZE == 5) */

856

857

src_addr += src_stride_z;

858

weights_addr += weights_stride_z;

859

}

860

861

#ifdef HAS_BIAS

Georgios Pinitas

540d008

2017-11-17 10:55:00 +0000

[diff] [blame]

862

Vector biases = CONVERT_TO_VECTOR_STRUCT_NO_STEP(biases);

863

__global int *bias_addr = ((__global int *)(vector_offset(&biases, kernel_index)));

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

864

values0 += (int8)(*bias_addr);

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

865

#endif /* defined(HAS_BIAS) */

866

Michele Di Giorgio

2019-10-23 10:53:10 +0100

[diff] [blame]

867

#if OUTPUT_SHIFT < 0

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

868

values0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_GREATER_THAN_ONE(values0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8);

Michele Di Giorgio

2019-10-23 10:53:10 +0100

[diff] [blame]

869

#else // OUTPUT_SHIFT < 0

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

870

values0 = ASYMM_MULT_BY_QUANT_MULTIPLIER_LESS_THAN_ONE(values0, OUTPUT_MULTIPLIER, OUTPUT_SHIFT, 8);

Michele Di Giorgio

2019-10-23 10:53:10 +0100

[diff] [blame]

871

#endif // OUTPUT_SHIFT < 0

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

872

values0 = values0 + output_offset;

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

873

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

874

vstore8(CONVERT_SAT(values0, DATA_TYPE), 0, (__global DATA_TYPE *)dst.ptr);

Chunosov

2017-11-03 17:33:15 +0700

[diff] [blame]

875

}

Giorgio Arena

2020-02-17 16:33:20 +0000

[diff] [blame]

876

877

#endif // defined(DATA_LAYOUT_NHWC)

Michele Di Giorgio