From patchwork Tue Jul 3 18:16:24 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Ruiling Song X-Patchwork-Id: 9596 Delivered-To: ffmpegpatchwork@gmail.com Received: by 2002:a02:104:0:0:0:0:0 with SMTP id c4-v6csp1417344jad; Tue, 3 Jul 2018 11:17:45 -0700 (PDT) X-Google-Smtp-Source: AAOMgpeUa8yvykynBRYL91cPpEmfLkXcVr//DzdWzhrSM+fQ7l9K4zqrI6O8SEEQoUXYdFzP7QP2 X-Received: by 2002:a1c:e541:: with SMTP id c62-v6mr9023534wmh.154.1530641865481; Tue, 03 Jul 2018 11:17:45 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1530641865; cv=none; d=google.com; s=arc-20160816; b=eJGifOVAmTg/rpT8bJgtKUPL0WP6AGCknTuJNtxMHLzU4IOfsAmrAN8xEqIhhbwK7i mN4RM77Tri4F8Sp7CCY2qvIiz5r++m7Ni6ieRjc/UuFDiGrvIOelivf6rOUDZX6dP1Pp /xmB+WJmT29Ket07L5WZctApe76+JpIxwLLyQP84x33VDYZHx6YRA5/HWLBzT9tJlUgw UyquXScVouN8BPm0wOTePi8D4DzH4kndGjXGn4PjocKFKRIaVLJ1v6+N/+fNSD6QAZ6V FpqN65//me2tdWA/R8rxX/xvmgoeYeQ4ouPiWyZ+cHv+icqHO60S0Dq/r0IAJJTxS/gw Bchw== 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=3JGRACr9vtWsCV02Pe80F0t3CbfNPzvwDJWZXvycxLI=; b=eTmGTAEcDxUlTh5rDR3jy1nX0SMa4qPqnTcANNSMQMPBnk0lRpL4VSu14RAy7J6WMt Uv5Tr+ga3sYkBze/RhLMnj68qUl2pDhK882b8gbTaU6IPXREnAvzbbA5GW2Lo5PMZJLD nR2kXwTX7w8QwJRac2FExJeiZDDNVGAuF71GosfF5QCTUcWZ0T9r4DVMYumQCETUYZYM i87M/3BPF4ZUkmODraFa14T/nNDl1LoW7G8zRSq27kw+YBNlBMzZLcuZeQDbCq+cdhK0 2YlJZPOKip5ruVgJyAJT1UZ2A9jlKNhEuhBRI4LIMDRX20Mmfd0/bT1PkgM7BqWJSG2e MdUQ== 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 y15-v6si1411726wrn.343.2018.07.03.11.17.44; Tue, 03 Jul 2018 11:17:45 -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 CD44068A8EF; Tue, 3 Jul 2018 21:17:38 +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 683BB68A8B2 for ; Tue, 3 Jul 2018 21:17:31 +0300 (EEST) X-Amp-Result: SKIPPED(no attachment in message) X-Amp-File-Uploaded: False Received: from fmsmga001.fm.intel.com ([10.253.24.23]) by orsmga101.jf.intel.com with ESMTP/TLS/DHE-RSA-AES256-GCM-SHA384; 03 Jul 2018 11:17:31 -0700 X-ExtLoop1: 1 X-IronPort-AV: E=Sophos;i="5.51,304,1526367600"; d="scan'208";a="69365531" Received: from ruiling-skl2.sh.intel.com ([10.239.160.154]) by fmsmga001.fm.intel.com with ESMTP; 03 Jul 2018 11:17:30 -0700 From: Ruiling Song To: ffmpeg-devel@ffmpeg.org Date: Wed, 4 Jul 2018 02:16:24 +0800 Message-Id: <1530641785-22104-1-git-send-email-ruiling.song@intel.com> X-Mailer: git-send-email 2.7.4 Subject: [FFmpeg-devel] [PATCH v2 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 | 11 +++++++++ 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, 47 insertions(+), 137 deletions(-) diff --git a/libavfilter/opencl.h b/libavfilter/opencl.h index 7441b11..0ed360b 100644 --- a/libavfilter/opencl.h +++ b/libavfilter/opencl.h @@ -112,5 +112,16 @@ int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx, size_t *work_size, AVFrame *frame, int plane, int block_alignment); +/** + * 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 CL_FAIL_ON_ERROR(errcode, ...) do {\ + if (cle != CL_SUCCESS) {\ + av_log(avctx, AV_LOG_ERROR, __VA_ARGS__);\ + err = errcode;\ + goto fail;\ + }\ +} while(0) #endif /* AVFILTER_OPENCL_H */ 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)