From patchwork Fri Jul 26 22:02:24 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Jarek Samic X-Patchwork-Id: 14091 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 217FE448383 for ; Sat, 27 Jul 2019 01:09:19 +0300 (EEST) Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id ECE6168AC57; Sat, 27 Jul 2019 01:09:18 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-io1-f65.google.com (mail-io1-f65.google.com [209.85.166.65]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 1630C68AB9D for ; Sat, 27 Jul 2019 01:09:12 +0300 (EEST) Received: by mail-io1-f65.google.com with SMTP id f4so107897785ioh.6 for ; Fri, 26 Jul 2019 15:09:12 -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=Ba7tIVvnQwZ69EvFPDcXlgyUtE+wC2ElWYCtCF/3oeU1YuFWwMTpSd/i/KIuN+LhbS GJVHf+AkGCRvPaIRxrMwUTwL4mcbzlZXknEKX1m8ajAhENndvldRxiQYgnlVDHs7/L4B 2Upa4F5+V9Hc4n5Wf9YRvvC+NbOwnsPvCSt36tUwZ3DNVgrq4+qSl4iBwhRYAKCfqYmG j/Byh6avluMviIGKu9GxnocedU9ceM3RBJitpnl87T5YIR6VXQO2GE9r50k11XLGd2BC xS9ZzJovC0HdTZmkZiOGr6YeI1SDcTjUDi/XFasVson7fgQqbiMDy/O5IczOLjQ6hwiF j3dQ== 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=hQA3XscMO9/IZExNJi1Vfa5IgfjLaUt90pjmL3qoOIIjZqDNDlXlQa0UxNPsXxcFSX bVwrfnjToan8eBEx9AUt+ZMvsjNmOGglJgO7Cul9ZoVfLyVsWvYIwI2KvyxgF+NkQSDF mKfMRTULi+YX3m4wxMRVGUVO5oAMWHH43kW6hYUTZmpRmPnioQyXSedChDe2q+5NX6ko nshOrshhPdzFcqcn0xXuW3mYBHgmcURsfYI1ieTiRhPZziSQyDI0ycKjONGNzdftl8z4 3jRGFElvC14AbF5Bd69ycKoA3gYFBapz3CONuAumboB/xQH+qVcMwiZKmn4QqJslxFVg RQLQ== X-Gm-Message-State: APjAAAU7uTWn1gaItgMgzu4+062UYDqtrlTbKhiS9DRRbKM7qLlC2IIa nfZPGRSKffKAOBVJEWI8CERdoYSKogo= X-Google-Smtp-Source: APXvYqzAA+EDyCfBqTyJvSiyWmaRdipXs1N7AY8i615xzUZFOgqtNWwgMSBGsRNbd3W654nkr1IvrA== X-Received: by 2002:a02:c492:: with SMTP id t18mr99931606jam.67.1564178592438; Fri, 26 Jul 2019 15:03:12 -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 t133sm82535226iof.21.2019.07.26.15.03.10 (version=TLS1_3 cipher=AEAD-AES256-GCM-SHA384 bits=256/256); Fri, 26 Jul 2019 15:03:11 -0700 (PDT) From: Jarek Samic To: ffmpeg-devel@ffmpeg.org Date: Fri, 26 Jul 2019 18:02:24 -0400 Message-Id: <20190726220226.28369-1-cldfire3@gmail.com> X-Mailer: git-send-email 2.22.0 MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH v2 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 */