FFmpeg: libavfilter/unsharp_opencl.c Source File

FFmpeg

[フレーム]

libavfilter

unsharp_opencl.c

Go to the documentation of this file.

1 /*

4 *

5 * This file is part of FFmpeg.

6 *

7 * FFmpeg is free software; you can redistribute it and/or

8 * modify it under the terms of the GNU Lesser General Public

9 * License as published by the Free Software Foundation; either

10 * version 2.1 of the License, or (at your option) any later version.

11 *

12 * FFmpeg is distributed in the hope that it will be useful,

13 * but WITHOUT ANY WARRANTY; without even the implied warranty of

14 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU

15 * Lesser General Public License for more details.

16 *

17 * You should have received a copy of the GNU Lesser General Public

18 * License along with FFmpeg; if not, write to the Free Software

19 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA

20 */

22 /**

23 * @file

24 * unsharp input video

25 */

27 #include "unsharp_opencl.h"

28 #include "libavutil/common.h"

29 #include "libavutil/opencl_internal.h"

31 #define PLANE_NUM 3

32 #define ROUND_TO_16(a) (((((a) - 1)/16)+1)*16)

34 static inline void add_mask_counter(uint32_t *dst, uint32_t *counter1, uint32_t *counter2, int len)

35 {

36 int i;

37 for (i = 0; i < len; i++) {

38 dst[i] = counter1[i] + counter2[i];

39 }

40 }

42 static int compute_mask(int step, uint32_t *mask)

43 {

44 int i, z, ret = 0;

45 int counter_size = sizeof(uint32_t) * (2 * step + 1);

46 uint32_t *temp1_counter, *temp2_counter, **counter;

47 temp1_counter = av_mallocz(counter_size);

48 if (!temp1_counter) {

49 ret = AVERROR(ENOMEM);

50 goto end;

51 }

52 temp2_counter = av_mallocz(counter_size);

53 if (!temp2_counter) {

54 ret = AVERROR(ENOMEM);

55 goto end;

56 }

57 counter = av_mallocz_array(2 * step + 1, sizeof(uint32_t *));

58 if (!counter) {

59 ret = AVERROR(ENOMEM);

60 goto end;

61 }

62 for (i = 0; i < 2 * step + 1; i++) {

63 counter[i] = av_mallocz(counter_size);

64 if (!counter[i]) {

65 ret = AVERROR(ENOMEM);

66 goto end;

67 }

68 }

69 for (i = 0; i < 2 * step + 1; i++) {

70 memset(temp1_counter, 0, counter_size);

71 temp1_counter[i] = 1;

72 for (z = 0; z < step * 2; z += 2) {

73 add_mask_counter(temp2_counter, counter[z], temp1_counter, step * 2);

74 memcpy(counter[z], temp1_counter, counter_size);

75 add_mask_counter(temp1_counter, counter[z + 1], temp2_counter, step * 2);

76 memcpy(counter[z + 1], temp2_counter, counter_size);

77 }

78 }

79 memcpy(mask, temp1_counter, counter_size);

80 end:

81 av_freep(&temp1_counter);

82 av_freep(&temp2_counter);

83 for (i = 0; i < 2 * step + 1; i++) {

84 av_freep(&counter[i]);

85 }

86 av_freep(&counter);

87 return ret;

88 }

90 static int copy_separable_masks(cl_mem cl_mask_x, cl_mem cl_mask_y, int step_x, int step_y)

91 {

92 int ret = 0;

93 uint32_t *mask_x, *mask_y;

94 size_t size_mask_x = sizeof(uint32_t) * (2 * step_x + 1);

95 size_t size_mask_y = sizeof(uint32_t) * (2 * step_y + 1);

96 mask_x = av_mallocz_array(2 * step_x + 1, sizeof(uint32_t));

97 if (!mask_x) {

98 ret = AVERROR(ENOMEM);

99 goto end;

100 }

101 mask_y = av_mallocz_array(2 * step_y + 1, sizeof(uint32_t));

102 if (!mask_y) {

103 ret = AVERROR(ENOMEM);

104 goto end;

105 }

106

107 ret = compute_mask(step_x, mask_x);

108 if (ret < 0)

109 goto end;

110 ret = compute_mask(step_y, mask_y);

111 if (ret < 0)

112 goto end;

113

114 ret = av_opencl_buffer_write(cl_mask_x, (uint8_t *)mask_x, size_mask_x);

115 ret = av_opencl_buffer_write(cl_mask_y, (uint8_t *)mask_y, size_mask_y);

116 end:

117 av_freep(&mask_x);

118 av_freep(&mask_y);

119

120 return ret;

121 }

122

123 static int generate_mask(AVFilterContext *ctx)

