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

Anthony Barbier

6ff3b19

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

[diff] [blame]

1

/*

Ioan-Cristian Szabo

ae3c8ab

2017-11-16 17:55:03 +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"

/** This function combines three planes to a single RGB image.

27

*

28

* @param[in] plane0_ptr Pointer to the first plane. Supported Format: U8

29

* @param[in] plane0_stride_x Stride of the first plane in X dimension (in bytes)

30

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

31

* @param[in] plane0_stride_y Stride of the first plane in Y dimension (in bytes)

32

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

33

* @param[in] plane0_offset_first_element_in_bytes The offset of the first element in the first plane

34

* @param[in] plane1_ptr Pointer to the second plane. Supported Format: U8

35

* @param[in] plane1_stride_x Stride of the second plane in X dimension (in bytes)

36

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

37

* @param[in] plane1_stride_y Stride of the second plane in Y dimension (in bytes)

38

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

39

* @param[in] plane1_offset_first_element_in_bytes The offset of the first element in the second plane

40

* @param[in] plane2_ptr Pointer to the third plane. Supported Format: U8

41

* @param[in] plane2_stride_x Stride of the third plane in X dimension (in bytes)

42

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

43

* @param[in] plane2_stride_y Stride of the third plane in Y dimension (in bytes)

44

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

45

* @param[in] plane2_offset_first_element_in_bytes The offset of the first element in the third plane

46

* @param[in] dst_ptr Pointer to the destination image. Supported Format: RGB

47

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

48

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

49

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

50

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

51

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

52

*/

53

__kernel void channel_combine_RGB888(

54

IMAGE_DECLARATION(plane0),

55

IMAGE_DECLARATION(plane1),

56

IMAGE_DECLARATION(plane2),

57

IMAGE_DECLARATION(dst))

58

{

59

// Get pixels pointer

60

Image plane0 = CONVERT_TO_IMAGE_STRUCT(plane0);

61

Image plane1 = CONVERT_TO_IMAGE_STRUCT(plane1);

62

Image plane2 = CONVERT_TO_IMAGE_STRUCT(plane2);

63

Image dst = CONVERT_TO_IMAGE_STRUCT(dst);

64

65

uchar16 data0 = vload16(0, plane0.ptr);

66

uchar16 data1 = vload16(0, plane1.ptr);

67

uchar16 data2 = vload16(0, plane2.ptr);

68

69

uchar16 out0 = (uchar16)(data0.s0, data1.s0, data2.s0,

70

data0.s1, data1.s1, data2.s1,

71

data0.s2, data1.s2, data2.s2,

72

data0.s3, data1.s3, data2.s3,

73

data0.s4, data1.s4, data2.s4,

74

data0.s5);

75

vstore16(out0, 0, dst.ptr);

76

77

uchar16 out1 = (uchar16)(data1.s5, data2.s5, data0.s6,

78

data1.s6, data2.s6, data0.s7,

79

data1.s7, data2.s7, data0.s8,

80

data1.s8, data2.s8, data0.s9,

81

data1.s9, data2.s9, data0.sA,

82

data1.sA);

83

vstore16(out1, 0, dst.ptr + 16);

84

85

uchar16 out2 = (uchar16)(data2.sA, data0.sB, data1.sB,

86

data2.sB, data0.sC, data1.sC,

87

data2.sC, data0.sD, data1.sD,

88

data2.sD, data0.sE, data1.sE,

89

data2.sE, data0.sF, data1.sF,

90

data2.sF);

91

vstore16(out2, 0, dst.ptr + 32);

92

}

93

94

/** This function combines three planes to a single RGBA image.

95

*

96

* @param[in] plane0_ptr Pointer to the first plane. Supported Format: U8

97

* @param[in] plane0_stride_x Stride of the first plane in X dimension (in bytes)

98

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

99

* @param[in] plane0_stride_y Stride of the first plane in Y dimension (in bytes)

100

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

101

* @param[in] plane0_offset_first_element_in_bytes The offset of the first element in the first plane

102

* @param[in] plane1_ptr Pointer to the second plane. Supported Format: U8

103

* @param[in] plane1_stride_x Stride of the second plane in X dimension (in bytes)

104

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

105

* @param[in] plane1_stride_y Stride of the second plane in Y dimension (in bytes)

106

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

107

* @param[in] plane1_offset_first_element_in_bytes The offset of the first element in the second plane

108

* @param[in] plane2_ptr Pointer to the third plane. Supported Format: U8

109

* @param[in] plane2_stride_x Stride of the third plane in X dimension (in bytes)

110

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

111

* @param[in] plane2_stride_y Stride of the third plane in Y dimension (in bytes)

112

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

113

* @param[in] plane2_offset_first_element_in_bytes The offset of the first element in the third plane

114

* @param[in] plane3_ptr Pointer to the fourth plane. Supported Format: U8

115

* @param[in] plane3_stride_x Stride of the fourth plane in X dimension (in bytes)

116

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

117

* @param[in] plane3_stride_y Stride of the fourth plane in Y dimension (in bytes)

118

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

119

* @param[in] plane3_offset_first_element_in_bytes The offset of the first element in the fourth plane

120

* @param[in] dst_ptr Pointer to the destination image. Supported Format: RGBA

121

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

122

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

123

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

124

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

125

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

126

*/

