[FFmpeg-devel] [PATCH]opencl: compile kernels separately
Wei Gao
highgod0401 at gmail.com
Tue Oct 29 09:19:01 CET 2013
2013/10/29 Lenny Wang <lenny at multicorewareinc.com>
> Currently opencl kernels in ffmpeg are compiled altogether at
> initialization stage, most of related data structures are maintained within
> the opencl framework of ffmpeg. This is very cumbersome to use and is not
> efficient. This patch uses distributed opencl programs/kernels for each
> filter (or potentially any other component that uses opencl), allowing
> kernels to be compiled separately at each component's initialization stage.
>
> Tests have been conducted successfully on mainstream Nvidia/AMD/Intel
> platforms with "-vf deshake=opencl=1,unsharp=opencl=1".
>
libavfilter/deshake.h | 9 +-
libavfilter/deshake_opencl.c | 17 ++--
libavfilter/opencl_allkernels.c | 8 +-
libavfilter/unsharp.h | 5 +-
libavfilter/unsharp_opencl.c | 20 ++--
libavutil/opencl.c | 213
++++++++++++++--------------------------
libavutil/opencl.h | 43 ++++----
7 files changed, 126 insertions(+), 189 deletions(-)
diff --git a/libavfilter/deshake.h b/libavfilter/deshake.h
index c24090e..e9ea00c 100644
--- a/libavfilter/deshake.h
+++ b/libavfilter/deshake.h
@@ -58,15 +58,16 @@ typedef struct {
size_t matrix_size;
float matrix_y[9];
float matrix_uv[9];
- cl_mem cl_matrix_y;
- cl_mem cl_matrix_uv;
int in_plane_size[8];
int out_plane_size[8];
int plane_num;
- cl_mem cl_inbuf;
size_t cl_inbuf_size;
- cl_mem cl_outbuf;
size_t cl_outbuf_size;
+ cl_mem cl_matrix_y;
+ cl_mem cl_matrix_uv;
+ cl_mem cl_inbuf;
+ cl_mem cl_outbuf;
+ cl_kernel kernel;
AVOpenCLKernelEnv kernel_env;
} DeshakeOpenclContext;
diff --git a/libavfilter/deshake_opencl.c b/libavfilter/deshake_opencl.c
index eea873e..4854e9b 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;
@@ -76,7 +76,7 @@ int ff_opencl_transform(AVFilterContext *ctx,
if (ret < 0)
return ret;
status =
clEnqueueNDRangeKernel(deshake->opencl_ctx.kernel_env.command_queue,
- deshake->opencl_ctx.kernel_env.kernel,
1, NULL,
+ 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));
@@ -108,10 +108,13 @@ 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");
+
+ av_opencl_compile(&deshake->opencl_ctx.kernel_env, "deshake", NULL);
+
+ if (!deshake->opencl_ctx.kernel) {
+ deshake->opencl_ctx.kernel =
clCreateKernel(deshake->opencl_ctx.kernel_env.program,
"avfilter_transform", &ret);
if (ret < 0) {
- av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel for
name 'avfilter_transform'\n");
+ av_log(ctx, AV_LOG_ERROR, "OpenCL: Unable to create kernel
'avfilter_transform'\n");
I think it should use the previous output print: av_log(ctx, AV_LOG_ERROR,
"OpenCL failed to create kernel for name 'avfilter_transform'\n");
return ret;
}
}
@@ -125,7 +128,9 @@ 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.kernel_env.program);
+ deshake->opencl_ctx.kernel_env.command_queue = NULL;
av_opencl_uninit();
}
diff --git a/libavfilter/opencl_allkernels.c
b/libavfilter/opencl_allkernels.c
index 6d80fa8..66b48e1 100644
--- a/libavfilter/opencl_allkernels.c
+++ b/libavfilter/opencl_allkernels.c
@@ -25,17 +25,17 @@
#include "unsharp_opencl_kernel.h"
#endif
-#define OPENCL_REGISTER_KERNEL_CODE(X, x)
\
+#define OPENCL_REGISTER_KERNEL_CODE(X, x, name)
\
{
\
if (CONFIG_##X##_FILTER) {
\
- av_opencl_register_kernel_code(ff_kernel_##x##_opencl);
\
+ av_opencl_register_kernel_code(ff_kernel_##x##_opencl, name);
\
}
\
}
I think the macro should be decleared as below:
#define OPENCL_REGISTER_KERNEL_CODE(X, x)
\
{
\
if (CONFIG_##X##_FILTER) {
\
av_opencl_register_kernel_code(ff_kernel_##x##_opencl, #x);
\
}
\
}
and no need to modify the below code:
void ff_opencl_register_filter_kernel_code_all(void)
{
#if CONFIG_OPENCL
- OPENCL_REGISTER_KERNEL_CODE(DESHAKE, deshake);
- OPENCL_REGISTER_KERNEL_CODE(UNSHARP, unsharp);
+ OPENCL_REGISTER_KERNEL_CODE(DESHAKE, deshake, "deshake");
+ OPENCL_REGISTER_KERNEL_CODE(UNSHARP, unsharp, "unsharp");
#endif
}
diff --git a/libavfilter/unsharp.h b/libavfilter/unsharp.h
index c225929..d754341 100644
--- a/libavfilter/unsharp.h
+++ b/libavfilter/unsharp.h
@@ -35,12 +35,13 @@
typedef struct {
cl_mem cl_luma_mask;
cl_mem cl_chroma_mask;
+ cl_mem cl_inbuf;
+ cl_mem cl_outbuf;
+ cl_kernel kernel;
int in_plane_size[8];
int out_plane_size[8];
int plane_num;
- cl_mem cl_inbuf;
size_t cl_inbuf_size;
- cl_mem cl_outbuf;
size_t cl_outbuf_size;
AVOpenCLKernelEnv kernel_env;
} UnsharpOpenclContext;
diff --git a/libavfilter/unsharp_opencl.c b/libavfilter/unsharp_opencl.c
index b373b66..260056c 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),
@@ -187,7 +187,7 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx,
AVFrame *in, AVFrame *out)
if (ret < 0)
return ret;
status =
clEnqueueNDRangeKernel(unsharp->opencl_ctx.kernel_env.command_queue,
- unsharp->opencl_ctx.kernel_env.kernel,
1, NULL,
+ 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));
@@ -220,10 +220,13 @@ 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");
+
+ av_opencl_compile(&unsharp->opencl_ctx.kernel_env, "unsharp", NULL);
+
+ if (!unsharp->opencl_ctx.kernel) {
+ unsharp->opencl_ctx.kernel =
clCreateKernel(unsharp->opencl_ctx.kernel_env.program, "unsharp", &ret);
+ if (ret != CL_SUCCESS) {
+ av_log(ctx, AV_LOG_ERROR, "OpenCL: Unable to create kernel
'unsharp'\n");
I think it should use the previous output print:av_log(ctx, AV_LOG_ERROR,
"OpenCL failed to create kernel with name 'unsharp'\n");
return ret;
}
}
@@ -232,12 +235,15 @@ int ff_opencl_unsharp_init(AVFilterContext *ctx)
void ff_opencl_unsharp_uninit(AVFilterContext *ctx)
{
+// #define RELEASE( a, f ) do { if( a ) { ocl->f( a ); a = NULL; } }
while( 0 )
UnsharpContext *unsharp = ctx->priv;
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_inbuf);
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.kernel_env.program);
+ unsharp->opencl_ctx.kernel_env.command_queue = NULL;
av_opencl_uninit();
}
diff --git a/libavutil/opencl.c b/libavutil/opencl.c
index e0b28c3..b02480c 100644
--- a/libavutil/opencl.c
+++ b/libavutil/opencl.c
@@ -1,7 +1,8 @@
/*
- * Copyright (C) 2012 Peng Gao <peng at multicorewareinc.com>
- * Copyright (C) 2012 Li Cao <li at multicorewareinc.com>
- * Copyright (C) 2012 Wei Gao <weigao at multicorewareinc.com>
+ * Copyright (C) 2012 Peng Gao <peng at multicorewareinc.com>
+ * Copyright (C) 2012 Li Cao <li at multicorewareinc.com>
+ * Copyright (C) 2012 Wei Gao <weigao at multicorewareinc.com>
+ * Copyright (C) 2013 Lenny Wang <lwanghpc at gmail.com>
*
* This file is part of FFmpeg.
*
@@ -39,13 +40,12 @@ static pthread_mutex_t atomic_opencl_lock =
PTHREAD_MUTEX_INITIALIZER;
#define UNLOCK_OPENCL
#endif
-
-#define MAX_KERNEL_NUM 500
-#define MAX_KERNEL_CODE_NUM 200
+#define MAX_KERNEL_CODE_NUM 32
typedef struct {
int is_compiled;
const char *kernel_string;
+ char program_name[32];
} KernelCode;
typedef struct {
@@ -61,17 +61,13 @@ typedef struct {
int is_user_created;
int platform_idx;
int device_idx;
- char *build_options;
cl_platform_id platform_id;
cl_device_type device_type;
cl_context context;
cl_device_id device_id;
cl_command_queue command_queue;
- int program_count;
- cl_program programs[MAX_KERNEL_CODE_NUM];
- int kernel_code_count;
KernelCode kernel_code[MAX_KERNEL_CODE_NUM];
- int kernel_count;
+ int kernel_code_count;
AVOpenCLDeviceList device_list;
} OpenclContext;
@@ -80,7 +76,6 @@ typedef struct {
static const AVOption opencl_options[] = {
{ "platform_idx", "set platform index value",
OFFSET(platform_idx), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX},
{ "device_idx", "set device index value",
OFFSET(device_idx), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX},
- { "build_options", "build options of opencl",
OFFSET(build_options), AV_OPT_TYPE_STRING, {.str="-I."}, CHAR_MIN,
CHAR_MAX},
{ NULL }
};
@@ -363,7 +358,7 @@ void av_opencl_free_external_env(AVOpenCLExternalEnv
**ext_opencl_env)
av_freep(ext_opencl_env);
}
-int av_opencl_register_kernel_code(const char *kernel_code)
+int av_opencl_register_kernel_code(const char *kernel_code, const char*
name)
{
int i, ret = 0;
LOCK_OPENCL;
@@ -382,71 +377,63 @@ int av_opencl_register_kernel_code(const char
*kernel_code)
}
opencl_ctx.kernel_code[opencl_ctx.kernel_code_count].kernel_string =
kernel_code;
opencl_ctx.kernel_code[opencl_ctx.kernel_code_count].is_compiled = 0;
+
strcpy(opencl_ctx.kernel_code[opencl_ctx.kernel_code_count].program_name,
name);
it should
be:av_strlcpy(opencl_ctx.kernel_code[opencl_ctx.kernel_code_count].program_name,
name,
sizeof(opencl_ctx.kernel_code[opencl_ctx.kernel_code_count].program_name));
+// av_log(&opencl_ctx, AV_LOG_VERBOSE, "program %d registered: %s\n",
opencl_ctx.kernel_code_count,
+//
opencl_ctx.kernel_code[opencl_ctx.kernel_code_count].program_name);
I think you should delete the above two lines
opencl_ctx.kernel_code_count++;
+
end:
UNLOCK_OPENCL;
return ret;
}
-int av_opencl_create_kernel(AVOpenCLKernelEnv *env, const char
*kernel_name)
+int av_opencl_compile(AVOpenCLKernelEnv *env, const char *name, const char
*buildopts)
{
- cl_int status;
int i, ret = 0;
+ cl_int status;
+ int kernel_code_idx = 0;
+ const char *kernel_source;
+ size_t kernel_code_len;
+
LOCK_OPENCL;
- if (strlen(kernel_name) + 1 > AV_OPENCL_MAX_KERNEL_NAME_SIZE) {
- av_log(&opencl_ctx, AV_LOG_ERROR, "Created kernel name %s is too
long\n", kernel_name);
+ for (i = 0; i < opencl_ctx.kernel_code_count; i++) {
+ if (!strcmp(name, opencl_ctx.kernel_code[i].program_name) &&
!opencl_ctx.kernel_code[i].is_compiled) {
+ kernel_source = opencl_ctx.kernel_code[i].kernel_string;
+ kernel_code_len =
strlen(opencl_ctx.kernel_code[i].kernel_string);
+ kernel_code_idx = i;
+ }
+ }
+ if (!kernel_source) {
+ av_log(&opencl_ctx, AV_LOG_ERROR,
+ "OpenCL: unable to find kernel source '%s'\n", name);
I think the module name is not make sence to user, it should be: "OpenCL
unable to find kernel source..........."
ret = AVERROR(EINVAL);
goto end;
}
- if (!env->kernel) {
- if (opencl_ctx.kernel_count >= MAX_KERNEL_NUM) {
- av_log(&opencl_ctx, AV_LOG_ERROR,
- "Could not create kernel with name '%s', maximum number
of kernels %d already reached\n",
- kernel_name, MAX_KERNEL_NUM);
- ret = AVERROR(EINVAL);
- goto end;
- }
- if (opencl_ctx.program_count == 0) {
- av_log(&opencl_ctx, AV_LOG_ERROR, "Program count of OpenCL is
0, can not create kernel\n");
- ret = AVERROR(EINVAL);
- goto end;
- }
- for (i = 0; i < opencl_ctx.program_count; i++) {
- env->kernel = clCreateKernel(opencl_ctx.programs[i],
kernel_name, &status);
- if (status == CL_SUCCESS)
- break;
- }
- if (status != CL_SUCCESS) {
- av_log(&opencl_ctx, AV_LOG_ERROR, "Could not create OpenCL
kernel: %s\n", av_opencl_errstr(status));
- ret = AVERROR_EXTERNAL;
- goto end;
- }
- opencl_ctx.kernel_count++;
- env->command_queue = opencl_ctx.command_queue;
- av_strlcpy(env->kernel_name, kernel_name,
sizeof(env->kernel_name));
- }
-end:
- UNLOCK_OPENCL;
- return ret;
-}
-void av_opencl_release_kernel(AVOpenCLKernelEnv *env)
-{
- cl_int status;
- LOCK_OPENCL;
- if (!env->kernel)
+ /* create a CL program from kernel source */
+ env->command_queue = opencl_ctx.command_queue;
+ env->program = clCreateProgramWithSource(opencl_ctx.context, 1,
&kernel_source, &kernel_code_len, &status);
+ if(status != CL_SUCCESS || !(env->program)) {
+ av_log(&opencl_ctx, AV_LOG_ERROR,
+ "OpenCL: Unable to create program '%s': %s\n", name,
av_opencl_errstr(status));
+ ret = AVERROR(EINVAL);
goto end;
- status = clReleaseKernel(env->kernel);
+ }
+ status = clBuildProgram(env->program, 1, &(opencl_ctx.device_id),
buildopts, NULL, NULL);
if (status != CL_SUCCESS) {
- av_log(&opencl_ctx, AV_LOG_ERROR, "Could not release kernel: %s\n",
- av_opencl_errstr(status));
+ av_log(&opencl_ctx, AV_LOG_ERROR,
+ "OpenCL: Compilation failed with program: %s\n", name);
the same as the previous comment.
+ ret = AVERROR(EINVAL);
+ goto end;
}
- env->kernel = NULL;
- env->command_queue = NULL;
- env->kernel_name[0] = 0;
- opencl_ctx.kernel_count--;
+
+ opencl_ctx.kernel_code[kernel_code_idx].is_compiled = 1;
+// av_log(&opencl_ctx, AV_LOG_VERBOSE, "program %s compiled
successfully: %s\n",
+// opencl_ctx.kernel_code[kernel_code_idx].program_name,
buildopts);
delete the above two lines
+
end:
UNLOCK_OPENCL;
+ return ret;
}
static int init_opencl_env(OpenclContext *opencl_ctx, AVOpenCLExternalEnv
*ext_opencl_env)
@@ -542,49 +529,6 @@ static int init_opencl_env(OpenclContext *opencl_ctx,
AVOpenCLExternalEnv *ext_o
return ret;
}
-static int compile_kernel_file(OpenclContext *opencl_ctx)
-{
- cl_int status;
- int i, kernel_code_count = 0;
- const char *kernel_code[MAX_KERNEL_CODE_NUM] = {NULL};
- size_t kernel_code_len[MAX_KERNEL_CODE_NUM] = {0};
-
- for (i = 0; i < opencl_ctx->kernel_code_count; i++) {
- if (!opencl_ctx->kernel_code[i].is_compiled) {
- kernel_code[kernel_code_count] =
opencl_ctx->kernel_code[i].kernel_string;
- kernel_code_len[kernel_code_count] =
strlen(opencl_ctx->kernel_code[i].kernel_string);
- opencl_ctx->kernel_code[i].is_compiled = 1;
- kernel_code_count++;
- }
- }
- if (!kernel_code_count)
- return 0;
- /* create a CL program using the kernel source */
- opencl_ctx->programs[opencl_ctx->program_count] =
clCreateProgramWithSource(opencl_ctx->context,
-
kernel_code_count,
-
kernel_code,
-
kernel_code_len,
-
&status);
- if(status != CL_SUCCESS) {
- av_log(opencl_ctx, AV_LOG_ERROR,
- "Could not create OpenCL program with source code: %s\n",
av_opencl_errstr(status));
- return AVERROR_EXTERNAL;
- }
- if (!opencl_ctx->programs[opencl_ctx->program_count]) {
- av_log(opencl_ctx, AV_LOG_ERROR, "Created program is NULL\n");
- return AVERROR_EXTERNAL;
- }
- status =
clBuildProgram(opencl_ctx->programs[opencl_ctx->program_count], 1,
&(opencl_ctx->device_id),
- opencl_ctx->build_options, NULL, NULL);
- if (status != CL_SUCCESS) {
- av_log(opencl_ctx, AV_LOG_ERROR,
- "Could not compile OpenCL kernel: %s\n",
av_opencl_errstr(status));
- return AVERROR_EXTERNAL;
- }
- opencl_ctx->program_count++;
- return 0;
-}
-
int av_opencl_init(AVOpenCLExternalEnv *ext_opencl_env)
{
int ret = 0;
@@ -597,18 +541,14 @@ int av_opencl_init(AVOpenCLExternalEnv
*ext_opencl_env)
ret = init_opencl_env(&opencl_ctx, ext_opencl_env);
if (ret < 0)
goto end;
- }
- ret = compile_kernel_file(&opencl_ctx);
- if (ret < 0)
- goto end;
- if (opencl_ctx.kernel_code_count <= 0) {
- av_log(&opencl_ctx, AV_LOG_ERROR,
- "No kernel code is registered, compile kernel file
failed\n");
- ret = AVERROR(EINVAL);
- goto end;
+ if (opencl_ctx.kernel_code_count <= 0) {
+ av_log(&opencl_ctx, AV_LOG_ERROR,
+ "No kernel code is registered, compile kernel file
failed\n");
+ ret = AVERROR(EINVAL);
+ goto end;
+ }
}
opencl_ctx.init_count++;
-
end:
UNLOCK_OPENCL;
return ret;
@@ -617,43 +557,34 @@ end:
void av_opencl_uninit(void)
{
cl_int status;
- int i;
LOCK_OPENCL;
- opencl_ctx.init_count--;
- if (opencl_ctx.is_user_created)
- goto end;
- if (opencl_ctx.init_count > 0 || opencl_ctx.kernel_count > 0)
+ if (opencl_ctx.is_user_created) {
+ av_opt_free(&opencl_ctx);
goto end;
- for (i = 0; i < opencl_ctx.program_count; i++) {
- if (opencl_ctx.programs[i]) {
- status = clReleaseProgram(opencl_ctx.programs[i]);
+ }
+ opencl_ctx.init_count--;
+ if ( opencl_ctx.init_count == 0 )
+ {
I think it should be wrote like this:
opencl_ctx.init_count--;
if (opencl_ctx.is_user_created) {
av_opt_free(&opencl_ctx);
goto end;
}
if (opencl_ctx.init_count > 0)
goto end;
if (opencl_ctx.command_queue) {
..................
.....................
and not use if ( opencl_ctx.init_count == 0 ) to judge.
+ if (opencl_ctx.command_queue) {
+ status = clReleaseCommandQueue(opencl_ctx.command_queue);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
- "Could not release OpenCL program: %s\n",
av_opencl_errstr(status));
+ "Could not release OpenCL command queue: %s\n",
av_opencl_errstr(status));
}
- opencl_ctx.programs[i] = NULL;
+ opencl_ctx.command_queue = NULL;
}
- }
- if (opencl_ctx.command_queue) {
- status = clReleaseCommandQueue(opencl_ctx.command_queue);
- if (status != CL_SUCCESS) {
- av_log(&opencl_ctx, AV_LOG_ERROR,
- "Could not release OpenCL command queue: %s\n",
av_opencl_errstr(status));
- }
- opencl_ctx.command_queue = NULL;
- }
- if (opencl_ctx.context) {
- status = clReleaseContext(opencl_ctx.context);
- if (status != CL_SUCCESS) {
- av_log(&opencl_ctx, AV_LOG_ERROR,
- "Could not release OpenCL context: %s\n",
av_opencl_errstr(status));
+ if (opencl_ctx.context) {
+ status = clReleaseContext(opencl_ctx.context);
+ if (status != CL_SUCCESS) {
+ av_log(&opencl_ctx, AV_LOG_ERROR,
+ "Could not release OpenCL context: %s\n",
av_opencl_errstr(status));
+ }
+ opencl_ctx.context = NULL;
}
- opencl_ctx.context = NULL;
+ free_device_list(&opencl_ctx.device_list);
+ av_opt_free(&opencl_ctx); //FIXME: free openclutils context
I think it should be use the previous code
if (opencl_ctx.init_count <= 0))
av_opt_free(&opencl_ctx); //FIXME: free openclutils context
}
- free_device_list(&opencl_ctx.device_list);
end:
- if ((opencl_ctx.init_count <= 0) && (opencl_ctx.kernel_count <= 0))
- av_opt_free(&opencl_ctx); //FIXME: free openclutils context
UNLOCK_OPENCL;
}
diff --git a/libavutil/opencl.h b/libavutil/opencl.h
index 094c108..2a6e6e5 100644
--- a/libavutil/opencl.h
+++ b/libavutil/opencl.h
@@ -1,7 +1,8 @@
/*
- * Copyright (C) 2012 Peng Gao <peng at multicorewareinc.com>
- * Copyright (C) 2012 Li Cao <li at multicorewareinc.com>
- * Copyright (C) 2012 Wei Gao <weigao at multicorewareinc.com>
+ * Copyright (C) 2012 Peng Gao <peng at multicorewareinc.com>
+ * Copyright (C) 2012 Li Cao <li at multicorewareinc.com>
+ * Copyright (C) 2012 Wei Gao <weigao at multicorewareinc.com>
+ * Copyright (C) 2013 Lenny Wang <lwanghpc at gmail.com>
*
* This file is part of FFmpeg.
*
@@ -41,11 +42,11 @@
#define AV_OPENCL_KERNEL( ... )# __VA_ARGS__
-#define AV_OPENCL_MAX_KERNEL_NAME_SIZE 150
+#define AV_OPENCL_MAX_KERNEL_NAME_SIZE 64
I think you can delete the AV_OPENCL_MAX_KERNEL_NAME_SIZE because no code
use it now
-#define AV_OPENCL_MAX_DEVICE_NAME_SIZE 100
+#define AV_OPENCL_MAX_DEVICE_NAME_SIZE 64
-#define AV_OPENCL_MAX_PLATFORM_NAME_SIZE 100
+#define AV_OPENCL_MAX_PLATFORM_NAME_SIZE 64
typedef struct {
int device_type;
@@ -67,8 +68,7 @@ typedef struct {
typedef struct {
cl_command_queue command_queue;
- cl_kernel kernel;
- char kernel_name[AV_OPENCL_MAX_KERNEL_NAME_SIZE];
+ cl_program program;
} AVOpenCLKernelEnv;
typedef struct {
@@ -107,7 +107,6 @@ void av_opencl_free_device_list(AVOpenCLDeviceList
**device_list);
* av_opencl_init() operation.
*
* The currently accepted options are:
- * - build_options: set options to compile registered kernels code
* - platform: set index of platform in device list
* - device: set index of device in device list
*
@@ -169,13 +168,13 @@ const char *av_opencl_errstr(cl_int status);
* in the runtime environment when av_opencl_init() is called.
*
* @param kernel_code kernel code to be compiled in the OpenCL runtime
environment
+ * @param name pointer to program name used for identification
* @return >=0 on success, a negative error code in case of failure
*/
-int av_opencl_register_kernel_code(const char *kernel_code);
+int av_opencl_register_kernel_code(const char *kernel_code, const char*
name);
I think const "char* program_name" can be more cleare for user.
/**
- * Initialize the run time OpenCL environment and compile the kernel
- * code registered with av_opencl_register_kernel_code().
+ * Initialize the run time OpenCL environment
*
* @param ext_opencl_env external OpenCL environment, created by an
* application program, ignored if set to NULL
@@ -184,14 +183,16 @@ int av_opencl_register_kernel_code(const char
*kernel_code);
int av_opencl_init(AVOpenCLExternalEnv *ext_opencl_env);
/**
- * Create kernel object in the specified kernel environment.
+ * compile specific kernel source with name registered in
+ * av_opencl_register_kernel_code()
*
- * @param env pointer to kernel environment which is filled
with
- * the environment used to run the kernel
- * @param kernel_name kernel function name
+ * @param env pointer to kernel environment of a filter
+ * @param name filter name
+ * @param buildopts pointer to a string that describes the
preprocessor
+ * build options to be used for building the program
* @return >=0 on success, a negative error code in case of failure
*/
-int av_opencl_create_kernel(AVOpenCLKernelEnv *env, const char
*kernel_name);
+int av_opencl_compile(AVOpenCLKernelEnv *env, const char *name, const
char* buildopts);
I think const "char* program_name" can be more cleare for user. also the
buildopts should be renamed to build_opts, the same as the + * @param
buildopts
/**
* Create OpenCL buffer.
@@ -269,14 +270,6 @@ int av_opencl_buffer_read_image(uint8_t **dst_data,
int *plane_size, int plane_n
void av_opencl_buffer_release(cl_mem *cl_buf);
/**
- * Release kernel object.
- *
- * @param env kernel environment where the kernel object was created
- * with av_opencl_create_kernel()
- */
-void av_opencl_release_kernel(AVOpenCLKernelEnv *env);
-
-/**
* Release OpenCL environment.
*
* The OpenCL environment is effectively released only if all the created
Also you should modify the file version.h in libavutil folder.
I think the main idea is that, you make the compile kernel as an API, then
it can make the ffmpeg just compile the kernel which is used, it is good.
And you delete the av_opencl_create_kernel and av_opencl_release_kernel, it
will make user to maintain the kernels they created
anyway, Look good to me
Thanks
>
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel at ffmpeg.org
> http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
>
More information about the ffmpeg-devel
mailing list