From 89a3be8921e265a487a5dfbd68671393b532edb7 Mon Sep 17 00:00:00 2001 From: Lenny Wang Date: Sun, 3 Nov 2013 21:58:09 -0600 Subject: avfilter/opencl: compile kernels separately Reviewed-by: Wei Gao Signed-off-by: Michael Niedermayer --- libavfilter/deshake.h | 4 +++- libavfilter/deshake_opencl.c | 33 ++++++++++++++++++++++----------- libavfilter/unsharp.h | 4 +++- libavfilter/unsharp_opencl.c | 32 ++++++++++++++++++++++---------- libavfilter/version.h | 2 +- 5 files changed, 51 insertions(+), 24 deletions(-) (limited to 'libavfilter') diff --git a/libavfilter/deshake.h b/libavfilter/deshake.h index c24090ee36..5792973957 100644 --- a/libavfilter/deshake.h +++ b/libavfilter/deshake.h @@ -55,6 +55,9 @@ typedef struct { #if CONFIG_OPENCL typedef struct { + cl_command_queue command_queue; + cl_program program; + cl_kernel kernel; size_t matrix_size; float matrix_y[9]; float matrix_uv[9]; @@ -67,7 +70,6 @@ typedef struct { size_t cl_inbuf_size; cl_mem cl_outbuf; size_t cl_outbuf_size; - AVOpenCLKernelEnv kernel_env; } DeshakeOpenclContext; #endif diff --git a/libavfilter/deshake_opencl.c b/libavfilter/deshake_opencl.c index eea873ea4a..e4e4df19e8 100644 --- a/libavfilter/deshake_opencl.c +++ b/libavfilter/deshake_opencl.c @@ -45,7 +45,7 @@ int ff_opencl_transform(AVFilterContext *ctx, FFOpenclParam opencl_param = {0}; opencl_param.ctx = ctx; - opencl_param.kernel = deshake->opencl_ctx.kernel_env.kernel; + opencl_param.kernel = deshake->opencl_ctx.kernel; ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_y, (uint8_t *)matrix_y, deshake->opencl_ctx.matrix_size * sizeof(cl_float)); if (ret < 0) return ret; @@ -75,14 +75,14 @@ int ff_opencl_transform(AVFilterContext *ctx, NULL); if (ret < 0) return ret; - status = clEnqueueNDRangeKernel(deshake->opencl_ctx.kernel_env.command_queue, - deshake->opencl_ctx.kernel_env.kernel, 1, NULL, + status = clEnqueueNDRangeKernel(deshake->opencl_ctx.command_queue, + deshake->opencl_ctx.kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL); if (status != CL_SUCCESS) { av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status)); return AVERROR_EXTERNAL; } - clFinish(deshake->opencl_ctx.kernel_env.command_queue); + clFinish(deshake->opencl_ctx.command_queue); ret = av_opencl_buffer_read_image(out->data, deshake->opencl_ctx.out_plane_size, deshake->opencl_ctx.plane_num, deshake->opencl_ctx.cl_outbuf, deshake->opencl_ctx.cl_outbuf_size); @@ -108,11 +108,21 @@ int ff_opencl_deshake_init(AVFilterContext *ctx) deshake->opencl_ctx.matrix_size*sizeof(cl_float), CL_MEM_READ_ONLY, NULL); if (ret < 0) return ret; - if (!deshake->opencl_ctx.kernel_env.kernel) { - ret = av_opencl_create_kernel(&deshake->opencl_ctx.kernel_env, "avfilter_transform"); - if (ret < 0) { - av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel for name 'avfilter_transform'\n"); - return ret; + deshake->opencl_ctx.command_queue = av_opencl_get_command_queue(); + if (!deshake->opencl_ctx.command_queue) { + av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'deshake'\n"); + return AVERROR(EINVAL); + } + deshake->opencl_ctx.program = av_opencl_compile("avfilter_transform", NULL); + if (!deshake->opencl_ctx.program) { + av_log(ctx, AV_LOG_ERROR, "OpenCL failed to compile program 'avfilter_transform'\n"); + return AVERROR(EINVAL); + } + if (!deshake->opencl_ctx.kernel) { + deshake->opencl_ctx.kernel = clCreateKernel(deshake->opencl_ctx.program, "avfilter_transform", &ret); + if (ret != CL_SUCCESS) { + av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'avfilter_transform'\n"); + return AVERROR(EINVAL); } } return ret; @@ -125,11 +135,12 @@ void ff_opencl_deshake_uninit(AVFilterContext *ctx) av_opencl_buffer_release(&deshake->opencl_ctx.cl_outbuf); av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_y); av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_uv); - av_opencl_release_kernel(&deshake->opencl_ctx.kernel_env); + clReleaseKernel(deshake->opencl_ctx.kernel); + clReleaseProgram(deshake->opencl_ctx.program); + deshake->opencl_ctx.command_queue = NULL; av_opencl_uninit(); } - int ff_opencl_deshake_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out) { int ret = 0; diff --git a/libavfilter/unsharp.h b/libavfilter/unsharp.h index c2259298f6..2738243cea 100644 --- a/libavfilter/unsharp.h +++ b/libavfilter/unsharp.h @@ -33,6 +33,9 @@ #if CONFIG_OPENCL typedef struct { + cl_command_queue command_queue; + cl_program program; + cl_kernel kernel; cl_mem cl_luma_mask; cl_mem cl_chroma_mask; int in_plane_size[8]; @@ -42,7 +45,6 @@ typedef struct { size_t cl_inbuf_size; cl_mem cl_outbuf; size_t cl_outbuf_size; - AVOpenCLKernelEnv kernel_env; } UnsharpOpenclContext; #endif diff --git a/libavfilter/unsharp_opencl.c b/libavfilter/unsharp_opencl.c index b373b66116..77b59371a8 100644 --- a/libavfilter/unsharp_opencl.c +++ b/libavfilter/unsharp_opencl.c @@ -159,7 +159,7 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out) FFOpenclParam opencl_param = {0}; opencl_param.ctx = ctx; - opencl_param.kernel = unsharp->opencl_ctx.kernel_env.kernel; + opencl_param.kernel = unsharp->opencl_ctx.kernel; ret = ff_opencl_set_parameter(&opencl_param, FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf), FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf), @@ -186,14 +186,14 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out) NULL); if (ret < 0) return ret; - status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.kernel_env.command_queue, - unsharp->opencl_ctx.kernel_env.kernel, 1, NULL, + status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue, + unsharp->opencl_ctx.kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL); if (status != CL_SUCCESS) { av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status)); return AVERROR_EXTERNAL; } - clFinish(unsharp->opencl_ctx.kernel_env.command_queue); + clFinish(unsharp->opencl_ctx.command_queue); return av_opencl_buffer_read_image(out->data, unsharp->opencl_ctx.out_plane_size, unsharp->opencl_ctx.plane_num, unsharp->opencl_ctx.cl_outbuf, unsharp->opencl_ctx.cl_outbuf_size); @@ -220,11 +220,21 @@ int ff_opencl_unsharp_init(AVFilterContext *ctx) if (ret < 0) return ret; unsharp->opencl_ctx.plane_num = PLANE_NUM; - if (!unsharp->opencl_ctx.kernel_env.kernel) { - ret = av_opencl_create_kernel(&unsharp->opencl_ctx.kernel_env, "unsharp"); - if (ret < 0) { - av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel with name 'unsharp'\n"); - return ret; + unsharp->opencl_ctx.command_queue = av_opencl_get_command_queue(); + if (!unsharp->opencl_ctx.command_queue) { + av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'unsharp'\n"); + return AVERROR(EINVAL); + } + unsharp->opencl_ctx.program = av_opencl_compile("unsharp", NULL); + if (!unsharp->opencl_ctx.program) { + av_log(ctx, AV_LOG_ERROR, "OpenCL failed to compile program 'unsharp'\n"); + return AVERROR(EINVAL); + } + if (!unsharp->opencl_ctx.kernel) { + unsharp->opencl_ctx.kernel = clCreateKernel(unsharp->opencl_ctx.program, "unsharp", &ret); + if (ret != CL_SUCCESS) { + av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp'\n"); + return AVERROR(EINVAL); } } return ret; @@ -237,7 +247,9 @@ void ff_opencl_unsharp_uninit(AVFilterContext *ctx) av_opencl_buffer_release(&unsharp->opencl_ctx.cl_outbuf); av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask); av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask); - av_opencl_release_kernel(&unsharp->opencl_ctx.kernel_env); + clReleaseKernel(unsharp->opencl_ctx.kernel); + clReleaseProgram(unsharp->opencl_ctx.program); + unsharp->opencl_ctx.command_queue = NULL; av_opencl_uninit(); } diff --git a/libavfilter/version.h b/libavfilter/version.h index a293d3bfd2..546ab5b338 100644 --- a/libavfilter/version.h +++ b/libavfilter/version.h @@ -31,7 +31,7 @@ #define LIBAVFILTER_VERSION_MAJOR 3 #define LIBAVFILTER_VERSION_MINOR 90 -#define LIBAVFILTER_VERSION_MICRO 101 +#define LIBAVFILTER_VERSION_MICRO 102 #define LIBAVFILTER_VERSION_INT AV_VERSION_INT(LIBAVFILTER_VERSION_MAJOR, \ LIBAVFILTER_VERSION_MINOR, \ -- cgit v1.2.3