From patchwork Fri Jun 15 02:55:34 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Danil Iashchenko X-Patchwork-Id: 9416 Delivered-To: ffmpegpatchwork@gmail.com Received: by 2002:a02:11c:0:0:0:0:0 with SMTP id c28-v6csp259137jad; Thu, 14 Jun 2018 20:01:33 -0700 (PDT) X-Google-Smtp-Source: ADUXVKJnhlLXLEW+H30dzaAAw0qBJybImAeTe5ogxIl1MOak2Gc/+ZrbLipcgHPeK3gzAFqYvypv X-Received: by 2002:adf:f045:: with SMTP id t5-v6mr4138309wro.260.1529031692979; Thu, 14 Jun 2018 20:01:32 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1529031692; cv=none; d=google.com; s=arc-20160816; b=C4W01GegvuCySILpjC3oQuKTyaHRKdLIgTWK3dbJIb6ZbPZZgNtzRYXJVwFcr4Afns OOJDC9cGXB92cmLtXHkesB/zORJMU1lTApvZAEk+ekfXA7x00Dlb8Dijhreoelen71O0 lFNcGAkgy5YK/VRzkqW0eA/+W1caUApOUgcUHuDIPu/dGzxLOtawRaqGhF4VaK8fsR5B w1yvM0MeGcbkQ99midCLTC65n7wGcTifW2j5yrCi/xU09xxUknyDIcvxf28sMDJ5ttPu MxltDkzlr0eKDJESWjv+f+8IkFevmfpr+hv9Ezud9EWCKrPBa3u4TBUw7RT+5in3CD2h naPg== 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:references:in-reply-to:message-id:date :to:from:dkim-signature:delivered-to:arc-authentication-results; bh=QMXsiqUP4NbEPpF6Kh8z4y3Fml9MQp0nT7sk9+ZO9Eo=; b=hyd+kK5H32sgeJJ/inejCkckvj2clcY8D8/5IXsxgoZ45ncjZ55QR2mSIEeeYCTLTm zWfOK6d6QN6EnQlvp61NvuHQiNzssCqQTpyZTRWrp+FlHAdvfaRkWuVCktJWPfxhpU0D wfCVf/OQHRzs/Hfy7JjNMyw1LKt6ndXkx/3e6vsobrqOyTc/zGatSZVNy0Gd8W+OZ63x miIN5UCBdgW5iqBoMcNLWcDoR936aUJd8ewHFh0TOd1vXofY5TgOuF0SLfNxGA8ZKdk6 2ribLxywyRbnCbKZdkrCbDerp9I1ppjQKygcOpNxTBAuuK5Dvle3v6uHAP4ixGFVmZmX N+Tg== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20161025 header.b=icVFFSKx; 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=QUARANTINE dis=NONE) header.from=gmail.com Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id y67-v6si593855wmb.111.2018.06.14.20.01.32; Thu, 14 Jun 2018 20:01:32 -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; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20161025 header.b=icVFFSKx; 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=QUARANTINE dis=NONE) header.from=gmail.com Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id C5FDB68A934; Fri, 15 Jun 2018 06:00:39 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-lf0-f65.google.com (mail-lf0-f65.google.com [209.85.215.65]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 5597468A45E for ; Fri, 15 Jun 2018 06:00:33 +0300 (EEST) Received: by mail-lf0-f65.google.com with SMTP id v135-v6so12447867lfa.9 for ; Thu, 14 Jun 2018 20:01:24 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20161025; h=from:to:cc:subject:date:message-id:in-reply-to:references; bh=3zyNstNfx/5pjqJ8nUckQCRl9xNo5yz37foBslXMVg0=; b=icVFFSKx+My2vokUFvFfgqYlN2wrPTOVw2dyiSFR32z15fK1OIoPSPAGKHVQ+T35SY Mxcj1Z/Ll62xlDyDL+kWCAz8NbPoE1yaDNCVH9qr8aSRMGzSLaWHlBjqQVr+2Tfr2hxJ mpGA9Z29Apk2MpduuWLKw/TRXaWcMr7Wp3dBn0c8kv83GT65m0YtpkVStGxGY6fG6PCI 3nVakMqAKfsIKhaaHx8n+D1FnlE5QLPo0gA4i8WLQXXsg3N4GsfnClnHDWuN/yE2Nw0Z ZTyPWGhS3ttLEIPtIYF7OJdeTLDF+5L6aSkSqQpyVEgM1lIsrS5Djp5shKpSweMGmEBm 1Bmw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20161025; h=x-gm-message-state:from:to:cc:subject:date:message-id:in-reply-to :references; bh=3zyNstNfx/5pjqJ8nUckQCRl9xNo5yz37foBslXMVg0=; b=sH/rnRSFafeJNFiNNWdOyv4sITmMZVMxlLMBrPg/pauQ5yC20UDGNvL2pW0KotvhOe vLncJLfRX6VQCdQlS4bXiFSxn+d12Elaxc1bhidZbtKAFPCvHXMgtA3HM60NkkIVctm0 42XatG2LUx+67I7Yk8a68JWQYEBvLCXN6h0DZeFuvKbBYKRhGxwrYqm7RRPjCbyVlZxZ RCOzyoXDEq0mKlJ7B7DOt5KVcFOZgrCKjQgMde+jjEoO9jx9weegQOzF3xxiX0HoIZ2j 6u5VSUGAl+FzZ3sZjU1Ix/JAFbplJ23BsiciipFGXa14yCYngB+1Ahp6w4N336q4bZMO zbJg== X-Gm-Message-State: APt69E2zrTLadJAFyUeFnMbBe4J3H+PzcOmt1DfQlhF8cLoiVqg5/9PX x4CZ8bdzalqa47DrFE1EjMkR6pYuDA== X-Received: by 2002:a2e:8257:: with SMTP id j23-v6mr3252447ljh.1.1529031339465; Thu, 14 Jun 2018 19:55:39 -0700 (PDT) Received: from dan-acer.lan (campus.ifmo.ru. [194.85.161.2]) by smtp.gmail.com with ESMTPSA id c67-v6sm1324351lfe.4.2018.06.14.19.55.38 (version=TLS1_2 cipher=ECDHE-RSA-AES128-SHA bits=128/128); Thu, 14 Jun 2018 19:55:38 -0700 (PDT) From: Danil Iashchenko To: ffmpeg-devel@ffmpeg.org Date: Fri, 15 Jun 2018 05:55:34 +0300 Message-Id: <1529031334-29103-1-git-send-email-danyaschenko@gmail.com> X-Mailer: git-send-email 2.7.4 In-Reply-To: <148B1B7A67D1C24B9EF0BE42EA49770684F9E86F@SHSMSX103.ccr.corp.intel.com> References: <148B1B7A67D1C24B9EF0BE42EA49770684F9E86F@SHSMSX103.ccr.corp.intel.com> Subject: [FFmpeg-devel] [PATCH] libavfilter/opencl.h: Add macro for setting opencl kernel 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: Danil Iashchenko MIME-Version: 1.0 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" --- Hi! I like your idea with OCL_FAIL_ON_ERR(), but still do not know which one is better. My idea relies on fact, that there are only few OpenCL functions which are used multiple times in filters: clSetKernelArg, clCreateKernel(in case when there are multiple kernels) and maybe clEnqueueNDRangeKernel. So that is why my purpose is totally wrap them and significantly reduce code, but yes, there are some restrictions, like you can not use kernel_arg++ when setting kernel arguments. And still most of cl-error checking statements appear after using cl-functions listed above. Thanks, Danil libavfilter/opencl.h | 15 ++++++++++ libavfilter/vf_convolution_opencl.c | 43 ++++------------------------ libavfilter/vf_overlay_opencl.c | 44 +++++++++++----------------- libavfilter/vf_unsharp_opencl.c | 57 ++++++------------------------------- 4 files changed, 46 insertions(+), 113 deletions(-) diff --git a/libavfilter/opencl.h b/libavfilter/opencl.h index c0a4519..7441b11 100644 --- a/libavfilter/opencl.h +++ b/libavfilter/opencl.h @@ -46,6 +46,21 @@ typedef struct OpenCLFilterContext { int output_height; } OpenCLFilterContext; + +/** + * set argument to specific Kernel. + * This macro relies on usage of local label "fail" and variables: + * avctx, cle and err. + */ +#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg) \ + cle = clSetKernelArg(kernel, arg_num, sizeof(type), arg); \ + if (cle != CL_SUCCESS) { \ + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " \ + "argument %d: error %d.\n", arg_num, cle); \ + err = AVERROR(EIO); \ + goto fail; \ + } + /** * Return that all inputs and outputs support only AV_PIX_FMT_OPENCL. */ diff --git a/libavfilter/vf_convolution_opencl.c b/libavfilter/vf_convolution_opencl.c index 2df51e0..4d0ecf8 100644 --- a/libavfilter/vf_convolution_opencl.c +++ b/libavfilter/vf_convolution_opencl.c @@ -204,43 +204,12 @@ static int convolution_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) if (!dst) break; - cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "destination image argument: %d.\n", cle); - goto fail; - } - cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_mem), &src); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "source image argument: %d.\n", cle); - goto fail; - } - cle = clSetKernelArg(ctx->kernel, 2, sizeof(cl_int), &ctx->dims[p]); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "matrix size argument: %d.\n", cle); - goto fail; - } - cle = clSetKernelArg(ctx->kernel, 3, sizeof(cl_mem), &ctx->matrix[p]); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "matrix argument: %d.\n", cle); - goto fail; - } - cle = clSetKernelArg(ctx->kernel, 4, sizeof(cl_float), &ctx->rdivs[p]); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "rdiv argument: %d.\n", cle); - goto fail; - } - cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_float), &ctx->biases[p]); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "bias argument: %d.\n", cle); - goto fail; - } - + CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem, &dst); + CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem, &src); + CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_int, &ctx->dims[p]); + CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_mem, &ctx->matrix[p]); + CL_SET_KERNEL_ARG(ctx->kernel, 4, cl_float, &ctx->rdivs[p]); + CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_float, &ctx->biases[p]); err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0); if (err < 0) diff --git a/libavfilter/vf_overlay_opencl.c b/libavfilter/vf_overlay_opencl.c index b43050d..556ce35 100644 --- a/libavfilter/vf_overlay_opencl.c +++ b/libavfilter/vf_overlay_opencl.c @@ -167,47 +167,39 @@ static int overlay_opencl_blend(FFFrameSync *fs) kernel_arg = 0; mem = (cl_mem)output->data[plane]; - cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem); - if (cle != CL_SUCCESS) - goto fail_kernel_arg; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; mem = (cl_mem)input_main->data[plane]; - cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem); - if (cle != CL_SUCCESS) - goto fail_kernel_arg; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; mem = (cl_mem)input_overlay->data[plane]; - cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem); - if (cle != CL_SUCCESS) - goto fail_kernel_arg; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; if (ctx->alpha_separate) { mem = (cl_mem)input_overlay->data[ctx->nb_planes]; - cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem); - if (cle != CL_SUCCESS) - goto fail_kernel_arg; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; } x = ctx->x_position / (plane == 0 ? 1 : ctx->x_subsample); y = ctx->y_position / (plane == 0 ? 1 : ctx->y_subsample); - cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &x); - if (cle != CL_SUCCESS) - goto fail_kernel_arg; - cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &y); - if (cle != CL_SUCCESS) - goto fail_kernel_arg; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &x); + kernel_arg++; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &y); + kernel_arg++; if (ctx->alpha_separate) { cl_int alpha_adj_x = plane == 0 ? 1 : ctx->x_subsample; cl_int alpha_adj_y = plane == 0 ? 1 : ctx->y_subsample; - cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &alpha_adj_x); - if (cle != CL_SUCCESS) - goto fail_kernel_arg; - cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &alpha_adj_y); - if (cle != CL_SUCCESS) - goto fail_kernel_arg; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &alpha_adj_x); + kernel_arg++; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &alpha_adj_y); + kernel_arg++; } err = ff_opencl_filter_work_size_from_image(avctx, global_work, @@ -241,10 +233,6 @@ static int overlay_opencl_blend(FFFrameSync *fs) return ff_filter_frame(outlink, output); -fail_kernel_arg: - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel arg %d: %d.\n", - kernel_arg, cle); - err = AVERROR(EIO); fail: av_frame_free(&output); return err; diff --git a/libavfilter/vf_unsharp_opencl.c b/libavfilter/vf_unsharp_opencl.c index 19c9185..385d851 100644 --- a/libavfilter/vf_unsharp_opencl.c +++ b/libavfilter/vf_unsharp_opencl.c @@ -268,56 +268,17 @@ static int unsharp_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) if (!dst) break; - cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "destination image argument: %d.\n", cle); - goto fail; - } - cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_mem), &src); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "source image argument: %d.\n", cle); - goto fail; - } - cle = clSetKernelArg(ctx->kernel, 2, sizeof(cl_int), &ctx->plane[p].size_x); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "matrix size argument: %d.\n", cle); - goto fail; - } - cle = clSetKernelArg(ctx->kernel, 3, sizeof(cl_int), &ctx->plane[p].size_y); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "matrix size argument: %d.\n", cle); - goto fail; - } - cle = clSetKernelArg(ctx->kernel, 4, sizeof(cl_float), &ctx->plane[p].amount); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "amount argument: %d.\n", cle); - goto fail; - } + CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem, &dst); + CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem, &src); + CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_int, &ctx->plane[p].size_x); + CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_int, &ctx->plane[p].size_y); + CL_SET_KERNEL_ARG(ctx->kernel, 4, cl_float, &ctx->plane[p].amount); + if (ctx->global) { - cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].matrix); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "matrix argument: %d.\n", cle); - goto fail; - } + CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_mem, &ctx->plane[p].matrix); } else { - cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].coef_x); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "x-coef argument: %d.\n", cle); - goto fail; - } - cle = clSetKernelArg(ctx->kernel, 6, sizeof(cl_mem), &ctx->plane[p].coef_y); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "y-coef argument: %d.\n", cle); - goto fail; - } + CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_mem, &ctx->plane[p].coef_x); + CL_SET_KERNEL_ARG(ctx->kernel, 6, cl_mem, &ctx->plane[p].coef_y); } err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p,