Go to the documentation of this file.
77 const char *kernel_name;
80 if (main_format != xfade_format) {
87 main_planes =
FFMAX(main_planes,
90 ctx->nb_planes = main_planes;
100 ctx->command_queue = clCreateCommandQueue(
ctx->ocf.hwctx->context,
101 ctx->ocf.hwctx->device_id,
104 "command queue %d.\n", cle);
106 switch (
ctx->transition) {
107 case CUSTOM: kernel_name =
ctx->kernel_name;
break;
108 case FADE: kernel_name =
"fade";
break;
109 case WIPELEFT: kernel_name =
"wipeleft";
break;
110 case WIPERIGHT: kernel_name =
"wiperight";
break;
111 case WIPEUP: kernel_name =
"wipeup";
break;
112 case WIPEDOWN: kernel_name =
"wipedown";
break;
113 case SLIDELEFT: kernel_name =
"slideleft";
break;
114 case SLIDERIGHT: kernel_name =
"slideright";
break;
115 case SLIDEUP: kernel_name =
"slideup";
break;
116 case SLIDEDOWN: kernel_name =
"slidedown";
break;
122 ctx->kernel = clCreateKernel(
ctx->ocf.program, kernel_name, &cle);
125 ctx->initialised = 1;
130 if (
ctx->command_queue)
131 clReleaseCommandQueue(
ctx->command_queue);
133 clReleaseKernel(
ctx->kernel);
143 cl_float progress =
av_clipf(1.
f - ((cl_float)(
ctx->pts -
ctx->first_pts -
ctx->offset_pts) /
ctx->duration_pts), 0.f, 1.f);
144 size_t global_work[2];
148 if (!
ctx->initialised) {
155 xfade_fc->sw_format);
166 for (plane = 0; plane <
ctx->nb_planes; plane++) {
170 mem = (cl_mem)
output->data[plane];
174 mem = (cl_mem)
ctx->xf[0]->data[plane];
178 mem = (cl_mem)
ctx->xf[1]->data[plane];
190 cle = clEnqueueNDRangeKernel(
ctx->command_queue,
ctx->kernel, 2,
NULL,
193 "for plane %d: %d.\n", plane, cle);
196 cle = clFinish(
ctx->command_queue);
226 if (inlink0->
w != inlink1->
w || inlink0->
h != inlink1->
h) {
228 "(size %dx%d) do not match the corresponding "
229 "second input link %s parameters (size %dx%d)\n",
238 "(%d/%d) do not match the corresponding "
239 "second input link %s timebase (%d/%d)\n",
269 if (
ctx->xfade_is_over) {
273 }
else if (
ret > 0) {
291 ctx->first_pts =
ctx->xf[0]->pts;
293 ctx->pts =
ctx->xf[0]->pts;
294 if (
ctx->first_pts +
ctx->offset_pts >
ctx->xf[0]->pts) {
296 ctx->need_second = 0;
303 ctx->need_second = 1;
317 ctx->last_pts =
ctx->xf[1]->pts;
318 ctx->pts =
ctx->xf[0]->pts;
319 if (
ctx->xf[0]->pts - (
ctx->first_pts +
ctx->offset_pts) >
ctx->duration_pts)
320 ctx->xfade_is_over = 1;
336 ctx->xfade_is_over = 1;
341 if (!
ctx->eof[0] && !
ctx->xf[0])
343 if (!
ctx->eof[1] && (
ctx->need_second ||
ctx->eof[0]))
345 if (
ctx->eof[0] &&
ctx->eof[1] && (
361 cle = clReleaseKernel(
ctx->kernel);
362 if (cle != CL_SUCCESS)
364 "kernel: %d.\n", cle);
367 if (
ctx->command_queue) {
368 cle = clReleaseCommandQueue(
ctx->command_queue);
369 if (cle != CL_SUCCESS)
371 "command queue: %d.\n", cle);
381 return s->xfade_is_over || !
s->need_second ?
386 #define OFFSET(x) offsetof(XFadeOpenCLContext, x)
387 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
434 .
name =
"xfade_opencl",
437 .priv_class = &xfade_opencl_class,
AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)
Request a picture buffer with a specific set of permissions.
AVPixelFormat
Pixel format.
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a all references to both lists are replaced with a reference to the intersection And when a single format is eventually chosen for a link amongst the remaining all references to the list are updated That means that if a filter requires that its input and output have the same format amongst a supported all it has to do is use a reference to the same list of formats query_formats can leave some formats unset and return AVERROR(EAGAIN) to cause the negotiation mechanism toagain later. That can be used by filters with complex requirements to use the format negotiated on one link to set the formats supported on another. Frame references ownership and permissions
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
#define AVERROR_EOF
End of file.
#define AV_TIME_BASE_Q
Internal time base represented as fractional value.
filter_frame For filters that do not use the this method is called when a frame is pushed to the filter s input It can be called at any time except in a reentrant way If the input frame is enough to produce output
The exact code depends on how similar the blocks are and how related they are to the and needs to apply these operations to the correct inlink or outlink if there are several Macros are available to factor that when no extra processing is inlink
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
#define FILTER_INPUTS(array)
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).
static int xfade_opencl_activate(AVFilterContext *avctx)
@ AV_OPT_TYPE_DURATION
Underlying C type is int64_t.
const char * ff_source_xfade_cl
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.
const char * name
Filter name.
A link between two filters.
Link properties exposed to filter code, but not external callers.
int ff_inlink_consume_frame(AVFilterLink *link, AVFrame **rframe)
Take a frame from the link's FIFO and update the link's stats.
AVFrame * ff_default_get_video_buffer(AVFilterLink *link, int w, int h)
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.
#define FF_FILTER_FORWARD_STATUS_BACK_ALL(outlink, filter)
Forward the status on an output link to all input links.
void * priv
private data for use by the filter
static av_cold void xfade_opencl_uninit(AVFilterContext *avctx)
int ff_opencl_filter_config_output(AVFilterLink *outlink)
Create a suitable hardware frames context for the output.
A filter pad used for either input or output.
AVFilterPad * input_pads
array of input pads
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
static void ff_outlink_set_status(AVFilterLink *link, int status, int64_t pts)
Set the status field of a link from the source filter.
void ff_inlink_request_frame(AVFilterLink *link)
Mark that a frame is wanted on the link.
AVRational sample_aspect_ratio
agreed upon sample aspect ratio
int64_t av_rescale_q(int64_t a, AVRational bq, AVRational cq)
Rescale a 64-bit integer by 2 rational numbers.
#define FILTER_OUTPUTS(array)
AVFrame * ff_inlink_peek_frame(AVFilterLink *link, size_t idx)
Access a frame in the link fifo without consuming it.
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
uint8_t nb_components
The number of components each pixel has, (1-4)
filter_frame For filters that do not use the activate() callback
AVFilterLink ** inputs
array of pointers to input links
int plane
Which of the 4 planes contains the component.
@ AV_PIX_FMT_OPENCL
Hardware surfaces for OpenCL.
int ff_inlink_acknowledge_status(AVFilterLink *link, int *rstatus, int64_t *rpts)
Test and acknowledge the change of status on the link.
size_t ff_inlink_queued_frames(AVFilterLink *link)
Get the number of frames available on the link.
static FilterLink * ff_filter_link(AVFilterLink *link)
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
int(* init)(AVBSFContext *ctx)
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
#define AV_NOPTS_VALUE
Undefined timestamp value.
AVFilterContext * src
source filter
The reader does not expect b to be semantically here and if the code is changed by maybe adding a a division or other the signedness will almost certainly be mistaken To avoid this confusion a new type was SUINT is the C unsigned type but it holds a signed int to use the same example SUINT a
static const AVFilterPad xfade_opencl_outputs[]
it s the only field you need to keep assuming you have a context There is some magic you don t need to care about around this just let it vf offset
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.
AVFrame * ff_null_get_video_buffer(AVFilterLink *link, int w, int h)
static int xfade_opencl_load(AVFilterContext *avctx, enum AVPixelFormat main_format, enum AVPixelFormat xfade_format)
static void uninit(AVBSFContext *ctx)
#define i(width, name, range_min, range_max)
int w
agreed upon image width
static const AVOption xfade_opencl_options[]
const char * name
Pad name.
This struct describes a set or pool of "hardware" frames (i.e.
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
int h
agreed upon image height
@ AV_OPT_TYPE_INT
Underlying C type is int.
int ff_opencl_filter_load_program_from_file(AVFilterContext *avctx, const char *filename)
Load a new OpenCL program from a file.
const AVFilter ff_vf_xfade_opencl
AVComponentDescriptor comp[4]
Parameters that describe how pixels are packed.
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
AVRational time_base
Define the time base used by the PTS of the frames/samples which will pass through this link.
int64_t duration
Duration of the stream, in AV_TIME_BASE fractional seconds.
int ff_outlink_get_status(AVFilterLink *link)
Get the status on an output link.
static AVFrame * get_video_buffer(AVFilterLink *inlink, int w, int h)
cl_command_queue command_queue
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
#define AVERROR_BUG
Internal bug, also see AVERROR_BUG2.
#define CL_FAIL_ON_ERROR(errcode,...)
A helper macro to handle OpenCL errors.
static int xfade_frame(AVFilterContext *avctx, AVFrame *a, AVFrame *b)
the definition of that something depends on the semantic of the filter The callback must examine the status of the filter s links and proceed accordingly The status of output links is stored in the status_in and status_out fields and tested by the ff_outlink_frame_wanted() function. If this function returns true
AVRational frame_rate
Frame rate of the stream on the link, or 1/0 if unknown or variable.
@ AV_OPT_TYPE_STRING
Underlying C type is a uint8_t* that is either NULL or points to a C string allocated with the av_mal...
static int xfade_opencl_config_output(AVFilterLink *outlink)
AVFILTER_DEFINE_CLASS(xfade_opencl)
@ AV_OPT_TYPE_CONST
Special option type for declaring named constants.
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
void ff_filter_set_ready(AVFilterContext *filter, unsigned priority)
Mark a filter ready and schedule it for activation.
static const AVFilterPad xfade_opencl_inputs[]
AVFilterLink ** outputs
array of pointers to output links