124 {

125 cl_mem masks[4];

126 cl_mem mask_matrix[2];

127 int i, ret = 0, step_x[2], step_y[2];

128

129 UnsharpContext *unsharp = ctx->priv;

130 mask_matrix[0] = unsharp->opencl_ctx.cl_luma_mask;

131 mask_matrix[1] = unsharp->opencl_ctx.cl_chroma_mask;

132 masks[0] = unsharp->opencl_ctx.cl_luma_mask_x;

133 masks[1] = unsharp->opencl_ctx.cl_luma_mask_y;

134 masks[2] = unsharp->opencl_ctx.cl_chroma_mask_x;

135 masks[3] = unsharp->opencl_ctx.cl_chroma_mask_y;

136 step_x[0] = unsharp->luma.steps_x;

137 step_x[1] = unsharp->chroma.steps_x;

138 step_y[0] = unsharp->luma.steps_y;

139 step_y[1] = unsharp->chroma.steps_y;

140

141 /* use default kernel if any matrix dim larger than 8 due to limited local mem size */

142 if (step_x[0]>8 || step_x[1]>8 || step_y[0]>8 || step_y[1]>8)

143 unsharp->opencl_ctx.use_fast_kernels = 0;

144 else

145 unsharp->opencl_ctx.use_fast_kernels = 1;

146

147 if (!masks[0] || !masks[1] || !masks[2] || !masks[3]) {

148 av_log(ctx, AV_LOG_ERROR, "Luma mask and chroma mask should not be NULL\n");

149 return AVERROR(EINVAL);

150 }

151 if (!mask_matrix[0] || !mask_matrix[1]) {

152 av_log(ctx, AV_LOG_ERROR, "Luma mask and chroma mask should not be NULL\n");

153 return AVERROR(EINVAL);

154 }

155 for (i = 0; i < 2; i++) {

156 ret = copy_separable_masks(masks[2*i], masks[2*i+1], step_x[i], step_y[i]);

157 if (ret < 0)

158 return ret;

159 }

160 return ret;

161 }

162

163 int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out)

164 {

165 int ret;

166 AVFilterLink *link = ctx->inputs[0];

167 UnsharpContext *unsharp = ctx->priv;

168 cl_int status;

169 FFOpenclParam kernel1 = {0};

170 FFOpenclParam kernel2 = {0};

171 int width = link->w;

172 int height = link->h;

173 int cw = FF_CEIL_RSHIFT(link->w, unsharp->hsub);

174 int ch = FF_CEIL_RSHIFT(link->h, unsharp->vsub);

175 size_t globalWorkSize1d = width * height + 2 * ch * cw;

176 size_t globalWorkSize2dLuma[2];

177 size_t globalWorkSize2dChroma[2];

178 size_t localWorkSize2d[2] = {16, 16};

179

180 if (unsharp->opencl_ctx.use_fast_kernels) {

181 globalWorkSize2dLuma[0] = (size_t)ROUND_TO_16(width);

182 globalWorkSize2dLuma[1] = (size_t)ROUND_TO_16(height);

183 globalWorkSize2dChroma[0] = (size_t)ROUND_TO_16(cw);

184 globalWorkSize2dChroma[1] = (size_t)(2*ROUND_TO_16(ch));

185

186 kernel1.ctx = ctx;

187 kernel1.kernel = unsharp->opencl_ctx.kernel_luma;

188 ret = avpriv_opencl_set_parameter(&kernel1,

189 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),

190 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),

191 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask_x),

192 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask_y),

193 FF_OPENCL_PARAM_INFO(unsharp->luma.amount),

194 FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),

195 FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),

196 FF_OPENCL_PARAM_INFO(in->linesize[0]),

197 FF_OPENCL_PARAM_INFO(out->linesize[0]),

198 FF_OPENCL_PARAM_INFO(width),

199 FF_OPENCL_PARAM_INFO(height),

200 NULL);

201 if (ret < 0)

202 return ret;

203

204 kernel2.ctx = ctx;

205 kernel2.kernel = unsharp->opencl_ctx.kernel_chroma;

206 ret = avpriv_opencl_set_parameter(&kernel2,

207 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),

208 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),

209 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask_x),

210 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask_y),

211 FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),

212 FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),

213 FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),

214 FF_OPENCL_PARAM_INFO(in->linesize[0]),

215 FF_OPENCL_PARAM_INFO(in->linesize[1]),

216 FF_OPENCL_PARAM_INFO(out->linesize[0]),

217 FF_OPENCL_PARAM_INFO(out->linesize[1]),

218 FF_OPENCL_PARAM_INFO(link->w),

219 FF_OPENCL_PARAM_INFO(link->h),

220 FF_OPENCL_PARAM_INFO(cw),

