X-Git-Url: https://git.sesse.net/?a=blobdiff_plain;f=libavfilter%2Fopencl.h;h=7487e602416ce00a62317c45ffc2aad78229d46c;hb=cb9dbc60db1847fcb594684b92334da54ea8757e;hp=1b7f117865304e95934ce93bd1fa91f0a73fdc0b;hpb=2d62e06ff6a9f3dbd78136c1dc4a315a727c6f00;p=ffmpeg diff --git a/libavfilter/opencl.h b/libavfilter/opencl.h index 1b7f1178653..7487e602416 100644 --- a/libavfilter/opencl.h +++ b/libavfilter/opencl.h @@ -25,6 +25,7 @@ // it was introduced in OpenCL 2.0. #define CL_USE_DEPRECATED_OPENCL_1_2_APIS +#include "libavutil/bprint.h" #include "libavutil/buffer.h" #include "libavutil/hwcontext.h" #include "libavutil/hwcontext_opencl.h" @@ -46,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 +79,163 @@ typedef struct OpenCLFilterContext { } \ } 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 + */ +#define CL_RELEASE_KERNEL(k) \ +do { \ + if (k) { \ + cle = clReleaseKernel(k); \ + if (cle != CL_SUCCESS) \ + av_log(avctx, AV_LOG_ERROR, "Failed to release " \ + "OpenCL kernel: %d.\n", cle); \ + } \ +} while(0) + +/** + * release an OpenCL Memory Object + */ +#define CL_RELEASE_MEMORY(m) \ +do { \ + if (m) { \ + cle = clReleaseMemObject(m); \ + if (cle != CL_SUCCESS) \ + av_log(avctx, AV_LOG_ERROR, "Failed to release " \ + "OpenCL memory: %d.\n", cle); \ + } \ +} while(0) + +/** + * release an OpenCL Command Queue + */ +#define CL_RELEASE_QUEUE(q) \ +do { \ + if (q) { \ + cle = clReleaseCommandQueue(q); \ + if (cle != CL_SUCCESS) \ + av_log(avctx, AV_LOG_ERROR, "Failed to release " \ + "OpenCL command queue: %d.\n", cle); \ + } \ +} 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. */ @@ -124,5 +287,18 @@ int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx, size_t *work_size, AVFrame *frame, int plane, int block_alignment); +/** + * Print a 3x3 matrix into a buffer as __constant array, which could + * be included in an OpenCL program. +*/ + +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 */