127

__kernel void channel_combine_RGBA8888(

128

IMAGE_DECLARATION(plane0),

129

IMAGE_DECLARATION(plane1),

130

IMAGE_DECLARATION(plane2),

131

IMAGE_DECLARATION(plane3),

132

IMAGE_DECLARATION(dst))

133

{

134

// Get pixels pointer

135

Image plane0 = CONVERT_TO_IMAGE_STRUCT(plane0);

136

Image plane1 = CONVERT_TO_IMAGE_STRUCT(plane1);

137

Image plane2 = CONVERT_TO_IMAGE_STRUCT(plane2);

138

Image plane3 = CONVERT_TO_IMAGE_STRUCT(plane3);

139

Image dst = CONVERT_TO_IMAGE_STRUCT(dst);

140

141

uchar16 data0 = vload16(0, plane0.ptr);

142

uchar16 data1 = vload16(0, plane1.ptr);

143

uchar16 data2 = vload16(0, plane2.ptr);

144

uchar16 data3 = vload16(0, plane3.ptr);

145

146

uchar16 out0 = (uchar16)(data0.s0, data1.s0, data2.s0, data3.s0,

147

data0.s1, data1.s1, data2.s1, data3.s1,

148

data0.s2, data1.s2, data2.s2, data3.s2,

149

data0.s3, data1.s3, data2.s3, data3.s3);

150

vstore16(out0, 0, dst.ptr);

151

152

uchar16 out1 = (uchar16)(data0.s4, data1.s4, data2.s4, data3.s4,

153

data0.s5, data1.s5, data2.s5, data3.s5,

154

data0.s6, data1.s6, data2.s6, data3.s6,

155

data0.s7, data1.s7, data2.s7, data3.s7);

156

vstore16(out1, 0, dst.ptr + 16);

157

158

uchar16 out2 = (uchar16)(data0.s8, data1.s8, data2.s8, data3.s8,

159

data0.s9, data1.s9, data2.s9, data3.s9,

160

data0.sA, data1.sA, data2.sA, data3.sA,

161

data0.sB, data1.sB, data2.sB, data3.sB);

162

vstore16(out2, 0, dst.ptr + 32);

163

164

uchar16 out3 = (uchar16)(data0.sC, data1.sC, data2.sC, data3.sC,

165

data0.sD, data1.sD, data2.sD, data3.sD,

166

data0.sE, data1.sE, data2.sE, data3.sE,

167

data0.sF, data1.sF, data2.sF, data3.sF);

168

vstore16(out3, 0, dst.ptr + 48);

169

}

170

171

/** This function combines three planes to a single YUYV image.

172

*

173

* @param[in] plane0_ptr Pointer to the first plane. Supported Format: U8

174

* @param[in] plane0_stride_x Stride of the first plane in X dimension (in bytes)

175

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

176

* @param[in] plane0_stride_y Stride of the first plane in Y dimension (in bytes)

177

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

178

* @param[in] plane0_offset_first_element_in_bytes The offset of the first element in the first plane

179

* @param[in] plane1_ptr Pointer to the second plane. Supported Format: U8

180

* @param[in] plane1_stride_x Stride of the second plane in X dimension (in bytes)

181

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

182

* @param[in] plane1_stride_y Stride of the second plane in Y dimension (in bytes)

183

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

184

* @param[in] plane1_offset_first_element_in_bytes The offset of the first element in the second plane

185

* @param[in] plane2_ptr Pointer to the third plane. Supported Format: U8

186

* @param[in] plane2_stride_x Stride of the third plane in X dimension (in bytes)

187

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

188

* @param[in] plane2_stride_y Stride of the third plane in Y dimension (in bytes)

189

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

190

* @param[in] plane2_offset_first_element_in_bytes The offset of the first element in the third plane

191

* @param[in] dst_ptr Pointer to the destination image. Supported Format: YUYV

192

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

193

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

194

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

195

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

196

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

197

*/

