From 03c08d59fb2b998106829063b3840e63c2837923 Mon Sep 17 00:00:00 2001 From: Mark Thompson Date: Wed, 3 Jan 2018 22:43:53 +0000 Subject: [PATCH] doc/filters: Document OpenCL program filters Include some example programs. --- doc/filters.texi | 202 +++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 202 insertions(+) diff --git a/doc/filters.texi b/doc/filters.texi index f6954c947c..a13aef2196 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -12524,6 +12524,136 @@ Set value which will be multiplied with filtered result. Set value which will be added to filtered result. @end table +@anchor{program_opencl} +@section program_opencl + +Filter video using an OpenCL program. + +@table @option + +@item source +OpenCL program source file. + +@item kernel +Kernel name in program. + +@item inputs +Number of inputs to the filter. Defaults to 1. + +@item size, s +Size of output frames. Defaults to the same as the first input. + +@end table + +The program source file must contain a kernel function with the given name, +which will be run once for each plane of the output. Each run on a plane +gets enqueued as a separate 2D global NDRange with one work-item for each +pixel to be generated. The global ID offset for each work-item is therefore +the coordinates of a pixel in the destination image. + +The kernel function needs to take the following arguments: +@itemize +@item +Destination image, @var{__write_only image2d_t}. + +This image will become the output; the kernel should write all of it. +@item +Frame index, @var{unsigned int}. + +This is a counter starting from zero and increasing by one for each frame. +@item +Source images, @var{__read_only image2d_t}. + +These are the most recent images on each input. The kernel may read from +them to generate the output, but they can't be written to. +@end itemize + +Example programs: + +@itemize +@item +Copy the input to the output (output must be the same size as the input). +@verbatim +__kernel void copy(__write_only image2d_t destination, + unsigned int index, + __read_only image2d_t source) +{ + const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE; + + int2 location = (int2)(get_global_id(0), get_global_id(1)); + + float4 value = read_imagef(source, sampler, location); + + write_imagef(destination, location, value); +} +@end verbatim + +@item +Apply a simple transformation, rotating the input by an amount increasing +with the index counter. Pixel values are linearly interpolated by the +sampler, and the output need not have the same dimensions as the input. +@verbatim +__kernel void rotate_image(__write_only image2d_t dst, + unsigned int index, + __read_only image2d_t src) +{ + const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | + CLK_FILTER_LINEAR); + + float angle = (float)index / 100.0f; + + float2 dst_dim = convert_float2(get_image_dim(dst)); + float2 src_dim = convert_float2(get_image_dim(src)); + + float2 dst_cen = dst_dim / 2.0f; + float2 src_cen = src_dim / 2.0f; + + int2 dst_loc = (int2)(get_global_id(0), get_global_id(1)); + + float2 dst_pos = convert_float2(dst_loc) - dst_cen; + float2 src_pos = { + cos(angle) * dst_pos.x - sin(angle) * dst_pos.y, + sin(angle) * dst_pos.x + cos(angle) * dst_pos.y + }; + src_pos = src_pos * src_dim / dst_dim; + + float2 src_loc = src_pos + src_cen; + + if (src_loc.x < 0.0f || src_loc.y < 0.0f || + src_loc.x > src_dim.x || src_loc.y > src_dim.y) + write_imagef(dst, dst_loc, 0.5f); + else + write_imagef(dst, dst_loc, read_imagef(src, sampler, src_loc)); +} +@end verbatim + +@item +Blend two inputs together, with the amount of each input used varying +with the index counter. +@verbatim +__kernel void blend_images(__write_only image2d_t dst, + unsigned int index, + __read_only image2d_t src1, + __read_only image2d_t src2) +{ + const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | + CLK_FILTER_LINEAR); + + float blend = (cos((float)index / 50.0f) + 1.0f) / 2.0f; + + int2 dst_loc = (int2)(get_global_id(0), get_global_id(1)); + int2 src1_loc = dst_loc * get_image_dim(src1) / get_image_dim(dst); + int2 src2_loc = dst_loc * get_image_dim(src2) / get_image_dim(dst); + + float4 val1 = read_imagef(src1, sampler, src1_loc); + float4 val2 = read_imagef(src2, sampler, src2_loc); + + write_imagef(dst, dst_loc, val1 * blend + val2 * (1.0f - blend)); +} +@end verbatim + +@end itemize + @section pseudocolor Alter frame colors in video with pseudocolors. @@ -17498,6 +17628,78 @@ Set the color of the created image. Accepts the same syntax of the corresponding @option{color} option. @end table +@section openclsrc + +Generate video using an OpenCL program. + +@table @option + +@item source +OpenCL program source file. + +@item kernel +Kernel name in program. + +@item size, s +Size of frames to generate. This must be set. + +@item format +Pixel format to use for the generated frames. This must be set. + +@item rate, r +Number of frames generated every second. Default value is '25'. + +@end table + +For details of how the program loading works, see the @ref{program_opencl} +filter. + +Example programs: + +@itemize +@item +Generate a colour ramp by setting pixel values from the position of the pixel +in the output image. (Note that this will work with all pixel formats, but +the generated output will not be the same.) +@verbatim +__kernel void ramp(__write_only image2d_t dst, + unsigned int index) +{ + int2 loc = (int2)(get_global_id(0), get_global_id(1)); + + float4 val; + val.xy = val.zw = convert_float2(loc) / convert_float2(get_image_dim(dst)); + + write_imagef(dst, loc, val); +} +@end verbatim + +@item +Generate a Sierpinski carpet pattern, panning by a single pixel each frame. +@verbatim +__kernel void sierpinski_carpet(__write_only image2d_t dst, + unsigned int index) +{ + int2 loc = (int2)(get_global_id(0), get_global_id(1)); + + float4 value = 0.0f; + int x = loc.x + index; + int y = loc.y + index; + while (x > 0 || y > 0) { + if (x % 3 == 1 && y % 3 == 1) { + value = 1.0f; + break; + } + x /= 3; + y /= 3; + } + + write_imagef(dst, loc, value); +} +@end verbatim + +@end itemize + @c man end VIDEO SOURCES @chapter Video Sinks