lavfi: add utilities to reduce OpenCL boilerplate code

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;
+}

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,6 +78,23 @@ 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_<kernel_name>.
+ *
+ * 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
  */
@@ -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 <buffer_name>.
+ *
+ * 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 <buffer_name>.
+ *
+ * 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 */