198

__kernel void channel_combine_YUYV422(

199

IMAGE_DECLARATION(plane0),

200

IMAGE_DECLARATION(plane1),

201

IMAGE_DECLARATION(plane2),

202

IMAGE_DECLARATION(dst))

203

{

204

// Get pixels pointer

205

Image plane0 = CONVERT_TO_IMAGE_STRUCT(plane0);

206

Image plane1 = CONVERT_TO_IMAGE_STRUCT(plane1);

207

Image plane2 = CONVERT_TO_IMAGE_STRUCT(plane2);

208

Image dst = CONVERT_TO_IMAGE_STRUCT(dst);

209

210

uchar16 data0 = vload16(0, plane0.ptr);

211

uchar8 data1 = vload8(0, plane1.ptr);

212

uchar8 data2 = vload8(0, plane2.ptr);

213

214

uchar16 out0 = (uchar16)(data0.s0, data1.s0, data0.s1, data2.s0,

215

data0.s2, data1.s1, data0.s3, data2.s1,

216

data0.s4, data1.s2, data0.s5, data2.s2,

217

data0.s6, data1.s3, data0.s7, data2.s3);

218

vstore16(out0, 0, dst.ptr);

219

uchar16 out1 = (uchar16)(data0.s8, data1.s4, data0.s9, data2.s4,

220

data0.sA, data1.s5, data0.sB, data2.s5,

221

data0.sC, data1.s6, data0.sD, data2.s6,

222

data0.sE, data1.s7, data0.sF, data2.s7);

223

vstore16(out1, 0, dst.ptr + 16);

224

}

225

226

/** This function combines three planes to a single UYUV image.

227

*

228

* @param[in] plane0_ptr Pointer to the first plane. Supported Format: U8

229

* @param[in] plane0_stride_x Stride of the first plane in X dimension (in bytes)

230

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

231

* @param[in] plane0_stride_y Stride of the first plane in Y dimension (in bytes)

232

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

233

* @param[in] plane0_offset_first_element_in_bytes The offset of the first element in the first plane

234

* @param[in] plane1_ptr Pointer to the second plane. Supported Format: U8

235

* @param[in] plane1_stride_x Stride of the second plane in X dimension (in bytes)

236

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

237

* @param[in] plane1_stride_y Stride of the second plane in Y dimension (in bytes)

238

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

239

* @param[in] plane1_offset_first_element_in_bytes The offset of the first element in the second plane

240

* @param[in] plane2_ptr Pointer to the third plane. Supported Format: U8

241

* @param[in] plane2_stride_x Stride of the third plane in X dimension (in bytes)

242

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

243

* @param[in] plane2_stride_y Stride of the third plane in Y dimension (in bytes)

244

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

245

* @param[in] plane2_offset_first_element_in_bytes The offset of the first element in the third plane

246

* @param[in] dst_ptr Pointer to the destination image. Supported Format: UYUV

247

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

248

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

249

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

250

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

251

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

252

*/

253

__kernel void channel_combine_UYVY422(

254

IMAGE_DECLARATION(plane0),

255

IMAGE_DECLARATION(plane1),

256

IMAGE_DECLARATION(plane2),

257

IMAGE_DECLARATION(dst))

258

{

259

// Get pixels pointer

260

Image plane0 = CONVERT_TO_IMAGE_STRUCT(plane0);

261

Image plane1 = CONVERT_TO_IMAGE_STRUCT(plane1);

262

Image plane2 = CONVERT_TO_IMAGE_STRUCT(plane2);

263

Image dst = CONVERT_TO_IMAGE_STRUCT(dst);

264

265

uchar16 data0 = vload16(0, plane0.ptr);

266

uchar8 data1 = vload8(0, plane1.ptr);

267

uchar8 data2 = vload8(0, plane2.ptr);

268

269

uchar16 out0 = (uchar16)(data1.s0, data0.s0, data2.s0, data0.s1,

270

data1.s1, data0.s2, data2.s1, data0.s3,

271

data1.s2, data0.s4, data2.s2, data0.s5,

272

data1.s3, data0.s6, data2.s3, data0.s7);

273

vstore16(out0, 0, dst.ptr);

274

uchar16 out1 = (uchar16)(data1.s4, data0.s8, data2.s4, data0.s9,

275

data1.s5, data0.sA, data2.s5, data0.sB,

276

data1.s6, data0.sC, data2.s6, data0.sD,

277

data1.s7, data0.sE, data2.s7, data0.sF);

278

vstore16(out1, 0, dst.ptr + 16);

279

}

