Signed-off-by: Ruiling Song <ruiling.song@intel.com>tags/n4.1
| @@ -61,6 +61,18 @@ typedef struct OpenCLFilterContext { | |||
| goto fail; \ | |||
| } | |||
| /** | |||
| * A helper macro to handle OpenCL errors. It will assign errcode to | |||
| * variable err, log error msg, and jump to fail label on error. | |||
| */ | |||
| #define CL_FAIL_ON_ERROR(errcode, ...) do { \ | |||
| if (cle != CL_SUCCESS) { \ | |||
| av_log(avctx, AV_LOG_ERROR, __VA_ARGS__); \ | |||
| err = errcode; \ | |||
| goto fail; \ | |||
| } \ | |||
| } while(0) | |||
| /** | |||
| * Return that all inputs and outputs support only AV_PIX_FMT_OPENCL. | |||
| */ | |||
| @@ -64,26 +64,16 @@ static int avgblur_opencl_init(AVFilterContext *avctx) | |||
| ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context, | |||
| ctx->ocf.hwctx->device_id, | |||
| 0, &cle); | |||
| if (!ctx->command_queue) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL " | |||
| "command queue: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " | |||
| "command queue %d.\n", cle); | |||
| ctx->kernel_horiz = clCreateKernel(ctx->ocf.program,"avgblur_horiz", &cle); | |||
| if (!ctx->kernel_horiz) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create horizontal " | |||
| "kernel %d.\n", cle); | |||
| ctx->kernel_vert = clCreateKernel(ctx->ocf.program,"avgblur_vert", &cle); | |||
| if (!ctx->kernel_vert) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create vertical " | |||
| "kernel %d.\n", cle); | |||
| ctx->initialised = 1; | |||
| return 0; | |||
| @@ -236,12 +226,8 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) | |||
| cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_horiz, 2, NULL, | |||
| global_work, NULL, | |||
| 0, NULL, NULL); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n", | |||
| cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue horizontal " | |||
| "kernel: %d.\n", cle); | |||
| cle = clFinish(ctx->command_queue); | |||
| err = ff_opencl_filter_work_size_from_image(avctx, global_work, | |||
| @@ -259,22 +245,13 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) | |||
| cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_vert, 2, NULL, | |||
| global_work, NULL, | |||
| 0, NULL, NULL); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n", | |||
| cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue vertical " | |||
| "kernel: %d.\n", cle); | |||
| } | |||
| } | |||
| cle = clFinish(ctx->command_queue); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n", | |||
| cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle); | |||
| err = av_frame_copy_props(output, input); | |||
| if (err < 0) | |||
| @@ -100,19 +100,11 @@ static int overlay_opencl_load(AVFilterContext *avctx, | |||
| ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context, | |||
| ctx->ocf.hwctx->device_id, | |||
| 0, &cle); | |||
| if (!ctx->command_queue) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL " | |||
| "command queue: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " | |||
| "command queue %d.\n", cle); | |||
| ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle); | |||
| if (!ctx->kernel) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle); | |||
| ctx->initialised = 1; | |||
| return 0; | |||
| @@ -209,21 +201,12 @@ static int overlay_opencl_blend(FFFrameSync *fs) | |||
| cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, | |||
| global_work, NULL, 0, NULL, NULL); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to enqueue " | |||
| "overlay kernel for plane %d: %d.\n", cle, plane); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue overlay kernel " | |||
| "for plane %d: %d.\n", plane, cle); | |||
| } | |||
| cle = clFinish(ctx->command_queue); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to finish " | |||
| "command queue: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle); | |||
| err = av_frame_copy_props(output, input_main); | |||
| @@ -148,21 +148,11 @@ static int program_opencl_run(AVFilterContext *avctx) | |||
| cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, | |||
| global_work, NULL, 0, NULL, NULL); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n", | |||
| cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle); | |||
| } | |||
| cle = clFinish(ctx->command_queue); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n", | |||
| cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle); | |||
| if (ctx->nb_inputs > 0) { | |||
| err = av_frame_copy_props(output, ctx->frames[0]); | |||
| @@ -262,29 +262,17 @@ static int tonemap_opencl_init(AVFilterContext *avctx) | |||
| ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context, | |||
| ctx->ocf.hwctx->device_id, | |||
| 0, &cle); | |||
| if (!ctx->command_queue) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL " | |||
| "command queue: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " | |||
| "command queue %d.\n", cle); | |||
| ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle); | |||
| if (!ctx->kernel) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle); | |||
| ctx->util_mem = | |||
| clCreateBuffer(ctx->ocf.hwctx->context, 0, | |||
| (2 * DETECTION_FRAMES + 7) * sizeof(unsigned), | |||
| NULL, &cle); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create util buffer: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle); | |||
| ctx->initialised = 1; | |||
| return 0; | |||
| @@ -349,11 +337,7 @@ static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel, | |||
| cle = clEnqueueNDRangeKernel(ctx->command_queue, kernel, 2, NULL, | |||
| global_work, local_work, | |||
| 0, NULL, NULL); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n", | |||
| cle); | |||
| return AVERROR(EIO); | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle); | |||
| return 0; | |||
| fail: | |||
| return err; | |||
| @@ -482,12 +466,7 @@ static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) | |||
| } | |||
| cle = clFinish(ctx->command_queue); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n", | |||
| cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle); | |||
| av_frame_free(&input); | |||
| @@ -76,12 +76,8 @@ static int unsharp_opencl_init(AVFilterContext *avctx) | |||
| ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context, | |||
| ctx->ocf.hwctx->device_id, | |||
| 0, &cle); | |||
| if (!ctx->command_queue) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL " | |||
| "command queue: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " | |||
| "command queue %d.\n", cle); | |||
| // Use global kernel if mask size will be too big for the local store.. | |||
| ctx->global = (ctx->luma_size_x > 17.0f || | |||
| @@ -92,11 +88,7 @@ static int unsharp_opencl_init(AVFilterContext *avctx) | |||
| ctx->kernel = clCreateKernel(ctx->ocf.program, | |||
| ctx->global ? "unsharp_global" | |||
| : "unsharp_local", &cle); | |||
| if (!ctx->kernel) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle); | |||
| ctx->initialised = 1; | |||
| return 0; | |||
| @@ -176,12 +168,8 @@ static int unsharp_opencl_make_filter_params(AVFilterContext *avctx) | |||
| CL_MEM_COPY_HOST_PTR | | |||
| CL_MEM_HOST_NO_ACCESS, | |||
| matrix_bytes, matrix, &cle); | |||
| if (!buffer) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: " | |||
| "%d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create matrix buffer: " | |||
| "%d.\n", cle); | |||
| ctx->plane[p].matrix = buffer; | |||
| } else { | |||
| buffer = clCreateBuffer(ctx->ocf.hwctx->context, | |||
| @@ -190,12 +178,8 @@ static int unsharp_opencl_make_filter_params(AVFilterContext *avctx) | |||
| CL_MEM_HOST_NO_ACCESS, | |||
| sizeof(ctx->plane[p].blur_x), | |||
| ctx->plane[p].blur_x, &cle); | |||
| if (!buffer) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create x-coef buffer: " | |||
| "%d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create x-coef buffer: " | |||
| "%d.\n", cle); | |||
| ctx->plane[p].coef_x = buffer; | |||
| buffer = clCreateBuffer(ctx->ocf.hwctx->context, | |||
| @@ -204,12 +188,8 @@ static int unsharp_opencl_make_filter_params(AVFilterContext *avctx) | |||
| CL_MEM_HOST_NO_ACCESS, | |||
| sizeof(ctx->plane[p].blur_y), | |||
| ctx->plane[p].blur_y, &cle); | |||
| if (!buffer) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to create y-coef buffer: " | |||
| "%d.\n", cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create y-coef buffer: " | |||
| "%d.\n", cle); | |||
| ctx->plane[p].coef_y = buffer; | |||
| } | |||
| @@ -296,21 +276,11 @@ static int unsharp_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) | |||
| cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, | |||
| global_work, ctx->global ? NULL : local_work, | |||
| 0, NULL, NULL); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n", | |||
| cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle); | |||
| } | |||
| cle = clFinish(ctx->command_queue); | |||
| if (cle != CL_SUCCESS) { | |||
| av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n", | |||
| cle); | |||
| err = AVERROR(EIO); | |||
| goto fail; | |||
| } | |||
| CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle); | |||
| err = av_frame_copy_props(output, input); | |||
| if (err < 0) | |||