88#define BRIEF_PATCH_SIZE 31
89#define BRIEF_PATCH_SIZE_HALF (BRIEF_PATCH_SIZE / 2)
91#define MATCHES_CONTIG_SIZE 2000
93#define ROUNDED_UP_DIV(a, b) ((a + (b - 1)) / b)
310 return (
double)total_time / (
double)num_frames / 1000000.0;
322 double x1 = point_pairs[0].
p.
p1.s[0];
323 double y1 = point_pairs[0].
p.
p1.s[1];
324 double x2 = point_pairs[1].
p.
p1.s[0];
325 double y2 = point_pairs[1].
p.
p1.s[1];
326 double x3 = point_pairs[2].
p.
p1.s[0];
327 double y3 = point_pairs[2].
p.
p1.s[1];
330 double X1 = point_pairs[0].
p.
p2.s[0];
331 double Y1 = point_pairs[0].
p.
p2.s[1];
332 double X2 = point_pairs[1].
p.
p2.s[0];
333 double Y2 = point_pairs[1].
p.
p2.s[1];
334 double X3 = point_pairs[2].
p.
p2.s[0];
335 double Y3 = point_pairs[2].
p.
p2.s[1];
337 double d = 1.0 / ( x1*(y2-y3) + x2*(y3-y1) + x3*(y1-y2) );
339 model[0] = d * ( X1*(y2-y3) + X2*(y3-y1) + X3*(y1-y2) );
340 model[1] = d * ( X1*(x3-x2) + X2*(x1-x3) + X3*(x2-x1) );
341 model[2] = d * ( X1*(x2*y3 - x3*y2) + X2*(x3*y1 - x1*y3) + X3*(x1*y2 - x2*y1) );
343 model[3] = d * ( Y1*(y2-y3) + Y2*(y3-y1) + Y3*(y1-y2) );
344 model[4] = d * ( Y1*(x3-x2) + Y2*(x1-x3) + Y3*(x2-x1) );
345 model[5] = d * ( Y1*(x2*y3 - x3*y2) + Y2*(x3*y1 - x1*y3) + Y3*(x1*y2 - x2*y1) );
353 for (j = 0; j <
i; j++) {
354 double dx1 = points[j]->s[0] - points[
i]->s[0];
355 double dy1 = points[j]->s[1] - points[
i]->s[1];
357 for (k = 0; k < j; k++) {
358 double dx2 = points[k]->s[0] - points[
i]->s[0];
359 double dy2 = points[k]->s[1] - points[
i]->s[1];
364 if (
fabs(dx2*dy1 - dy2*dx1) <= 1.0) {
377 const cl_float2 *prev_points[] = {
378 &pairs_subset[0].
p.
p1,
379 &pairs_subset[1].
p.
p1,
380 &pairs_subset[2].
p.
p1
383 const cl_float2 *curr_points[] = {
384 &pairs_subset[0].
p.
p2,
385 &pairs_subset[1].
p.
p2,
386 &pairs_subset[2].
p.
p2
396 const int num_point_pairs,
401 int i = 0, j, iters = 0;
403 for (; iters < max_attempts; iters++) {
404 for (
i = 0;
i < 3 && iters < max_attempts;) {
408 idx_i = idx[
i] =
rand_in(0, num_point_pairs, alfg);
410 for (j = 0; j <
i; j++) {
411 if (idx_i == idx[j]) {
421 pairs_subset[
i] = point_pairs[idx[
i]];
431 return i == 3 && iters < max_attempts;
437 const int num_point_pairs,
441 double F0 = model[0],
F1 = model[1],
F2 = model[2];
442 double F3 = model[3], F4 = model[4], F5 = model[5];
444 for (
int i = 0;
i < num_point_pairs;
i++) {
445 const cl_float2 *
f = &point_pairs[
i].
p.
p1;
446 const cl_float2 *t = &point_pairs[
i].
p.
p2;
448 double a = F0*
f->s[0] +
F1*
f->s[1] +
F2 - t->s[0];
449 double b =
F3*
f->s[0] + F4*
f->s[1] + F5 - t->s[1];
461 const int num_point_pairs,
466 float t = (
float)(thresh * thresh);
467 int i, n = num_point_pairs, num_inliers = 0;
471 for (
i = 0;
i < n;
i++) {
499 confidence =
av_clipd(confidence, 0.0, 1.0);
500 num_outliers =
av_clipd(num_outliers, 0.0, 1.0);
503 num =
FFMAX(1.0 - confidence, DBL_MIN);
504 denom = 1.0 - pow(1.0 - num_outliers, 3);
505 if (denom < DBL_MIN) {
512 return denom >= 0 || -num >= max_iters * (-denom) ? max_iters : (int)
round(num / denom);
521 const int num_point_pairs,
523 const double threshold,
525 const double confidence
528 double best_model[6], model[6];
531 int iter, niters =
FFMAX(max_iters, 1);
532 int good_count, max_good_count = 0;
535 if (num_point_pairs < 3) {
537 }
else if (num_point_pairs == 3) {
541 for (
int i = 0;
i < 3; ++
i) {
548 for (iter = 0; iter < niters; ++iter) {
549 int found =
get_subset(&deshake_ctx->
alfg, point_pairs, num_point_pairs, pairs_subset, 10000);
562 if (good_count >
FFMAX(max_good_count, 2)) {
563 for (
int mi = 0;
mi < 6; ++
mi) {
564 best_model[
mi] = model[
mi];
567 for (
int pi = 0; pi < 3; pi++) {
568 best_pairs[pi] = pairs_subset[pi];
571 max_good_count = good_count;
574 (
double)(num_point_pairs - good_count) / num_point_pairs,
580 if (max_good_count > 0) {
581 for (
int mi = 0;
mi < 6; ++
mi) {
582 model_out[
mi] = best_model[
mi];
585 for (
int pi = 0; pi < 3; ++pi) {
604 const int num_inliers,
608 float move_x_val = 0.01;
609 float move_y_val = 0.01;
611 float old_move_x_val = 0;
613 int last_changed = 0;
615 for (
int iters = 0; iters < 200; iters++) {
619 best_pairs[0].
p.
p2.s[0] += move_x_val;
621 best_pairs[0].
p.
p2.s[0] += move_y_val;
627 for (
int j = 0; j < num_inliers; j++) {
631 if (total_err < best_err) {
632 for (
int mi = 0;
mi < 6; ++
mi) {
633 model_out[
mi] = model[
mi];
636 best_err = total_err;
637 last_changed = iters;
641 best_pairs[0].
p.
p2.s[0] -= move_x_val;
643 best_pairs[0].
p.
p2.s[0] -= move_y_val;
646 if (iters - last_changed > 4) {
651 old_move_x_val = move_x_val;
659 if (old_move_x_val < 0) {
677 const int num_inliers,
682 float best_err = FLT_MAX;
683 double best_model[6], model[6];
686 for (
int i = 0;
i < max_iters;
i++) {
688 int found =
get_subset(&deshake_ctx->
alfg, inliers, num_inliers, pairs_subset, 10000);
701 for (
int j = 0; j < num_inliers; j++) {
705 if (
i == 0 || total_err < best_err) {
706 for (
int mi = 0;
mi < 6; ++
mi) {
707 best_model[
mi] = model[
mi];
710 for (
int pi = 0; pi < 3; pi++) {
711 best_pairs[pi] = pairs_subset[pi];
714 best_err = total_err;
718 for (
int mi = 0;
mi < 6; ++
mi) {
719 model_out[
mi] = best_model[
mi];
722 for (
int pi = 0; pi < 3; ++pi) {
728 optimize_model(deshake_ctx, best_pairs, inliers, num_inliers, best_err, model_out);
749 memset(&ret, 0,
sizeof(ret));
755 if (
a != 0 ||
b != 0) {
761 ret.
skew.s[0] = atan((
a *
c +
b * d) / (
r *
r));
763 }
else if (
c != 0 || d != 0) {
764 double s = sqrt(
c *
c + d * d);
770 ret.
skew.s[1] = atan((
a *
c +
b * d) / (
s *
s));
784 for (
int i = 0;
i < size_y; ++
i) {
785 for (
int j = 0; j < size_x; ++j) {
804 return 1.0f /
expf(((
float)x * (
float)x) / (2.0f * sigma * sigma));
812 int window_half = length / 2;
814 for (
int i = 0;
i < length; ++
i) {
818 gauss_kernel[
i] =
val;
822 for (
int i = 0;
i < length; ++
i) {
823 gauss_kernel[
i] /= gauss_sum;
850 int clip_start, clip_end, offset_clipped;
894 float new_large_s = 0, new_small_s = 0, new_best = 0, old, diff_between,
895 percent_of_max, inverted_percent;
897 float large_sigma = 40.0f;
898 float small_sigma = 2.0f;
902 best_sigma = (large_sigma - 0.5f) * deshake_ctx->
smooth_percent + 0.5f;
914 for (
int i = indices.
start, j = 0;
i < indices.
end; ++
i, ++j) {
916 new_large_s += old * gauss_kernel[j];
920 for (
int i = indices.
start, j = 0;
i < indices.
end; ++
i, ++j) {
922 new_small_s += old * gauss_kernel[j];
925 diff_between =
fabsf(new_large_s - new_small_s);
926 percent_of_max = diff_between / max_val;
927 inverted_percent = 1 - percent_of_max;
928 best_sigma = large_sigma *
powf(inverted_percent, 40);
932 for (
int i = indices.
start, j = 0;
i < indices.
end; ++
i, ++j) {
934 new_best += old * gauss_kernel[j];
962 float center_s_w, center_s_h;
974 center_s_w = center_w - center_s.s[0];
975 center_s_h = center_h - center_s.s[1];
978 x_shift + center_s_w,
979 y_shift + center_s_h,
995 float new_width, new_height, adjusted_width, adjusted_height, adjusted_x, adjusted_y;
1001 float ar_h = frame_height / frame_width;
1002 float ar_w = frame_width / frame_height;
1038 adjusted_width = new_height * ar_w;
1041 if (adjusted_x >= crop->
top_left.s[0]) {
1044 adjusted_height = new_width * ar_h;
1045 adjusted_y = crop->
bottom_right.s[1] - adjusted_height;
1061 if (
ctx->gauss_kernel)
1064 if (
ctx->ransac_err)
1067 if (
ctx->matches_host)
1070 if (
ctx->matches_contig_host)
1101 if (
ctx->debug_on) {
1120 cl_ulong8 zeroed_ulong8;
1122 cl_image_format grayscale_format;
1123 cl_image_desc grayscale_desc;
1124 cl_command_queue_properties queue_props;
1146 const int descriptor_buf_size = image_grid_32 * (
BREIFN / 8);
1147 const int features_buf_size = image_grid_32 *
sizeof(cl_float2);
1159 ctx->curr_frame = 0;
1161 memset(&zeroed_ulong8, 0,
sizeof(cl_ulong8));
1164 if (!
ctx->gauss_kernel) {
1170 if (!
ctx->ransac_err) {
1179 if (!
ctx->abs_motion.ringbuffers[
i]) {
1185 if (
ctx->debug_on) {
1187 ctx->smooth_window / 2,
1191 if (!
ctx->abs_motion.debug_matches) {
1197 ctx->abs_motion.curr_frame_offset = 0;
1198 ctx->abs_motion.data_start_offset = -1;
1199 ctx->abs_motion.data_end_offset = -1;
1202 if (!pattern_host) {
1208 if (!
ctx->matches_host) {
1214 if (!
ctx->matches_contig_host) {
1220 if (!
ctx->inliers) {
1230 for (
int j = 0; j < 2; ++j) {
1235 pattern_host[
i] = pair;
1238 for (
int i = 0;
i < 14;
i++) {
1239 if (
ctx->sw_format == disallowed_formats[
i]) {
1251 ctx->sw_format = hw_frames_ctx->sw_format;
1257 if (
ctx->debug_on) {
1258 queue_props = CL_QUEUE_PROFILING_ENABLE;
1262 ctx->command_queue = clCreateCommandQueue(
1263 ctx->ocf.hwctx->context,
1264 ctx->ocf.hwctx->device_id,
1281 grayscale_format.image_channel_order = CL_R;
1282 grayscale_format.image_channel_data_type = CL_FLOAT;
1284 grayscale_desc = (cl_image_desc) {
1285 .image_type = CL_MEM_OBJECT_IMAGE2D,
1286 .image_width = outlink->
w,
1287 .image_height = outlink->
h,
1289 .image_array_size = 0,
1290 .image_row_pitch = 0,
1291 .image_slice_pitch = 0,
1292 .num_mip_levels = 0,
1297 ctx->grayscale = clCreateImage(
1298 ctx->ocf.hwctx->context,
1314 CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
1324 if (
ctx->debug_on) {
1329 ctx->initialized = 1;
1343 "\tframe moved from: %f x, %f y\n"
1344 "\t to: %f x, %f y\n"
1345 "\t rotated from: %f degrees\n"
1346 "\t to: %f degrees\n"
1347 "\t scaled from: %f x, %f y\n"
1348 "\t to: %f x, %f y\n"
1350 "\tframe moved by: %f x, %f y\n"
1351 "\t rotated by: %f degrees\n"
1352 "\t scaled by: %f x, %f y\n",
1380 float transform_y[9];
1382 float transform_uv[9];
1384 float transform_crop_y[9];
1386 float transform_crop_uv[9];
1387 float transform_debug_rgb[9];
1388 size_t global_work[2];
1390 cl_mem
src, transformed,
dst;
1393 cl_event transform_event, crop_upscale_event;
1395 cl_int num_model_matches;
1397 const float center_w = (
float)input_frame->
width / 2;
1398 const float center_h = (
float)input_frame->
height / 2;
1404 const float center_w_chroma = (
float)chroma_width / 2;
1405 const float center_h_chroma = (
float)chroma_height / 2;
1407 const float luma_w_over_chroma_w = ((
float)input_frame->
width / (
float)chroma_width);
1408 const float luma_h_over_chroma_h = ((
float)input_frame->
height / (
float)chroma_height);
1508 if (!cropped_frame) {
1514 if (!transformed_frame) {
1524 src = (cl_mem)input_frame->
data[p];
1525 transformed = (cl_mem)transformed_frame->data[p];
1540 { sizeof(cl_mem), &src },
1541 { sizeof(cl_mem), &transformed },
1542 { sizeof(cl_mem), &transforms[p] },
1579 transformed = (cl_mem)transformed_frame->data[0];
1583 (
size_t[]){ debug_matches.num_matches },
1586 { sizeof(cl_mem), &transformed },
1587 { sizeof(cl_mem), &deshake_ctx->debug_matches },
1588 { sizeof(cl_mem), &deshake_ctx->debug_model_matches },
1589 { sizeof(cl_int), &num_model_matches },
1590 { sizeof(cl_mem), &deshake_ctx->transform_y }
1620 crops[0] = deshake_ctx->
crop_y;
1621 crops[1] = crops[2] = deshake_ctx->
crop_uv;
1625 dst = (cl_mem)cropped_frame->
data[p];
1626 transformed = (cl_mem)transformed_frame->data[p];
1640 &crop_upscale_event,
1641 { sizeof(cl_mem), &transformed },
1642 { sizeof(cl_mem), &dst },
1643 { sizeof(cl_float2), &crops[p].top_left },
1644 { sizeof(cl_float2), &crops[p].bottom_right },
1732 int num_inliers = 0;
1736 size_t global_work[2];
1737 size_t harris_global_work[2];
1738 size_t grid_32_global_work[2];
1739 int grid_32_h, grid_32_w;
1740 size_t local_work[2];
1744 cl_event grayscale_event, harris_response_event, refine_features_event,
1745 brief_event, match_descriptors_event, read_buf_event;
1766 grid_32_global_work[0] /= 32;
1767 grid_32_global_work[1] /= 32;
1772 if (deshake_ctx->
is_yuv) {
1775 src = (cl_mem)input_frame->
data[0];
1783 { sizeof(cl_mem), &src },
1784 { sizeof(cl_mem), &deshake_ctx->grayscale }
1793 &harris_response_event,
1794 { sizeof(cl_mem), &deshake_ctx->grayscale },
1795 { sizeof(cl_mem), &deshake_ctx->harris_buf }
1801 grid_32_global_work,
1803 &refine_features_event,
1804 { sizeof(cl_mem), &deshake_ctx->grayscale },
1805 { sizeof(cl_mem), &deshake_ctx->harris_buf },
1806 { sizeof(cl_mem), &deshake_ctx->refined_features },
1807 { sizeof(cl_int), &deshake_ctx->refine_features }
1813 grid_32_global_work,
1816 { sizeof(cl_mem), &deshake_ctx->grayscale },
1817 { sizeof(cl_mem), &deshake_ctx->refined_features },
1818 { sizeof(cl_mem), &deshake_ctx->descriptors },
1819 { sizeof(cl_mem), &deshake_ctx->brief_pattern}
1826 goto no_motion_data;
1832 grid_32_global_work,
1834 &match_descriptors_event,
1835 { sizeof(cl_mem), &deshake_ctx->prev_refined_features },
1836 { sizeof(cl_mem), &deshake_ctx->refined_features },
1837 { sizeof(cl_mem), &deshake_ctx->descriptors },
1838 { sizeof(cl_mem), &deshake_ctx->prev_descriptors },
1839 { sizeof(cl_mem), &deshake_ctx->matches }
1842 cle = clEnqueueReadBuffer(
1857 if (num_vectors < 10) {
1875 goto no_motion_data;
1888 goto no_motion_data;
1891 for (
int i = 0;
i < num_vectors;
i++) {
1906 goto no_motion_data;
1927 if (!deshake_ctx->
is_yuv) {
1946 for (
int i = 0;
i < num_vectors;
i++) {
1953 "\n[ALERT] No motion data found in queue_frame, motion reset to 0\n\n"
1972 if (num_vectors == 0) {
1983 for (
int i = 0;
i < num_vectors;
i++) {
2016 if (!deshake_ctx->
eof) {
2021 if (!
frame->hw_frames_ctx)
2053 deshake_ctx->
eof = 1;
2057 if (deshake_ctx->
eof) {
2072 "Average kernel execution times:\n"
2073 "\t grayscale: %0.3f ms\n"
2074 "\t harris_response: %0.3f ms\n"
2075 "\t refine_features: %0.3f ms\n"
2076 "\tbrief_descriptors: %0.3f ms\n"
2077 "\tmatch_descriptors: %0.3f ms\n"
2078 "\t transform: %0.3f ms\n"
2079 "\t crop_upscale: %0.3f ms\n"
2080 "Average buffer read times:\n"
2081 "\t features buf: %0.3f ms\n",
2097 if (!deshake_ctx->
eof) {
2120#define OFFSET(x) offsetof(DeshakeOpenCLContext, x)
2121#define FLAGS AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM
2125 "tripod",
"simulates a tripod by preventing any camera movement whatsoever "
2126 "from the original frame",
2130 "debug",
"turn on additional debugging information",
2134 "adaptive_crop",
"attempt to subtly crop borders to reduce mirrored content",
2138 "refine_features",
"refine feature point locations at a sub-pixel level",
2142 "smooth_strength",
"smoothing strength (0 attempts to adaptively determine optimal strength)",
2146 "smooth_window_multiplier",
"multiplier for number of frames to buffer for motion data",
2155 .p.name =
"deshake_opencl",
2157 .p.priv_class = &deshake_opencl_class,
uint8_t ptrdiff_t const uint8_t ptrdiff_t int intptr_t intptr_t int int16_t * dst
static double val(void *priv, double ch)
const FFFilter ff_vf_deshake_opencl
simple assert() macros that are a bit more flexible than ISO C assert().
#define av_assert0(cond)
assert() equivalent, that is always enabled.
int ff_inlink_acknowledge_status(AVFilterLink *link, int *rstatus, int64_t *rpts)
Test and acknowledge the change of status on the link.
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
int ff_inlink_consume_frame(AVFilterLink *link, AVFrame **rframe)
Take a frame from the link's FIFO and update the link's stats.
Main libavfilter public API header.
#define i(width, name, range_min, range_max)
static IPT relative(const CmsCtx *ctx, IPT ipt)
common internal and external API header
#define AV_CEIL_RSHIFT(a, b)
static __device__ float fabsf(float a)
static __device__ float fabs(float a)
int(* init)(AVBSFContext *ctx)
AVFrame * ff_framequeue_take(FFFrameQueue *fq)
Take the first frame in the queue.
void ff_framequeue_init(FFFrameQueue *fq, FFFrameQueueGlobal *fqg)
Init a frame queue and attach it to a global structure.
int ff_framequeue_add(FFFrameQueue *fq, AVFrame *frame)
Add a frame.
void ff_framequeue_free(FFFrameQueue *fq)
Free the queue and all queued frames.
void ff_framequeue_global_init(FFFrameQueueGlobal *fqg)
Init a global structure.
static size_t ff_framequeue_queued_frames(const FFFrameQueue *fq)
Get the number of queued frames.
@ AV_OPT_TYPE_FLOAT
Underlying C type is float.
@ AV_OPT_TYPE_BOOL
Underlying C type is int.
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
#define AVERROR_EOF
End of file.
AVFifo * av_fifo_alloc2(size_t nb_elems, size_t elem_size, unsigned int flags)
Allocate and initialize an AVFifo with a given element size.
void av_fifo_freep2(AVFifo **f)
Free an AVFifo and reset pointer to NULL.
size_t av_fifo_can_write(const AVFifo *f)
size_t av_fifo_can_read(const AVFifo *f)
int av_fifo_peek(const AVFifo *f, void *buf, size_t nb_elems, size_t offset)
Read data from a FIFO without modifying FIFO state.
int av_fifo_write(AVFifo *f, const void *buf, size_t nb_elems)
Write data into a FIFO.
void av_fifo_drain2(AVFifo *f, size_t size)
Discard the specified amount of data from an AVFifo.
int av_fifo_read(AVFifo *f, void *buf, size_t nb_elems)
Read data from a FIFO.
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_VERBOSE
Detailed information.
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
static double av_q2d(AVRational a)
Convert an AVRational to a double.
static av_always_inline AVRational av_inv_q(AVRational q)
Invert a rational.
int64_t av_rescale_q(int64_t a, AVRational bq, AVRational cq)
Rescale a 64-bit integer by 2 rational numbers.
static const int8_t transform[32][32]
static const struct @264265015021240254217117147104053036241137144013 transforms[18]
static av_cold void uninit(AVBitStreamFilterContext *ctx)
static int activate(AVBitStreamFilterContext *ctx)
av_cold void av_lfg_init(AVLFG *c, unsigned int seed)
static unsigned int av_lfg_get(AVLFG *c)
Get the next random unsigned 32-bit number using an ALFG.
#define FILTER_INPUTS(array)
#define FILTER_OUTPUTS(array)
#define FF_FILTER_FORWARD_WANTED(outlink, inlink)
Forward the frame_wanted_out flag from an output link to an input link.
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
static void ff_outlink_set_status(AVFilterLink *link, int status, int64_t pts)
Set the status field of a link from the source filter.
#define FFERROR_NOT_READY
Filters implementation helper functions and internal structures.
#define FF_FILTER_FORWARD_STATUS_BACK(outlink, inlink)
Forward the status on an output link to an input link.
static FilterLink * ff_filter_link(AVFilterLink *link)
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
#define AVFILTER_DEFINE_CLASS(fname)
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
static av_const double hypot(double x, double y)
static av_always_inline av_const double round(double x)
Memory handling functions.
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.
cl_ulong ff_opencl_get_event_time(cl_event event)
Gets the command start and end times for the given event and returns the difference (the time that th...
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.
#define CL_RELEASE_KERNEL(k)
release an OpenCL Kernel
#define CL_RUN_KERNEL_WITH_ARGS(queue, kernel, global_work_size, local_work_size, event,...)
Uses the above macro to enqueue the given kernel and then additionally runs it to completion via clFi...
#define CL_CREATE_KERNEL(ctx, kernel_name)
Create a kernel with the given name.
#define CL_CREATE_BUFFER_FLAGS(ctx, buffer_name, flags, size, host_ptr)
Create a buffer with the given information.
#define CL_CREATE_BUFFER(ctx, buffer_name, size)
Create a buffer with the given information.
#define CL_BLOCKING_WRITE_BUFFER(queue, buffer, size, host_ptr, event)
Perform a blocking write to a buffer.
#define CL_RELEASE_MEMORY(m)
release an OpenCL Memory Object
#define CL_RELEASE_QUEUE(q)
release an OpenCL Command Queue
#define CL_FAIL_ON_ERROR(errcode,...)
A helper macro to handle OpenCL errors.
const char * ff_source_deshake_cl
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
#define AV_PIX_FMT_FLAG_RGB
The pixel format contains RGB-like data (as opposed to YUV/grayscale).
#define AV_VIDEO_MAX_PLANES
Maximum number of planes in any pixel format.
AVPixelFormat
Pixel format.
@ AV_PIX_FMT_GBRP10BE
planar GBR 4:4:4 30bpp, big-endian
@ AV_PIX_FMT_GBRP9LE
planar GBR 4:4:4 27bpp, little-endian
@ AV_PIX_FMT_GBRAP12BE
planar GBR 4:4:4:4 48bpp, big-endian
@ AV_PIX_FMT_GBRAP16BE
planar GBRA 4:4:4:4 64bpp, big-endian
@ AV_PIX_FMT_GBRAP
planar GBRA 4:4:4:4 32bpp
@ AV_PIX_FMT_OPENCL
Hardware surfaces for OpenCL.
@ AV_PIX_FMT_GBRP16BE
planar GBR 4:4:4 48bpp, big-endian
@ AV_PIX_FMT_GBRAP12LE
planar GBR 4:4:4:4 48bpp, little-endian
@ AV_PIX_FMT_GBRP9BE
planar GBR 4:4:4 27bpp, big-endian
@ AV_PIX_FMT_GBRP10LE
planar GBR 4:4:4 30bpp, little-endian
@ AV_PIX_FMT_GBRAP10BE
planar GBR 4:4:4:4 40bpp, big-endian
@ AV_PIX_FMT_GBRAP10LE
planar GBR 4:4:4:4 40bpp, little-endian
@ AV_PIX_FMT_GBRP
planar GBR 4:4:4 24bpp
@ AV_PIX_FMT_GBRAP16LE
planar GBRA 4:4:4:4 64bpp, little-endian
@ AV_PIX_FMT_GBRP16LE
planar GBR 4:4:4 48bpp, little-endian
uint8_t * data
The data buffer.
AVFilterLink ** inputs
array of pointers to input links
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
AVRational time_base
Define the time base used by the PTS of the frames/samples which will pass through this link.
AVFilterContext * dst
dest filter
A filter pad used for either input or output.
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.
int64_t duration
Duration of the frame, in the same units as pts.
This struct describes a set or pool of "hardware" frames (i.e.
Context structure for the Lagged Fibonacci PRNG.
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
AVFifo * ringbuffers[RingbufCount]
MotionVector model_matches[3]
cl_kernel kernel_refine_features
unsigned long long harris_response_time
unsigned long long brief_descriptors_time
cl_kernel kernel_brief_descriptors
cl_kernel kernel_transform
unsigned long long read_buf_time
unsigned long long refine_features_time
cl_kernel kernel_match_descriptors
unsigned long long grayscale_time
cl_mem debug_model_matches
cl_mem prev_refined_features
unsigned long long match_descriptors_time
cl_command_queue command_queue
cl_kernel kernel_crop_upscale
MotionVector * matches_host
float smooth_window_multiplier
unsigned long long transform_time
cl_kernel kernel_draw_debug_info
cl_kernel kernel_harris_response
cl_kernel kernel_grayscale
MotionVector * matches_contig_host
unsigned long long crop_upscale_time
AbsoluteFrameMotion abs_motion
Structure to hold global options and statistics for frame queues.
Queue of AVFrame pointers.
Link properties exposed to filter code, but not external callers.
AVRational frame_rate
Frame rate of the stream on the link, or 1/0 if unknown or variable.
AVBufferRef * hw_frames_ctx
For hwaccel pixel formats, this should be a reference to the AVHWFramesContext describing the frames.
#define av_malloc_array(a, b)
static AVFormatContext * ctx
static void optimize_model(DeshakeOpenCLContext *deshake_ctx, MotionVector *best_pairs, MotionVector *inliers, const int num_inliers, float best_err, double *model_out)
#define BRIEF_PATCH_SIZE_HALF
static int deshake_opencl_init(AVFilterContext *avctx)
#define MATCHES_CONTIG_SIZE
static float gaussian_for(int x, float sigma)
static float smooth(DeshakeOpenCLContext *deshake_ctx, float *gauss_kernel, int length, float max_val, AVFifo *values)
static FrameDelta decompose_transform(double *model)
static void run_estimate_kernel(const MotionVector *point_pairs, double *model)
static const AVFilterPad deshake_opencl_outputs[]
static void compute_error(const MotionVector *point_pairs, const int num_point_pairs, const double *model, float *err)
static int filter_frame(AVFilterLink *link, AVFrame *input_frame)
static int minimize_error(DeshakeOpenCLContext *deshake_ctx, MotionVector *inliers, DebugMatches *debug_matches, const int num_inliers, double *model_out, const int max_iters)
static int queue_frame(AVFilterLink *link, AVFrame *input_frame)
static int rand_in(int low, int high, AVLFG *alfg)
static int points_not_collinear(const cl_float2 **points)
static void ringbuf_float_at(DeshakeOpenCLContext *deshake_ctx, AVFifo *values, float *val, int offset)
static int make_vectors_contig(DeshakeOpenCLContext *deshake_ctx, int size_y, int size_x)
static const AVOption deshake_opencl_options[]
static double averaged_event_time_ms(unsigned long long total_time, int num_frames)
#define ROUNDED_UP_DIV(a, b)
static av_cold void deshake_opencl_uninit(AVFilterContext *avctx)
static int check_subset(const MotionVector *pairs_subset)
static void transform_center_scale(float x_shift, float y_shift, float angle, float scale_x, float scale_y, float center_w, float center_h, float *matrix)
static int ransac_update_num_iters(double confidence, double num_outliers, int max_iters)
static int estimate_affine_2d(DeshakeOpenCLContext *deshake_ctx, MotionVector *point_pairs, DebugMatches *debug_matches, const int num_point_pairs, double *model_out, const double threshold, const int max_iters, const double confidence)
static void update_needed_crop(CropInfo *crop, float *transform, float frame_width, float frame_height)
static IterIndices start_end_for(DeshakeOpenCLContext *deshake_ctx, int length)
static int activate(AVFilterContext *ctx)
static const AVFilterPad deshake_opencl_inputs[]
static int get_subset(AVLFG *alfg, const MotionVector *point_pairs, const int num_point_pairs, MotionVector *pairs_subset, int max_attempts)
static void make_gauss_kernel(float *gauss_kernel, float length, float sigma)
static void free_debug_matches(AbsoluteFrameMotion *afm)
static void transform_debug(AVFilterContext *avctx, float *new_vals, float *old_vals, int curr_frame)
static cl_float2 transformed_point(float x, float y, float *transform)
static int find_inliers(MotionVector *point_pairs, const int num_point_pairs, const double *model, float *err, double thresh)
AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)
Request a picture buffer with a specific set of permissions.