280

281

/** This function combines three planes to a single NV12/NV21 image.

282

*

283

* @note NV12 or NV21 has to be specified through preprocessor macro. eg. -DNV12 performs NV12 channel combine.

284

*

285

* @param[in] src_plane0_ptr Pointer to the first plane. Supported Format: U8

286

* @param[in] src_plane0_stride_x Stride of the first plane in X dimension (in bytes)

287

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

288

* @param[in] src_plane0_stride_y Stride of the first plane in Y dimension (in bytes)

289

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

290

* @param[in] src_plane0_offset_first_element_in_bytes The offset of the first element in the first plane

291

* @param[in] src_plane1_ptr Pointer to the second plane. Supported Format: U8

292

* @param[in] src_plane1_stride_x Stride of the second plane in X dimension (in bytes)

293

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

294

* @param[in] src_plane1_stride_y Stride of the second plane in Y dimension (in bytes)

295

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

296

* @param[in] src_plane1_offset_first_element_in_bytes The offset of the first element in the second plane

297

* @param[in] src_plane2_ptr Pointer to the third plane. Supported Format: U8

298

* @param[in] src_plane2_stride_x Stride of the third plane in X dimension (in bytes)

299

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

300

* @param[in] src_plane2_stride_y Stride of the third plane in Y dimension (in bytes)

301

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

302

* @param[in] src_plane2_offset_first_element_in_bytes The offset of the first element in the third plane

303

* @param[in] dst_plane0_ptr Pointer to the first plane of the destination image. Supported Format: U8

304

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

305

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

306

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

307

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

308

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

309

* @param[in] dst_plane1_ptr Pointer to the second plane of the destination image. Supported Format: UV88

310

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

311

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

312

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

313

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

314

* @param[in] dst_plane1_offset_first_element_in_bytes The offset of the first element in the second plane of the destination image

315

* @param[in] height Sub-sampled height

316

*/

317

__kernel void channel_combine_NV(

318

IMAGE_DECLARATION(src_plane0),

319

IMAGE_DECLARATION(src_plane1),

320

IMAGE_DECLARATION(src_plane2),

321

IMAGE_DECLARATION(dst_plane0),

322

IMAGE_DECLARATION(dst_plane1),

323

uint height)

324

{

325

// Get pixels pointer

326

Image src_plane0 = CONVERT_TO_IMAGE_STRUCT(src_plane0);

327

Image src_plane1 = CONVERT_TO_IMAGE_STRUCT(src_plane1);

328

Image src_plane2 = CONVERT_TO_IMAGE_STRUCT(src_plane2);

329

Image dst_plane0 = CONVERT_TO_IMAGE_STRUCT(dst_plane0);

330

Image dst_plane1 = CONVERT_TO_IMAGE_STRUCT(dst_plane1);

331

332

// Copy plane data

333

vstore16(vload16(0, src_plane0.ptr), 0, dst_plane0.ptr);

334

vstore16(vload16(0, offset(&src_plane0, 0, height)), 0, (__global uchar *)offset(&dst_plane0, 0, height));

335

336

// Create UV place

337

uchar8 data1 = vload8(0, src_plane1.ptr);

338

uchar8 data2 = vload8(0, src_plane2.ptr);

339

Anthony Barbier

ac69aa1

2017-07-03 17:39:37 +0100

[diff] [blame]

340

#ifdef NV12

Ioan-Cristian Szabo

ae3c8ab

2017-11-16 17:55:03 +0000

[diff] [blame]

341

vstore16(shuffle2(data1, data2, (uchar16)(0, 8, 1, 9, 2, 10, 3, 11, 4, 12, 5, 13, 6, 14, 7, 15)), 0, dst_plane1.ptr);

Anthony Barbier

ac69aa1

2017-07-03 17:39:37 +0100

[diff] [blame]

342

#elif defined(NV21)

Ioan-Cristian Szabo

ae3c8ab

2017-11-16 17:55:03 +0000

[diff] [blame]

343

vstore16(shuffle2(data2, data1, (uchar16)(0, 8, 1, 9, 2, 10, 3, 11, 4, 12, 5, 13, 6, 14, 7, 15)), 0, dst_plane1.ptr);

Anthony Barbier

ac69aa1

2017-07-03 17:39:37 +0100

[diff] [blame]

344

#endif /* NV12 or NV21 */

Anthony Barbier

6ff3b19

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

[diff] [blame]

345

}

