64 ctx->command_queue = clCreateCommandQueue(
ctx->ocf.hwctx->context,
65 ctx->ocf.hwctx->device_id,
68 "command queue %d.\n", cle);
70 ctx->kernel_horiz = clCreateKernel(
ctx->ocf.program,
"avgblur_horiz", &cle);
74 ctx->kernel_vert = clCreateKernel(
ctx->ocf.program,
"avgblur_vert", &cle);
82 if (
ctx->command_queue)
83 clReleaseCommandQueue(
ctx->command_queue);
84 if (
ctx->kernel_horiz)
85 clReleaseKernel(
ctx->kernel_horiz);
87 clReleaseKernel(
ctx->kernel_vert);
98 if (
s->radiusV <= 0) {
99 s->radiusV =
s->radiusH;
102 for (
i = 0;
i < 4;
i++) {
122 "filter params: %d.\n", err);
126 s->radius[
Y] =
s->luma_param.radius;
127 s->radius[
U] =
s->radius[
V] =
s->chroma_param.radius;
128 s->radius[
A] =
s->alpha_param.radius;
130 s->power[
Y] =
s->luma_param.power;
131 s->power[
U] =
s->power[
V] =
s->chroma_param.power;
132 s->power[
A] =
s->alpha_param.power;
134 for (
i = 0;
i < 4;
i++) {
135 if (
s->power[
i] == 0) {
153 size_t global_work[2];
154 cl_mem
src, dst, inter;
155 int err, p, radius_x, radius_y,
i;
164 if (!
ctx->initialised) {
169 if (!strcmp(avctx->
filter->
name,
"avgblur_opencl")) {
173 }
else if (!strcmp(avctx->
filter->
name,
"boxblur_opencl")) {
194 dst = (cl_mem) output->
data[p];
195 inter = (cl_mem)intermediate->
data[p];
200 radius_x =
ctx->radiusH;
201 radius_y =
ctx->radiusV;
203 if (!(
ctx->planes & (1 << p))) {
208 for (
i = 0;
i <
ctx->power[p];
i++) {
211 if (!strcmp(avctx->
filter->
name,
"avgblur_opencl")) {
213 }
else if (!strcmp(avctx->
filter->
name,
"boxblur_opencl")) {
218 i == 0 ? intermediate : output, p, 0);
224 p, global_work[0], global_work[1]);
226 cle = clEnqueueNDRangeKernel(
ctx->command_queue,
ctx->kernel_horiz, 2,
NULL,
230 "kernel: %d.\n", cle);
233 i == 0 ? output : intermediate, p, 0);
238 if (!strcmp(avctx->
filter->
name,
"avgblur_opencl")) {
240 }
else if (!strcmp(avctx->
filter->
name,
"boxblur_opencl")) {
244 cle = clEnqueueNDRangeKernel(
ctx->command_queue,
ctx->kernel_vert, 2,
NULL,
248 "kernel: %d.\n", cle);
252 cle = clFinish(
ctx->command_queue);
269 clFinish(
ctx->command_queue);
282 if (
ctx->kernel_horiz) {
283 cle = clReleaseKernel(
ctx->kernel_horiz);
284 if (cle != CL_SUCCESS)
286 "kernel: %d.\n", cle);
289 if (
ctx->kernel_vert) {
290 cle = clReleaseKernel(
ctx->kernel_vert);
291 if (cle != CL_SUCCESS)
293 "kernel: %d.\n", cle);
296 if (
ctx->command_queue) {
297 cle = clReleaseCommandQueue(
ctx->command_queue);
298 if (cle != CL_SUCCESS)
300 "command queue: %d.\n", cle);
328 #define OFFSET(x) offsetof(AverageBlurOpenCLContext, x)
329 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
331 #if CONFIG_AVGBLUR_OPENCL_FILTER
333 static const AVOption avgblur_opencl_options[] = {
344 .
name =
"avgblur_opencl",
347 .priv_class = &avgblur_opencl_class,
359 #if CONFIG_BOXBLUR_OPENCL_FILTER
361 static const AVOption boxblur_opencl_options[] = {
364 {
"luma_power",
"How many times should the boxblur be applied to luma",
OFFSET(luma_param.power),
AV_OPT_TYPE_INT, {.i64=2}, 0, INT_MAX, .flags =
FLAGS },
365 {
"lp",
"How many times should the boxblur be applied to luma",
OFFSET(luma_param.power),
AV_OPT_TYPE_INT, {.i64=2}, 0, INT_MAX, .flags =
FLAGS },
369 {
"chroma_power",
"How many times should the boxblur be applied to chroma",
OFFSET(chroma_param.power),
AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags =
FLAGS },
370 {
"cp",
"How many times should the boxblur be applied to chroma",
OFFSET(chroma_param.power),
AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags =
FLAGS },
374 {
"alpha_power",
"How many times should the boxblur be applied to alpha",
OFFSET(alpha_param.power),
AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags =
FLAGS },
375 {
"ap",
"How many times should the boxblur be applied to alpha",
OFFSET(alpha_param.power),
AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags =
FLAGS },
383 .
name =
"boxblur_opencl",
386 .priv_class = &boxblur_opencl_class,
static int query_formats(AVFilterContext *ctx)
static const AVFilterPad inputs[]
static const AVFilterPad outputs[]
AVFilter ff_vf_avgblur_opencl
AVFilter ff_vf_boxblur_opencl
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Main libavfilter public API header.
const char * ff_opencl_source_avgblur
static av_cold int init(AVCodecContext *avctx)
int ff_boxblur_eval_filter_params(AVFilterLink *inlink, FilterParam *luma_param, FilterParam *chroma_param, FilterParam *alpha_param)
common internal and external API header
static av_cold int uninit(AVCodecContext *avctx)
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
#define AVFILTER_DEFINE_CLASS(fname)
common internal API header
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
static const struct @322 planes[]
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
int ff_opencl_filter_load_program(AVFilterContext *avctx, const char **program_source_array, int nb_strings)
Load a new OpenCL program from strings in memory.
int ff_opencl_filter_config_input(AVFilterLink *inlink)
Check that the input link contains a suitable hardware frames context and extract the device from it.
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx, size_t *work_size, AVFrame *frame, int plane, int block_alignment)
Find the work size needed needed for a given plane of an image.
int ff_opencl_filter_config_output(AVFilterLink *outlink)
Create a suitable hardware frames context for the output.
int ff_opencl_filter_query_formats(AVFilterContext *avctx)
Return that all inputs and outputs support only AV_PIX_FMT_OPENCL.
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
#define CL_FAIL_ON_ERROR(errcode,...)
A helper macro to handle OpenCL errors.
const char * av_get_pix_fmt_name(enum AVPixelFormat pix_fmt)
Return the short name for a pixel format, NULL in case pix_fmt is unknown.
#define FF_ARRAY_ELEMS(a)
const AVFilter * filter
the AVFilter of which this is an instance
void * priv
private data for use by the filter
AVFilterLink ** outputs
array of pointers to output links
A link between two filters.
int w
agreed upon image width
int h
agreed upon image height
AVFilterContext * dst
dest filter
A filter pad used for either input or output.
const char * name
Pad name.
const char * name
Filter name.
This structure describes decoded (raw) audio or video data.
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame.
int format
format of the frame, -1 if unknown or unset Values correspond to enum AVPixelFormat for video frames,...
cl_command_queue command_queue
static const AVFilterPad avgblur_opencl_inputs[]
static const AVFilterPad avgblur_opencl_outputs[]
static int boxblur_opencl_make_filter_params(AVFilterLink *inlink)
static int avgblur_opencl_make_filter_params(AVFilterLink *inlink)
static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
static av_cold void avgblur_opencl_uninit(AVFilterContext *avctx)
static int avgblur_opencl_init(AVFilterContext *avctx)
AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)
Request a picture buffer with a specific set of permissions.