From patchwork Wed Jan 3 22:47:22 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Mark Thompson X-Patchwork-Id: 7108 Delivered-To: ffmpegpatchwork@gmail.com Received: by 10.2.79.195 with SMTP id r64csp16445555jad; Wed, 3 Jan 2018 14:47:52 -0800 (PST) X-Google-Smtp-Source: ACJfBotN7EVeVZYxeJ9MsU7AlMr/x92vXxus9oih2Kf+t+NHmcxaFpA22f34hkUu9Wi+hISVIDpc X-Received: by 10.223.184.200 with SMTP id c8mr2805064wrg.268.1515019672546; Wed, 03 Jan 2018 14:47:52 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1515019672; cv=none; d=google.com; s=arc-20160816; b=z6V18tI0VE5d5Wwzz2dYEFVbSJhCvxZ6O96AS4AXYiVEf2tX+wTYY8qQnonDhpN/8h Zeau6wwKGF5ZcwMFPiIEEnhmvHCUcuKDyoopeo8I5JDphp9ETb7fST+oIK2t/UKaRop1 wOLBgIgLdQo87Z12KHdWD5e11jD/FYzEVaXn+XPQ0NJKO8waqLFVDYM/pIU9i3f4BbP3 AfpGIP8HkSplyTp1ebbzIxwt9+IgyB6gVidovbw0nGCYbcPoc5ktjVf4OAFh9HVongBA ETqThorQtgtoMfo5OZbM2i+Lju4paPwgqa60pbEPpmbercCT8WuImsDj0MZ7/FdUFA0b g+IA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:content-transfer-encoding:mime-version:reply-to :list-subscribe:list-help:list-post:list-archive:list-unsubscribe :list-id:precedence:subject:references:in-reply-to:message-id:date :to:from:dkim-signature:delivered-to:arc-authentication-results; bh=y2aS/rOK6oxWu6zvXEpqsdYfOhQ1UPuT16rX8FSrPW0=; b=Vh9nOdmdSLRxHxmKJnI9UsU8yhDq4jgRE7jOnAlL9BnpufYyQjnvuNXe8k3A9O32Le iXeg9TFx94pId5kR2A6iYcSxryWIptZDnbkdkroDdNfj8NdpoJjHqNyAjFTzf5LjYUXZ WAeui9ZV/6kWfultSK06msT+oDGbyoS3JbWVrJjQXb+clO3UxRObTqVS7hHtmmJIKLPj IEbdQ8uQ09lQMg6KAlFHCuTJyusT1t/agDbYUYHyGCQEz76Ym/Cs/AkLBo8htRcYk1/X Q2a94E9PP282H2XCHaCd3s/lcPwzLR1kk9iBoyM2CAUoBG+odbY4AMdQ/LfWeIT9AUwe fKfQ== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@jkqxz-net.20150623.gappssmtp.com header.s=20150623 header.b=bqCp2Zj+; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id 192si1357477wmg.178.2018.01.03.14.47.52; Wed, 03 Jan 2018 14:47:52 -0800 (PST) Received-SPF: pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) client-ip=79.124.17.100; Authentication-Results: mx.google.com; dkim=neutral (body hash did not verify) header.i=@jkqxz-net.20150623.gappssmtp.com header.s=20150623 header.b=bqCp2Zj+; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id CF2946898CF; Thu, 4 Jan 2018 00:47:22 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-wr0-f177.google.com (mail-wr0-f177.google.com [209.85.128.177]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id A6FCA6883E7 for ; Thu, 4 Jan 2018 00:47:16 +0200 (EET) Received: by mail-wr0-f177.google.com with SMTP id l19so3329725wrc.2 for ; Wed, 03 Jan 2018 14:47:33 -0800 (PST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=jkqxz-net.20150623.gappssmtp.com; s=20150623; h=from:to:subject:date:message-id:in-reply-to:references; bh=aTZcaQcnQWJckwxTMua+/rHcKJU+fM6jXVSHfdUUHVA=; b=bqCp2Zj+PlvN07buB2+/+oFixE4pFKlwEacHAm4/1ZBdJowvoZIuU3HYQ25XY6sztG ItX7lweY5G54RieGQmVgkO3Mae8cxB93PizAvVmJhykVyM3Xc9nL7xyCIfmg/Q3+Zgqo /AJv77mh6H6qfDdOuDndtSwCAlwopSwB2fAGPhVLMNc63stpERoK4ahAMOqBloQSbRG4 +YjMUSxxAUItvNnc/xQ28b8PacE6JWHAYRErftx5uicCeox7kuL+LWTbAC0PvMnwXlzh gDhVTSUmNDKAhyhwj9pSkTZLNUyIp6qYuRP7bbT5ArTqiY3dQ3qsVTDoah4/SDWenUJz GY8w== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20161025; h=x-gm-message-state:from:to:subject:date:message-id:in-reply-to :references; bh=aTZcaQcnQWJckwxTMua+/rHcKJU+fM6jXVSHfdUUHVA=; b=fEuXHga0jvKpICbc6yEpS05gLn1tTvaEJXNCI5dcZOxX7MvZYJ4hyrrSE0X+l7MCdc 9JXa9v1FKptxtJ47OnGbcSsqG5MMmIAt8xO29TA9lhF5rC3WxmjyeZi05tFNbUtpXwdU UqLhKG/UuSRrhj0BAJqUS6r6+lRPQyZF105UlV8x0wVD9iTEXMgGgRaLtX2vg6CV8piK FUk7IPRR6O08WCeC0DlUWEObvOolGpNNYKstRMvAJlWpTzONlN7xGBuCODJR0PGeqSiA ZM3VRtJhnGTQPpMbKR93QpCsyUBEf91p+hdKQgS75XsUYjcHxuYafL9vp3N3dmRnsxsJ 9mkQ== X-Gm-Message-State: AKGB3mLEwhmur4ORbIC4ZZ3enxD8Ng3n58ovJP+jXoMpG4t2efFFy1W1 AO58emrhSRBByPrO+J16RyCrg3S/ X-Received: by 10.223.134.115 with SMTP id 48mr2805621wrw.213.1515019652545; Wed, 03 Jan 2018 14:47:32 -0800 (PST) Received: from rywe.jkqxz.net (cpc91242-cmbg18-2-0-cust650.5-4.cable.virginm.net. [82.8.130.139]) by smtp.gmail.com with ESMTPSA id j10sm2983170wrb.49.2018.01.03.14.47.31 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Wed, 03 Jan 2018 14:47:31 -0800 (PST) From: Mark Thompson To: ffmpeg-devel@ffmpeg.org Date: Wed, 3 Jan 2018 22:47:22 +0000 Message-Id: <20180103224722.6040-3-sw@jkqxz.net> X-Mailer: git-send-email 2.11.0 In-Reply-To: <20180103224722.6040-1-sw@jkqxz.net> References: <20180103224722.6040-1-sw@jkqxz.net> Subject: [FFmpeg-devel] [PATCH 3/3] doc/filters: Document OpenCL program filters X-BeenThere: ffmpeg-devel@ffmpeg.org X-Mailman-Version: 2.1.20 Precedence: list List-Id: FFmpeg development discussions and patches List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Reply-To: FFmpeg development discussions and patches MIME-Version: 1.0 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Include some example programs. --- doc/filters.texi | 202 +++++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 202 insertions(+) diff --git a/doc/filters.texi b/doc/filters.texi index 477f833449..d2ddb2f0c0 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -12442,6 +12442,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. @@ -17416,6 +17546,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