Skip to content

F
ffmpeg.wasm-core
Commit 89a3be89 authored by Lenny Wang's avatar Lenny Wang Committed by Michael Niedermayer
Browse files
Options

avfilter/opencl: compile kernels separately 

Reviewed-by: 's avatarWei Gao <highgod0401@gmail.com>
Signed-off-by: 's avatarMichael Niedermayer <michaelni@gmx.at>
parent 66825547
Show whitespace changes
Inline Side-by-side
Showing with 51 additions and 24 deletions
libavfilter/deshake.h
View file @ 89a3be89
...@@ -55,6 +55,9 @@ typedef struct { ...@@ -55,6 +55,9 @@ typedef struct {
#if CONFIG_OPENCL #if CONFIG_OPENCL
typedef struct { typedef struct {
cl_command_queue command_queue;
cl_program program;
cl_kernel kernel;
size_t matrix_size; size_t matrix_size;
float matrix_y[9]; float matrix_y[9];
float matrix_uv[9]; float matrix_uv[9];
...@@ -67,7 +70,6 @@ typedef struct { ...@@ -67,7 +70,6 @@ typedef struct {
size_t cl_inbuf_size; size_t cl_inbuf_size;
cl_mem cl_outbuf; cl_mem cl_outbuf;
size_t cl_outbuf_size; size_t cl_outbuf_size;
AVOpenCLKernelEnv kernel_env;
} DeshakeOpenclContext; } DeshakeOpenclContext;
#endif #endif
......
libavfilter/deshake_opencl.c
View file @ 89a3be89
...@@ -45,7 +45,7 @@ int ff_opencl_transform(AVFilterContext *ctx, ...@@ -45,7 +45,7 @@ int ff_opencl_transform(AVFilterContext *ctx,
FFOpenclParam opencl_param = {0}; FFOpenclParam opencl_param = {0};
opencl_param.ctx = ctx; 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)); 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) if (ret < 0)
return ret; return ret;
...@@ -75,14 +75,14 @@ int ff_opencl_transform(AVFilterContext *ctx, ...@@ -75,14 +75,14 @@ int ff_opencl_transform(AVFilterContext *ctx,
NULL); NULL);
if (ret < 0) if (ret < 0)
return ret; return ret;
status = clEnqueueNDRangeKernel(deshake->opencl_ctx.kernel_env.command_queue, status = clEnqueueNDRangeKernel(deshake->opencl_ctx.command_queue,
deshake->opencl_ctx.kernel_env.kernel, 1, NULL, deshake->opencl_ctx.kernel, 1, NULL,
&global_work_size, NULL, 0, NULL, NULL); &global_work_size, NULL, 0, NULL, NULL);
if (status != CL_SUCCESS) { if (status != CL_SUCCESS) {
av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status)); av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL; 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, 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.plane_num, deshake->opencl_ctx.cl_outbuf,
deshake->opencl_ctx.cl_outbuf_size); deshake->opencl_ctx.cl_outbuf_size);
...@@ -108,11 +108,21 @@ int ff_opencl_deshake_init(AVFilterContext *ctx) ...@@ -108,11 +108,21 @@ int ff_opencl_deshake_init(AVFilterContext *ctx)
deshake->opencl_ctx.matrix_size*sizeof(cl_float), CL_MEM_READ_ONLY, NULL); deshake->opencl_ctx.matrix_size*sizeof(cl_float), CL_MEM_READ_ONLY, NULL);
if (ret < 0) if (ret < 0)
return ret; return ret;
if (!deshake->opencl_ctx.kernel_env.kernel) { deshake->opencl_ctx.command_queue = av_opencl_get_command_queue();
ret = av_opencl_create_kernel(&deshake->opencl_ctx.kernel_env, "avfilter_transform"); if (!deshake->opencl_ctx.command_queue) {
if (ret < 0) { av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'deshake'\n");
av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel for name 'avfilter_transform'\n"); return AVERROR(EINVAL);
return ret; }
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; return ret;
...@@ -125,11 +135,12 @@ void ff_opencl_deshake_uninit(AVFilterContext *ctx) ...@@ -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_outbuf);
av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_y); av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_y);
av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_uv); 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(); av_opencl_uninit();
} }
int ff_opencl_deshake_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out) int ff_opencl_deshake_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
{ {
int ret = 0; int ret = 0;
......
libavfilter/unsharp.h
View file @ 89a3be89
...@@ -33,6 +33,9 @@ ...@@ -33,6 +33,9 @@
#if CONFIG_OPENCL #if CONFIG_OPENCL
typedef struct { typedef struct {
cl_command_queue command_queue;
cl_program program;
cl_kernel kernel;
cl_mem cl_luma_mask; cl_mem cl_luma_mask;
cl_mem cl_chroma_mask; cl_mem cl_chroma_mask;
int in_plane_size[8]; int in_plane_size[8];
...@@ -42,7 +45,6 @@ typedef struct { ...@@ -42,7 +45,6 @@ typedef struct {
size_t cl_inbuf_size; size_t cl_inbuf_size;
cl_mem cl_outbuf; cl_mem cl_outbuf;
size_t cl_outbuf_size; size_t cl_outbuf_size;
AVOpenCLKernelEnv kernel_env;
} UnsharpOpenclContext; } UnsharpOpenclContext;
#endif #endif
......
libavfilter/unsharp_opencl.c
View file @ 89a3be89
...@@ -159,7 +159,7 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out) ...@@ -159,7 +159,7 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
FFOpenclParam opencl_param = {0}; FFOpenclParam opencl_param = {0};
opencl_param.ctx = ctx; 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, ret = ff_opencl_set_parameter(&opencl_param,
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf), FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf), FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
...@@ -186,14 +186,14 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out) ...@@ -186,14 +186,14 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
NULL); NULL);
if (ret < 0) if (ret < 0)
return ret; return ret;
status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.kernel_env.command_queue, status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
unsharp->opencl_ctx.kernel_env.kernel, 1, NULL, unsharp->opencl_ctx.kernel, 1, NULL,
&global_work_size, NULL, 0, NULL, NULL); &global_work_size, NULL, 0, NULL, NULL);
if (status != CL_SUCCESS) { if (status != CL_SUCCESS) {
av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status)); av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL; 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, 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.plane_num, unsharp->opencl_ctx.cl_outbuf,
unsharp->opencl_ctx.cl_outbuf_size); unsharp->opencl_ctx.cl_outbuf_size);
...@@ -220,11 +220,21 @@ int ff_opencl_unsharp_init(AVFilterContext *ctx) ...@@ -220,11 +220,21 @@ int ff_opencl_unsharp_init(AVFilterContext *ctx)
if (ret < 0) if (ret < 0)
return ret; return ret;
unsharp->opencl_ctx.plane_num = PLANE_NUM; unsharp->opencl_ctx.plane_num = PLANE_NUM;
if (!unsharp->opencl_ctx.kernel_env.kernel) { unsharp->opencl_ctx.command_queue = av_opencl_get_command_queue();
ret = av_opencl_create_kernel(&unsharp->opencl_ctx.kernel_env, "unsharp"); if (!unsharp->opencl_ctx.command_queue) {
if (ret < 0) { av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'unsharp'\n");
av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel with name 'unsharp'\n"); return AVERROR(EINVAL);
return ret; }
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; return ret;
...@@ -237,7 +247,9 @@ void ff_opencl_unsharp_uninit(AVFilterContext *ctx) ...@@ -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_outbuf);
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask); av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask);
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_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(); av_opencl_uninit();
} }
......
libavfilter/version.h
View file @ 89a3be89
...@@ -31,7 +31,7 @@ ...@@ -31,7 +31,7 @@
#define LIBAVFILTER_VERSION_MAJOR 3 #define LIBAVFILTER_VERSION_MAJOR 3
#define LIBAVFILTER_VERSION_MINOR 90 #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, \ #define LIBAVFILTER_VERSION_INT AV_VERSION_INT(LIBAVFILTER_VERSION_MAJOR, \
LIBAVFILTER_VERSION_MINOR, \ LIBAVFILTER_VERSION_MINOR, \
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment