From patchwork Thu Aug 8 13:24:30 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Jarek Samic X-Patchwork-Id: 14316 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 72D8F448D75 for ; Thu, 8 Aug 2019 16:30:57 +0300 (EEST) Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id 4BCAA68AB83; Thu, 8 Aug 2019 16:30:57 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-ot1-f67.google.com (mail-ot1-f67.google.com [209.85.210.67]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 83F0A68A9D8 for ; Thu, 8 Aug 2019 16:30:50 +0300 (EEST) Received: by mail-ot1-f67.google.com with SMTP id l15so118551872otn.9 for ; Thu, 08 Aug 2019 06:30:50 -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:mime-version :content-transfer-encoding; bh=tBUWJHRn2KQYY6l7+cWHOxDpCiyFEfUgjIGcPbgX0sU=; b=XhAXmzjZ2erD1z8ydaDH97UVB/jgQ1iOr3QW2p6R0PXwFi+fKCrYrYfXClr6gD07EU /ky9RiYwlODSFbrqqgV3cxBlVZF/1mIo/S4rPXRh+SHl1QX2aAwrAKfwBagY3Vlz/nZK wW+1uJg2ZZ29BrvxPZEkrJR6reRskt3w5YDzhQ/BzXzeG3f87CotMHsvzws7AHdjnCUp NQXkvXTW2j73L2GaLvqOsCbXBzogD3S+66pu+sn4bHx4GX6J58lQyVR0D48E4fy/47W/ PWLRE5eVibpgV6fszR5WJIK/4lNo29PkWd1mWsFxY8fgBRfzPJuNvLsgrgckmZPIuQkn eZQA== 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:mime-version :content-transfer-encoding; bh=tBUWJHRn2KQYY6l7+cWHOxDpCiyFEfUgjIGcPbgX0sU=; b=RYJaEq/CrnW/1dY5pnsB7S0N/Ky8GqInkgDCw6a3O8YhPXJzheDfG7yDgfjz4z+XVp 5TQ0XfmEBDv/4ASfwAKKPPaucbEMdM7VSbWtl91Xeq+HfNcV1M2roAb3qBoYfwkthCzr RHA3uD29RjOhRzXb8pUcQhrcdNY4Y77b29PVut8bYqQs0Z+w+4jUDelMZeKkIglDXwck Y96r107gPKsnJOPXUvjnGHfkeWPqQApmqsxloLwNevyFO2lekom8qDySpcsGDNwuyZTA uX7Lt1TxvzpDKeYvPFlPq4CxNE2DlQ0CUuBzmyDpspWSATq43u17sMOYJe81vozjNFwC Wqog== X-Gm-Message-State: APjAAAU5Ra//RQ7WYRw+SyeDtTQVZ0EIhICz4Mk53J2CyBfk68DtNEs7 9cZb/QY3x8H7kz7yEZ4GbdATbmiBBcI= X-Google-Smtp-Source: APXvYqwvQcX7Lk9kY/92rxhpml9YiKR7Nl4i5ljNGMFmKqqW1CaJ8grc/nXFMFE2cy8VtyCEu0OtDA== X-Received: by 2002:a6b:bbc1:: with SMTP id l184mr15746384iof.232.1565270684820; Thu, 08 Aug 2019 06:24:44 -0700 (PDT) Received: from cldire-arch.stormhome.local (rrcs-70-61-229-139.central.biz.rr.com. [70.61.229.139]) by smtp.gmail.com with ESMTPSA id n21sm66346197ioh.30.2019.08.08.06.24.43 (version=TLS1_3 cipher=AEAD-AES256-GCM-SHA384 bits=256/256); Thu, 08 Aug 2019 06:24:44 -0700 (PDT) From: Jarek Samic To: ffmpeg-devel@ffmpeg.org Date: Thu, 8 Aug 2019 09:24:30 -0400 Message-Id: <20190808132432.29544-1-cldfire3@gmail.com> X-Mailer: git-send-email 2.22.0 MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH v3 1/3] lavfi: add utilities to reduce OpenCL boilerplate code 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: Jarek Samic Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" --- libavfilter/opencl.c | 10 +++ libavfilter/opencl.h | 142 +++++++++++++++++++++++++++++++++++++++++-- 2 files changed, 146 insertions(+), 6 deletions(-) diff --git a/libavfilter/opencl.c b/libavfilter/opencl.c index 95f0bfc604..8e96543467 100644 --- a/libavfilter/opencl.c +++ b/libavfilter/opencl.c @@ -350,3 +350,13 @@ void ff_opencl_print_const_matrix_3x3(AVBPrint *buf, const char *name_str, } av_bprintf(buf, "};\n"); } + +cl_ulong ff_opencl_get_event_time(cl_event event) { + cl_ulong time_start; + cl_ulong time_end; + + clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL); + clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL); + + return time_end - time_start; +} diff --git a/libavfilter/opencl.h b/libavfilter/opencl.h index 973b6d82dd..7487e60241 100644 --- a/libavfilter/opencl.h +++ b/libavfilter/opencl.h @@ -47,6 +47,11 @@ typedef struct OpenCLFilterContext { int output_height; } OpenCLFilterContext; +// Groups together information about a kernel argument +typedef struct OpenCLKernelArg { + size_t arg_size; + const void *arg_val; +} OpenCLKernelArg; /** * set argument to specific Kernel. @@ -73,9 +78,26 @@ typedef struct OpenCLFilterContext { goto fail; \ } \ } while(0) + +/** + * Create a kernel with the given name. + * + * The kernel variable in the context structure must have a name of the form + * kernel_. + * + * The OpenCLFilterContext variable in the context structure must be named ocf. + * + * Requires the presence of a local cl_int variable named cle and a fail label for error + * handling. + */ +#define CL_CREATE_KERNEL(ctx, kernel_name) do { \ + ctx->kernel_ ## kernel_name = clCreateKernel(ctx->ocf.program, #kernel_name, &cle); \ + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create %s kernel: %d.\n", #kernel_name, cle); \ +} while(0) + /** - * release an OpenCL Kernel - */ + * release an OpenCL Kernel + */ #define CL_RELEASE_KERNEL(k) \ do { \ if (k) { \ @@ -87,8 +109,8 @@ do { \ } while(0) /** - * release an OpenCL Memory Object - */ + * release an OpenCL Memory Object + */ #define CL_RELEASE_MEMORY(m) \ do { \ if (m) { \ @@ -100,8 +122,8 @@ do { \ } while(0) /** - * release an OpenCL Command Queue - */ + * release an OpenCL Command Queue + */ #define CL_RELEASE_QUEUE(q) \ do { \ if (q) { \ @@ -112,6 +134,108 @@ do { \ } \ } while(0) +/** + * Enqueue a kernel with the given information. + * + * Kernel arguments are provided as KernelArg structures and are set in the order + * that they are passed. + * + * Requires the presence of a local cl_int variable named cle and a fail label for error + * handling. + */ +#define CL_ENQUEUE_KERNEL_WITH_ARGS(queue, kernel, global_work_size, local_work_size, event, ...) \ +do { \ + OpenCLKernelArg args[] = {__VA_ARGS__}; \ + for (int i = 0; i < FF_ARRAY_ELEMS(args); i++) { \ + cle = clSetKernelArg(kernel, i, args[i].arg_size, args[i].arg_val); \ + if (cle != CL_SUCCESS) { \ + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " \ + "argument %d: error %d.\n", i, cle); \ + err = AVERROR(EIO); \ + goto fail; \ + } \ + } \ + \ + cle = clEnqueueNDRangeKernel( \ + queue, \ + kernel, \ + FF_ARRAY_ELEMS(global_work_size), \ + NULL, \ + global_work_size, \ + local_work_size, \ + 0, \ + NULL, \ + event \ + ); \ + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle); \ +} while (0) + +/** + * Uses the above macro to enqueue the given kernel and then additionally runs it to + * completion via clFinish. + * + * Requires the presence of a local cl_int variable named cle and a fail label for error + * handling. + */ +#define CL_RUN_KERNEL_WITH_ARGS(queue, kernel, global_work_size, local_work_size, event, ...) do { \ + CL_ENQUEUE_KERNEL_WITH_ARGS( \ + queue, kernel, global_work_size, local_work_size, event, __VA_ARGS__ \ + ); \ + \ + cle = clFinish(queue); \ + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle); \ +} while (0) + +/** + * Create a buffer with the given information. + * + * The buffer variable in the context structure must be named . + * + * Requires the presence of a local cl_int variable named cle and a fail label for error + * handling. + */ +#define CL_CREATE_BUFFER_FLAGS(ctx, buffer_name, flags, size, host_ptr) do { \ + ctx->buffer_name = clCreateBuffer( \ + ctx->ocf.hwctx->context, \ + flags, \ + size, \ + host_ptr, \ + &cle \ + ); \ + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create buffer %s: %d.\n", #buffer_name, cle); \ +} while(0) + +/** + * Perform a blocking write to a buffer. + * + * Requires the presence of a local cl_int variable named cle and a fail label for error + * handling. + */ +#define CL_BLOCKING_WRITE_BUFFER(queue, buffer, size, host_ptr, event) do { \ + cle = clEnqueueWriteBuffer( \ + queue, \ + buffer, \ + CL_TRUE, \ + 0, \ + size, \ + host_ptr, \ + 0, \ + NULL, \ + event \ + ); \ + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to write buffer to device: %d.\n", cle); \ +} while(0) + +/** + * Create a buffer with the given information. + * + * The buffer variable in the context structure must be named . + * + * Requires the presence of a local cl_int variable named cle and a fail label for error + * handling. + */ +#define CL_CREATE_BUFFER(ctx, buffer_name, size) CL_CREATE_BUFFER_FLAGS(ctx, buffer_name, 0, size, NULL) + /** * Return that all inputs and outputs support only AV_PIX_FMT_OPENCL. */ @@ -171,4 +295,10 @@ int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx, void ff_opencl_print_const_matrix_3x3(AVBPrint *buf, const char *name_str, double mat[3][3]); +/** + * Gets the command start and end times for the given event and returns the + * difference (the time that the event took). + */ +cl_ulong ff_opencl_get_event_time(cl_event event); + #endif /* AVFILTER_OPENCL_H */