From patchwork Mon Feb 24 10:01:21 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul B Mahol X-Patchwork-Id: 17902 Return-Path: X-Original-To: patchwork@ffaux-bg.ffmpeg.org Delivered-To: patchwork@ffaux-bg.ffmpeg.org Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org [79.124.17.100]) by ffaux.localdomain (Postfix) with ESMTP id 3150C44B56F for ; Mon, 24 Feb 2020 12:09:03 +0200 (EET) Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id 06FA068B0C7; Mon, 24 Feb 2020 12:09:03 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-wm1-f67.google.com (mail-wm1-f67.google.com [209.85.128.67]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 34F7F68B041 for ; Mon, 24 Feb 2020 12:08:56 +0200 (EET) Received: by mail-wm1-f67.google.com with SMTP id q9so8346808wmj.5 for ; Mon, 24 Feb 2020 02:08:56 -0800 (PST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20161025; h=from:to:subject:date:message-id; bh=NN4QsU+/yyhMDX+jT3xqQ/Wp2s5M0YvcEw6beuWjn04=; b=JryCJbYwWUdNXLtmShnooZnwgWOTyaAIsO/J+TEOZ2WoX7ueK6rsS43IhTfd4obSpU /XH6Hk2GQjZtXcI1vbUfL8zoKSEdaSClA2Z8ro8eo59LUP5pCvTesRZ08o/H1HoAPQDb tmhJzaxgSkllQj1iDdQtotSpi0RETg4Atz4t+ceJe6jO/BeT4awJt1+twXiFdpz5YtxS yAD6c68oOqlRcxQCcRX7BgUE58s9TF+CfyNuGYjXxmSLk/+wx4shtWqYjnDwW8Lepfkv N2nZ/dMzL6Tg9PExxzWqGS4setm+rrYWWHXSvJ/rNLMeU8ok1dIzztd43vzNiFNWzy4V 8nTQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20161025; h=x-gm-message-state:from:to:subject:date:message-id; bh=NN4QsU+/yyhMDX+jT3xqQ/Wp2s5M0YvcEw6beuWjn04=; b=Wwhofy4lPzqJPKVH7L8T2nhPZszZOx4FdLfFOAQ4n9CrgNQJ6/dv+KpuxRtxFZ1L88 k5xErfeNovMTnQKTaHeAbpRjrf4OXNuS5rd8tvo2DtMSzU2NV1UG97mCPbMB0u7je/hZ h3FYQCKLPRMSJf4QC9+QwwrVXWn2piZAz9gRgdPqgh4huGkZOXdCddqfqVWxNiNr4Eb9 PVGXfqF8q2fdZ3KVvhIYnn0JYrS5vQdb8f8iTDXtJ+witoVc0sW5JLxLHebTe20zPQP8 jSBUPXUl0oEbo4G6F+gP0oz3QX1IY+7577KuCCa3s1syaKEaM/PQd7CdilP4WOLKT5uB E6Og== X-Gm-Message-State: APjAAAU0idUawq6xKAjNDMBdJm9O1GelVFynEWKi1KO6rce+fT0Wivta uNJrjz8f/MIK8/lvxTcTtL9M4iVj X-Google-Smtp-Source: APXvYqxilnLCNyajjuxFWatpB8Uds6vNURMSFKjuOf06FH94kGliaYuKa9rt6wU99MCueKOHzk687A== X-Received: by 2002:a1c:3b0a:: with SMTP id i10mr22322839wma.177.1582538490182; Mon, 24 Feb 2020 02:01:30 -0800 (PST) Received: from localhost.localdomain ([31.45.248.240]) by smtp.gmail.com with ESMTPSA id a7sm10972842wrm.29.2020.02.24.02.01.29 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 24 Feb 2020 02:01:29 -0800 (PST) From: Paul B Mahol To: ffmpeg-devel@ffmpeg.org Date: Mon, 24 Feb 2020 11:01:21 +0100 Message-Id: <20200224100121.19281-1-onemda@gmail.com> X-Mailer: git-send-email 2.17.1 Subject: [FFmpeg-devel] [PATCH] avfilter/vf_program_opencl: allow setting kernel per plane 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 MIME-Version: 1.0 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Fixes #7190 Signed-off-by: Paul B Mahol --- doc/filters.texi | 22 ++++++++++++ libavfilter/vf_program_opencl.c | 64 ++++++++++++++++++++++----------- 2 files changed, 65 insertions(+), 21 deletions(-) diff --git a/doc/filters.texi b/doc/filters.texi index 70fd7a4cc7..6b10f649b9 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -21302,6 +21302,17 @@ Number of inputs to the filter. Defaults to 1. @item size, s Size of output frames. Defaults to the same as the first input. +@item kernel2 +Kernel name in program for 2nd plane, if not set kernel from option +@var{kernel} is used. + +@item kernel3 +Kernel name in program for 3nd plane, if not set kernel from option +@var{kernel} is used. + +@item kernel4 +Kernel name in program for 4nd plane, if not set kernel from option +@var{kernel} is used. @end table The program source file must contain a kernel function with the given name, @@ -22488,6 +22499,17 @@ Pixel format to use for the generated frames. This must be set. @item rate, r Number of frames generated every second. Default value is '25'. +@item kernel2 +Kernel name in program for 2nd plane, if not set kernel from option +@var{kernel} is used. + +@item kernel3 +Kernel name in program for 3nd plane, if not set kernel from option +@var{kernel} is used. + +@item kernel4 +Kernel name in program for 4nd plane, if not set kernel from option +@var{kernel} is used. @end table For details of how the program loading works, see the @ref{program_opencl} diff --git a/libavfilter/vf_program_opencl.c b/libavfilter/vf_program_opencl.c index ec25e931f5..f748b15037 100644 --- a/libavfilter/vf_program_opencl.c +++ b/libavfilter/vf_program_opencl.c @@ -33,14 +33,14 @@ typedef struct ProgramOpenCLContext { int loaded; cl_uint index; - cl_kernel kernel; + cl_kernel kernel[4]; cl_command_queue command_queue; FFFrameSync fs; AVFrame **frames; const char *source_file; - const char *kernel_name; + const char *kernel_name[4]; int nb_inputs; int width, height; enum AVPixelFormat source_format; @@ -66,15 +66,17 @@ static int program_opencl_load(AVFilterContext *avctx) return AVERROR(EIO); } - ctx->kernel = clCreateKernel(ctx->ocf.program, ctx->kernel_name, &cle); - if (!ctx->kernel) { - if (cle == CL_INVALID_KERNEL_NAME) { - av_log(avctx, AV_LOG_ERROR, "Kernel function '%s' not found in " - "program.\n", ctx->kernel_name); - } else { - av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle); + for (int i = 0; i < 4; i++) { + ctx->kernel[i] = clCreateKernel(ctx->ocf.program, ctx->kernel_name[i] ? ctx->kernel_name[i] : ctx->kernel_name[0], &cle); + if (!ctx->kernel[i]) { + if (cle == CL_INVALID_KERNEL_NAME) { + av_log(avctx, AV_LOG_ERROR, "Kernel function '%s' not found in " + "program.\n", ctx->kernel_name[i] ? ctx->kernel_name[i] : ctx->kernel_name[0]); + } else { + av_log(avctx, AV_LOG_ERROR, "Failed to create kernel%d: %d.\n", i, cle); + } + return AVERROR(EIO); } - return AVERROR(EIO); } ctx->loaded = 1; @@ -108,14 +110,14 @@ static int program_opencl_run(AVFilterContext *avctx) if (!dst) break; - cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst); + cle = clSetKernelArg(ctx->kernel[plane], 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); err = AVERROR_UNKNOWN; goto fail; } - cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_uint), &ctx->index); + cle = clSetKernelArg(ctx->kernel[plane], 1, sizeof(cl_uint), &ctx->index); if (cle != CL_SUCCESS) { av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " "index argument: %d.\n", cle); @@ -129,7 +131,7 @@ static int program_opencl_run(AVFilterContext *avctx) src = (cl_mem)ctx->frames[input]->data[plane]; av_assert0(src); - cle = clSetKernelArg(ctx->kernel, 2 + input, sizeof(cl_mem), &src); + cle = clSetKernelArg(ctx->kernel[plane], 2 + input, sizeof(cl_mem), &src); if (cle != CL_SUCCESS) { av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " "source image argument %d: %d.\n", input, cle); @@ -147,7 +149,7 @@ static int program_opencl_run(AVFilterContext *avctx) "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n", plane, global_work[0], global_work[1]); - cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel[plane], 2, NULL, global_work, NULL, 0, NULL, NULL); CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle); } @@ -312,11 +314,13 @@ static av_cold void program_opencl_uninit(AVFilterContext *avctx) av_freep(&avctx->input_pads[i].name); } - if (ctx->kernel) { - cle = clReleaseKernel(ctx->kernel); - if (cle != CL_SUCCESS) - av_log(avctx, AV_LOG_ERROR, "Failed to release " - "kernel: %d.\n", cle); + for (i = 0; i < 4; i++) { + if (ctx->kernel[i]) { + cle = clReleaseKernel(ctx->kernel[i]); + if (cle != CL_SUCCESS) + av_log(avctx, AV_LOG_ERROR, "Failed to release " + "kernel%d: %d.\n", i, cle); + } } if (ctx->command_queue) { @@ -337,7 +341,7 @@ static av_cold void program_opencl_uninit(AVFilterContext *avctx) static const AVOption program_opencl_options[] = { { "source", "OpenCL program source file", OFFSET(source_file), AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS }, - { "kernel", "Kernel name in program", OFFSET(kernel_name), + { "kernel", "Kernel name in program", OFFSET(kernel_name[0]), AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS }, { "inputs", "Number of inputs", OFFSET(nb_inputs), @@ -348,6 +352,15 @@ static const AVOption program_opencl_options[] = { { "s", "Video size", OFFSET(width), AV_OPT_TYPE_IMAGE_SIZE, { .str = NULL }, 0, 0, FLAGS }, + { "kernel2", "Kernel name in program for 2nd plane", OFFSET(kernel_name[1]), + AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS }, + + { "kernel3", "Kernel name in program for 3rd plane", OFFSET(kernel_name[2]), + AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS }, + + { "kernel4", "Kernel name in program for 4th plane", OFFSET(kernel_name[3]), + AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS }, + { NULL }, }; @@ -384,7 +397,7 @@ AVFilter ff_vf_program_opencl = { static const AVOption openclsrc_options[] = { { "source", "OpenCL program source file", OFFSET(source_file), AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS }, - { "kernel", "Kernel name in program", OFFSET(kernel_name), + { "kernel", "Kernel name in program", OFFSET(kernel_name[0]), AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS }, { "size", "Video size", OFFSET(width), @@ -400,6 +413,15 @@ static const AVOption openclsrc_options[] = { { "r", "Video frame rate", OFFSET(source_rate), AV_OPT_TYPE_VIDEO_RATE, { .str = "25" }, 0, INT_MAX, FLAGS }, + { "kernel2", "Kernel name in program for 2nd plane", OFFSET(kernel_name[1]), + AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS }, + + { "kernel3", "Kernel name in program for 3rd plane", OFFSET(kernel_name[2]), + AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS }, + + { "kernel4", "Kernel name in program for 4th plane", OFFSET(kernel_name[3]), + AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS }, + { NULL }, };