From patchwork Mon Jul 2 20:55:52 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Ruiling Song X-Patchwork-Id: 9575 Delivered-To: ffmpegpatchwork@gmail.com Received: by 2002:a02:104:0:0:0:0:0 with SMTP id c4-v6csp389310jad; Mon, 2 Jul 2018 13:57:07 -0700 (PDT) X-Google-Smtp-Source: AAOMgpd1fLvXZCAMu8WOjxUtK5nRvcxgvgTLSAc1N7ldGNeYXrZb7citBudjY4i84uu/k3f0hh4Z X-Received: by 2002:adf:ff49:: with SMTP id u9-v6mr3279812wrs.15.1530565027264; Mon, 02 Jul 2018 13:57:07 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1530565027; cv=none; d=google.com; s=arc-20160816; b=uqSjwNX7dHS/b26EVB5PWstTQBqFAnuER6KsrHR66twHLkSSa2TkrJ2jhIt56GWSaK s0pFOwkoMau3QttR+wp1XAqgYE6NPDAh/Wm5E2qVwGDanDYoCSnZtt7p7LxuofHhMweE Y5Q7jIHTsZoAFe6dHGUvAibgrJalGwafFJmFxibMN9GCyBjk3mJpn0BQRGLlTmBU6vW0 us+O7a84p7Li3UbrOO1gwe7FPlRKLyUS+59TwYAVf+vvV/zCvVhsVsI5+F6bwNQj9ZC3 UkkAKIVzZj7AV/ppvW7oxuqATLvYTgIzOtKBDhCoPguKIO7ZbgOYROTLfSb+LUXhQmte y7PA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:content-transfer-encoding:mime-version:cc:reply-to :list-subscribe:list-help:list-post:list-archive:list-unsubscribe :list-id:precedence:subject:message-id:date:to:from:delivered-to :arc-authentication-results; bh=mGibD+HcXZxWaEmvP6gVO55R/3Pcflg3SF55eLTinJo=; b=qLxd6VkBpXwefeonEf34IQaJZGkuCd/XFj10MOvjgIxbNsbLKU/sjkb98AXaUGfB2K yNJRcFezh8XFHSYgnyI8AOZBRXKIUeUiUl9xk9AvX34XYEx/2Rtfx6Vbr5bV8Z1rSjpW 5/v6DXLQxrJaU1G7aYloTKgSUoXRZR0Mp8+XkVcay4fDUxQ9U8kC89ShkQFNekvN2U/u m0EOHwpgoSN8YJisNCc6TCVu3sQ4N3cab3NUNTDt9M6wKqL+imQ/ebRC6Hlym8YC/GhC aRTnFOKuJ3fj5FL7Z/NPHbbla5ZElmgTI3ZC5BF0vZSFcPYhPdJGARWaTezMVRwv39zA 7Tgw== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org; dmarc=fail (p=NONE sp=NONE dis=NONE) header.from=intel.com Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id o12-v6si4842926wmh.0.2018.07.02.13.57.06; Mon, 02 Jul 2018 13:57:07 -0700 (PDT) Received-SPF: pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) client-ip=79.124.17.100; Authentication-Results: mx.google.com; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org; dmarc=fail (p=NONE sp=NONE dis=NONE) header.from=intel.com Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id 4B61268A6D9; Mon, 2 Jul 2018 23:57:01 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mga02.intel.com (mga02.intel.com [134.134.136.20]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id B534B68A6D4 for ; Mon, 2 Jul 2018 23:56:54 +0300 (EEST) X-Amp-Result: SKIPPED(no attachment in message) X-Amp-File-Uploaded: False Received: from fmsmga008.fm.intel.com ([10.253.24.58]) by orsmga101.jf.intel.com with ESMTP/TLS/DHE-RSA-AES256-GCM-SHA384; 02 Jul 2018 13:56:55 -0700 X-ExtLoop1: 1 X-IronPort-AV: E=Sophos;i="5.51,300,1526367600"; d="scan'208";a="52140838" Received: from ruiling-skl2.sh.intel.com ([10.239.160.154]) by fmsmga008.fm.intel.com with ESMTP; 02 Jul 2018 13:56:54 -0700 From: Ruiling Song To: ffmpeg-devel@ffmpeg.org Date: Tue, 3 Jul 2018 04:55:52 +0800 Message-Id: <1530564953-4747-1-git-send-email-ruiling.song@intel.com> X-Mailer: git-send-email 2.7.4 Subject: [FFmpeg-devel] [PATCH 1/2] lavfi/opencl: add macro for opencl error handling. X-BeenThere: ffmpeg-devel@ffmpeg.org X-Mailman-Version: 2.1.20 Precedence: list List-Id: FFmpeg development discussions and patches List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Reply-To: FFmpeg development discussions and patches Cc: Ruiling Song MIME-Version: 1.0 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Signed-off-by: Ruiling Song --- libavfilter/opencl.h | 4 ++-- libavfilter/vf_avgblur_opencl.c | 45 +++++++++-------------------------- libavfilter/vf_overlay_opencl.c | 29 +++++------------------ libavfilter/vf_program_opencl.c | 14 ++--------- libavfilter/vf_tonemap_opencl.c | 33 +++++--------------------- libavfilter/vf_unsharp_opencl.c | 52 +++++++++-------------------------------- 6 files changed, 38 insertions(+), 139 deletions(-) diff --git a/libavfilter/opencl.h b/libavfilter/opencl.h index fd76f72..0ed360b 100644 --- a/libavfilter/opencl.h +++ b/libavfilter/opencl.h @@ -116,9 +116,9 @@ int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx, * A helper macro to handle OpenCL error. It will assign errcode to * variable err, log error msg, and jump to fail label on error. */ -#define OCL_FAIL_ON_ERR(logctx, cle, errcode, ...) do {\ +#define CL_FAIL_ON_ERROR(errcode, ...) do {\ if (cle != CL_SUCCESS) {\ - av_log(logctx, AV_LOG_ERROR, __VA_ARGS__);\ + av_log(avctx, AV_LOG_ERROR, __VA_ARGS__);\ err = errcode;\ goto fail;\ }\ diff --git a/libavfilter/vf_avgblur_opencl.c b/libavfilter/vf_avgblur_opencl.c index d1d3eb1..bc6bcab 100644 --- a/libavfilter/vf_avgblur_opencl.c +++ b/libavfilter/vf_avgblur_opencl.c @@ -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) diff --git a/libavfilter/vf_overlay_opencl.c b/libavfilter/vf_overlay_opencl.c index 556ce35..e9c8532 100644 --- a/libavfilter/vf_overlay_opencl.c +++ b/libavfilter/vf_overlay_opencl.c @@ -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); diff --git a/libavfilter/vf_program_opencl.c b/libavfilter/vf_program_opencl.c index a002792..dfb2565 100644 --- a/libavfilter/vf_program_opencl.c +++ b/libavfilter/vf_program_opencl.c @@ -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]); diff --git a/libavfilter/vf_tonemap_opencl.c b/libavfilter/vf_tonemap_opencl.c index 36c7fbe..241f95e 100644 --- a/libavfilter/vf_tonemap_opencl.c +++ b/libavfilter/vf_tonemap_opencl.c @@ -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); diff --git a/libavfilter/vf_unsharp_opencl.c b/libavfilter/vf_unsharp_opencl.c index 5b1eb59..d76d1b1 100644 --- a/libavfilter/vf_unsharp_opencl.c +++ b/libavfilter/vf_unsharp_opencl.c @@ -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)