221 FF_OPENCL_PARAM_INFO(ch),

222 NULL);

223 if (ret < 0)

224 return ret;

225 status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,

226 unsharp->opencl_ctx.kernel_luma, 2, NULL,

227 globalWorkSize2dLuma, localWorkSize2d, 0, NULL, NULL);

228 status |=clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,

229 unsharp->opencl_ctx.kernel_chroma, 2, NULL,

230 globalWorkSize2dChroma, localWorkSize2d, 0, NULL, NULL);

231 if (status != CL_SUCCESS) {

232 av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));

233 return AVERROR_EXTERNAL;

234 }

235 } else { /* use default kernel */

236 kernel1.ctx = ctx;

237 kernel1.kernel = unsharp->opencl_ctx.kernel_default;

238

239 ret = avpriv_opencl_set_parameter(&kernel1,

240 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),

241 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),

242 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask),

243 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask),

244 FF_OPENCL_PARAM_INFO(unsharp->luma.amount),

245 FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),

246 FF_OPENCL_PARAM_INFO(unsharp->luma.steps_x),

247 FF_OPENCL_PARAM_INFO(unsharp->luma.steps_y),

248 FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_x),

249 FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_y),

250 FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),

251 FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),

252 FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),

253 FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),

254 FF_OPENCL_PARAM_INFO(in->linesize[0]),

255 FF_OPENCL_PARAM_INFO(in->linesize[1]),

256 FF_OPENCL_PARAM_INFO(out->linesize[0]),

257 FF_OPENCL_PARAM_INFO(out->linesize[1]),

258 FF_OPENCL_PARAM_INFO(link->h),

259 FF_OPENCL_PARAM_INFO(link->w),

260 FF_OPENCL_PARAM_INFO(ch),

261 FF_OPENCL_PARAM_INFO(cw),

262 NULL);

263 if (ret < 0)

264 return ret;

265 status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,

266 unsharp->opencl_ctx.kernel_default, 1, NULL,

267 &globalWorkSize1d, NULL, 0, NULL, NULL);

268 if (status != CL_SUCCESS) {

269 av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));

270 return AVERROR_EXTERNAL;

271 }

272 }

273 //blocking map is suffficient, no need for clFinish

274 //clFinish(unsharp->opencl_ctx.command_queue);

275

276 return av_opencl_buffer_read_image(out->data, unsharp->opencl_ctx.out_plane_size,

277 unsharp->opencl_ctx.plane_num, unsharp->opencl_ctx.cl_outbuf,

278 unsharp->opencl_ctx.cl_outbuf_size);

279 }

280

281 int ff_opencl_unsharp_init(AVFilterContext *ctx)

282 {

283 int ret = 0;

284 char build_opts[96];

285 UnsharpContext *unsharp = ctx->priv;

286 ret = av_opencl_init(NULL);

287 if (ret < 0)

288 return ret;

289 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask,

290 sizeof(uint32_t) * (2 * unsharp->luma.steps_x + 1) * (2 * unsharp->luma.steps_y + 1),

291 CL_MEM_READ_ONLY, NULL);

292 if (ret < 0)

293 return ret;

294 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask,

295 sizeof(uint32_t) * (2 * unsharp->chroma.steps_x + 1) * (2 * unsharp->chroma.steps_y + 1),

296 CL_MEM_READ_ONLY, NULL);

297 // separable filters

298 if (ret < 0)

299 return ret;

300 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask_x,

301 sizeof(uint32_t) * (2 * unsharp->luma.steps_x + 1),

302 CL_MEM_READ_ONLY, NULL);

303 if (ret < 0)

304 return ret;

305 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask_y,

306 sizeof(uint32_t) * (2 * unsharp->luma.steps_y + 1),

307 CL_MEM_READ_ONLY, NULL);

308 if (ret < 0)

309 return ret;

310 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask_x,

311 sizeof(uint32_t) * (2 * unsharp->chroma.steps_x + 1),

312 CL_MEM_READ_ONLY, NULL);

313 if (ret < 0)

314 return ret;

315 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask_y,

316 sizeof(uint32_t) * (2 * unsharp->chroma.steps_y + 1),

317 CL_MEM_READ_ONLY, NULL);

318 if (ret < 0)

319 return ret;

320 ret = generate_mask(ctx);

321 if (ret < 0)

322 return ret;

323 unsharp->opencl_ctx.plane_num = PLANE_NUM;

324 unsharp->opencl_ctx.command_queue = av_opencl_get_command_queue();

325 if (!unsharp->opencl_ctx.command_queue) {

326 av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'unsharp'\n");

327 return AVERROR(EINVAL);

328 }

