X-Git-Url: https://git.sesse.net/?a=blobdiff_plain;f=libavfilter%2Fopencl.h;h=7487e602416ce00a62317c45ffc2aad78229d46c;hb=60e7021064ea9e99854c585d452e89f38635480e;hp=973b6d82dd48eb48b7fa2184036788fd91d8b55e;hpb=0fc464631aafa2741e97ce00fc8f81378fcf441a;p=ffmpeg diff --git a/libavfilter/opencl.h b/libavfilter/opencl.h index 973b6d82dd4..7487e602416 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 */