346

347

/** This function combines three planes to a single YUV444 or IYUV image.

348

*

349

* @note YUV444 or IYUV has to be specified through preprocessor macro. eg. -DIYUV performs IYUV channel combine.

350

*

351

* @param[in] src_plane0_ptr Pointer to the first plane. Supported Format: U8

352

* @param[in] src_plane0_stride_x Stride of the first plane in X dimension (in bytes)

353

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

354

* @param[in] src_plane0_stride_y Stride of the first plane in Y dimension (in bytes)

355

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

356

* @param[in] src_plane0_offset_first_element_in_bytes The offset of the first element in the first plane

357

* @param[in] src_plane1_ptr Pointer to the second plane. Supported Format: U8

358

* @param[in] src_plane1_stride_x Stride of the second plane in X dimension (in bytes)

359

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

360

* @param[in] src_plane1_stride_y Stride of the second plane in Y dimension (in bytes)

361

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

362

* @param[in] src_plane1_offset_first_element_in_bytes The offset of the first element in the second plane

363

* @param[in] src_plane2_ptr Pointer to the third plane. Supported Format: U8

364

* @param[in] src_plane2_stride_x Stride of the third plane in X dimension (in bytes)

365

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

366

* @param[in] src_plane2_stride_y Stride of the third plane in Y dimension (in bytes)

367

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

368

* @param[in] src_plane2_offset_first_element_in_bytes The offset of the first element in the third plane

369

* @param[in] dst_plane0_ptr Pointer to the first plane of the destination image. Supported Format: U8

370

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

371

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

372

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

373

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

374

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

375

* @param[in] dst_plane1_ptr Pointer to the second plane of the destination image. Supported Format: U8

376

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

377

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

378

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

379

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

380

* @param[in] dst_plane1_offset_first_element_in_bytes The offset of the first element in the second plane of the destination image

381

* @param[in] dst_plane2_ptr Pointer to the third plane of the destination image. Supported Format: U8

382

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

383

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

384

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

385

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

386

* @param[in] dst_plane2_offset_first_element_in_bytes The offset of the first element in the third plane of the destination image

387

* @param[in] height Sub-sampled height

388

*/

389

__kernel void copy_planes_3p(

390

IMAGE_DECLARATION(src_plane0),

391

IMAGE_DECLARATION(src_plane1),

392

IMAGE_DECLARATION(src_plane2),

393

IMAGE_DECLARATION(dst_plane0),

394

IMAGE_DECLARATION(dst_plane1),

395

IMAGE_DECLARATION(dst_plane2),

396

uint height)

397

{

398

// Get pixels pointer

399

Image src_plane0 = CONVERT_TO_IMAGE_STRUCT(src_plane0);

400

Image src_plane1 = CONVERT_TO_IMAGE_STRUCT(src_plane1);

401

Image src_plane2 = CONVERT_TO_IMAGE_STRUCT(src_plane2);

402

Image dst_plane0 = CONVERT_TO_IMAGE_STRUCT(dst_plane0);

403

Image dst_plane1 = CONVERT_TO_IMAGE_STRUCT(dst_plane1);

404

Image dst_plane2 = CONVERT_TO_IMAGE_STRUCT(dst_plane2);

405

406

// Copy plane data

407

vstore16(vload16(0, src_plane0.ptr), 0, dst_plane0.ptr);

Anthony Barbier

ac69aa1

2017-07-03 17:39:37 +0100

[diff] [blame]

408

#ifdef YUV444

Anthony Barbier

6ff3b19

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

[diff] [blame]

409

vstore16(vload16(0, src_plane1.ptr), 0, dst_plane1.ptr);

410

vstore16(vload16(0, src_plane2.ptr), 0, dst_plane2.ptr);

Anthony Barbier

ac69aa1

2017-07-03 17:39:37 +0100

[diff] [blame]

411

#elif defined(IYUV)

Anthony Barbier

6ff3b19

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

[diff] [blame]

412

vstore16(vload16(0, offset(&src_plane0, 0, height)), 0, (__global uchar *)offset(&dst_plane0, 0, height));

413

vstore8(vload8(0, src_plane1.ptr), 0, dst_plane1.ptr);

414

vstore8(vload8(0, src_plane2.ptr), 0, dst_plane2.ptr);

Anthony Barbier

ac69aa1

2017-07-03 17:39:37 +0100

[diff] [blame]

415

#endif /* YUV444 or IYUV */

Anthony Barbier

6ff3b19

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

[diff] [blame]

416

}