329 snprintf(build_opts, 96, "-D LU_RADIUS_X=%d -D LU_RADIUS_Y=%d -D CH_RADIUS_X=%d -D CH_RADIUS_Y=%d",

330 2*unsharp->luma.steps_x+1, 2*unsharp->luma.steps_y+1, 2*unsharp->chroma.steps_x+1, 2*unsharp->chroma.steps_y+1);

331 unsharp->opencl_ctx.program = av_opencl_compile("unsharp", build_opts);

332 if (!unsharp->opencl_ctx.program) {

333 av_log(ctx, AV_LOG_ERROR, "OpenCL failed to compile program 'unsharp'\n");

334 return AVERROR(EINVAL);

335 }

336 if (unsharp->opencl_ctx.use_fast_kernels) {

337 if (!unsharp->opencl_ctx.kernel_luma) {

338 unsharp->opencl_ctx.kernel_luma = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_luma", &ret);

339 if (ret != CL_SUCCESS) {

340 av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp_luma'\n");

341 return ret;

342 }

343 }

344 if (!unsharp->opencl_ctx.kernel_chroma) {

345 unsharp->opencl_ctx.kernel_chroma = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_chroma", &ret);

346 if (ret < 0) {

347 av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp_chroma'\n");

348 return ret;

349 }

350 }

351 }

352 else {

353 if (!unsharp->opencl_ctx.kernel_default) {

354 unsharp->opencl_ctx.kernel_default = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_default", &ret);

355 if (ret < 0) {

356 av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp_default'\n");

357 return ret;

358 }

359 }

360 }

361 return ret;

362 }

363

364 void ff_opencl_unsharp_uninit(AVFilterContext *ctx)

365 {

366 UnsharpContext *unsharp = ctx->priv;

367 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_inbuf);

368 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_outbuf);

369 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask);

370 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask);

371 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask_x);

372 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask_x);

373 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask_y);

374 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask_y);

375 clReleaseKernel(unsharp->opencl_ctx.kernel_default);

376 clReleaseKernel(unsharp->opencl_ctx.kernel_luma);

377 clReleaseKernel(unsharp->opencl_ctx.kernel_chroma);

378 clReleaseProgram(unsharp->opencl_ctx.program);

379 unsharp->opencl_ctx.command_queue = NULL;

380 av_opencl_uninit();

381 }

382

383 int ff_opencl_unsharp_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)

384 {

385 int ret = 0;

386 AVFilterLink *link = ctx->inputs[0];

387 UnsharpContext *unsharp = ctx->priv;

388 int ch = FF_CEIL_RSHIFT(link->h, unsharp->vsub);

389

390 if ((!unsharp->opencl_ctx.cl_inbuf) || (!unsharp->opencl_ctx.cl_outbuf)) {

391 unsharp->opencl_ctx.in_plane_size[0] = (in->linesize[0] * in->height);

392 unsharp->opencl_ctx.in_plane_size[1] = (in->linesize[1] * ch);

393 unsharp->opencl_ctx.in_plane_size[2] = (in->linesize[2] * ch);

394 unsharp->opencl_ctx.out_plane_size[0] = (out->linesize[0] * out->height);

395 unsharp->opencl_ctx.out_plane_size[1] = (out->linesize[1] * ch);

396 unsharp->opencl_ctx.out_plane_size[2] = (out->linesize[2] * ch);

397 unsharp->opencl_ctx.cl_inbuf_size = unsharp->opencl_ctx.in_plane_size[0] +

398 unsharp->opencl_ctx.in_plane_size[1] +

399 unsharp->opencl_ctx.in_plane_size[2];

400 unsharp->opencl_ctx.cl_outbuf_size = unsharp->opencl_ctx.out_plane_size[0] +

401 unsharp->opencl_ctx.out_plane_size[1] +

402 unsharp->opencl_ctx.out_plane_size[2];

403 if (!unsharp->opencl_ctx.cl_inbuf) {

404 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_inbuf,

405 unsharp->opencl_ctx.cl_inbuf_size,

406 CL_MEM_READ_ONLY, NULL);

407 if (ret < 0)

408 return ret;

409 }

410 if (!unsharp->opencl_ctx.cl_outbuf) {

411 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_outbuf,

412 unsharp->opencl_ctx.cl_outbuf_size,

413 CL_MEM_READ_WRITE, NULL);

414 if (ret < 0)

415 return ret;

416 }

417 }

418 return av_opencl_buffer_write_image(unsharp->opencl_ctx.cl_inbuf,

419 unsharp->opencl_ctx.cl_inbuf_size,

420 0, in->data, unsharp->opencl_ctx.in_plane_size,

421 unsharp->opencl_ctx.plane_num);

422 }

Generated on Sun Mar 8 2015 02:35:06 for FFmpeg by doxygen 1.8.2