From patchwork Fri Jul 20 16:31:21 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Danil Iashchenko X-Patchwork-Id: 9767 Delivered-To: ffmpegpatchwork@gmail.com Received: by 2002:a02:104:0:0:0:0:0 with SMTP id c4-v6csp3039648jad; Fri, 20 Jul 2018 09:32:23 -0700 (PDT) X-Google-Smtp-Source: AAOMgpdSihvLiKvISvB9gJ4Qs5V6udGgoRBakUbLihJxa7CKxxJVPrrAm9cqI5OI4ym9/ES6mM3Z X-Received: by 2002:a1c:9755:: with SMTP id z82-v6mr2029894wmd.104.1532104343504; Fri, 20 Jul 2018 09:32:23 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1532104343; cv=none; d=google.com; s=arc-20160816; b=zgx/sICY34KV3L35xy/lunH8Jc6POORAcv1U4TzA5ecnz+zXcuR6OwW8tQiARKHCz3 hCrao+ZvbQHWOPGkJu0R4zK4zC34MkJrw7/4df3uxh2i9NOSVTUHfW8nnIdA/FytoDqy SEAPjmKhS52KA5CONBbL2HnSP2jdNIZQayjADK0/JZ0UDvdA2r4K43SKh+dIQZPKaVhJ kqwn6AuJNz1cxZ+HTE0UActmZt+xmaba21YdazWjSHr6uFdoq4nQu8LQiFcPwBf6T707 J7IbRw07ZIG3Ktsiue8bi8hEAL4coa2ZTavdMU7L3F3lFgOqjDwEful234LHM2v9O2TQ 5eZQ== 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:cc: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=gA8ipv7FkJ5WtMtXbhPe1AgV/JAkcpOrx4MGaUfuIUY=; b=YGgWp8RXVgHI9dFTi5AhSK4xqoiMgKn0Ok84+gGG0mOKQzFQlOf/QsSFdkCm1h195i zNIzEwUuZBYsUSTtt/FxcobVLJYEHoJRLbov9+BBc5So8qTHGVoXy88dRioAa4BcFvr0 +SDoy6SjHQXzqZ77zb2dkejEZFIdn9YfpcOACYnSNpEg3Gu6qYGltCZ+XotEgmRzBlT7 NhYstBxHBMq2I5n48P887n+PwBdBO0/4WWnG9lDRo1aezZgA5IEq4XJiPxcerhKnLANe Eu+WVS4t//B/JeYbOGqnRvsMJ8Ll9mlExNkqaOXiMqnBXJxGTZ1/0bGS8DfC2O7l0UrR 9/MA== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20161025 header.b=WP2l4civ; 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; dmarc=fail (p=NONE sp=QUARANTINE dis=NONE) header.from=gmail.com Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id 6-v6si1398554wmd.52.2018.07.20.09.32.23; Fri, 20 Jul 2018 09:32:23 -0700 (PDT) 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=@gmail.com header.s=20161025 header.b=WP2l4civ; 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; dmarc=fail (p=NONE sp=QUARANTINE dis=NONE) header.from=gmail.com Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id C44CD68A03A; Fri, 20 Jul 2018 19:32:09 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-lf1-f41.google.com (mail-lf1-f41.google.com [209.85.167.41]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 56AF2680026 for ; Fri, 20 Jul 2018 19:32:03 +0300 (EEST) Received: by mail-lf1-f41.google.com with SMTP id n96-v6so2202765lfi.1 for ; Fri, 20 Jul 2018 09:32:15 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20161025; h=from:to:cc:subject:date:message-id:in-reply-to:references; bh=+goZzV4LByOeMXbLro+wJZL0cGIIu2Q5syF35aM1vew=; b=WP2l4civdtIQZfqlzCwcYHswoh4P0iqLHVeZ0+MbrwAQTjcZrU00Wve2PShHuSH9d5 dWFM1rcEjzeYU7G5GSxJ4v+iir/CACi/RZ2K+Nmu3rhHxG4rDWwsWhJA/HnRC7v9pI0W y5jRelYBU6tk4uVyRZEb0aO97xuWi/DOzs3WGj7lc7GKngPKSL4aSoylRZb3KOXU4gt+ +TEqRncBZnu8qyYuF8xcnk7/FjtxJO7zBErkRdbX91sCm+1H4eNRziBMKIpBMCPwFPry ROOc8M7bWPk3FhWHGlRJNwkHO/1STgqq6ElTrrR3ucLO/+vEvLUQ2MRqBjXC01poAXMR guWw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20161025; h=x-gm-message-state:from:to:cc:subject:date:message-id:in-reply-to :references; bh=+goZzV4LByOeMXbLro+wJZL0cGIIu2Q5syF35aM1vew=; b=Tm9frL14BSkVALV1jN6T/7oaTWs26cXj7vOmdeT1swKcsAeHfA9E6q2ea4S8eu2wdU 5I0V74f+qckJwXlGvooTY7vL3enLUqdxdu+0ZnzjFL8DOMPUXmA32pETXB85WdoVfWhL xUXAd+9eo+xUlAe8eikfUZrudLlb6mi2yPEPILSdx/ec0rmjd0dlruRl7sr4JEfItImK STNplrquNazimjE+Lz4qzQohXEMmY0WRh+OMiLtLgd/8/ZX67HDKAQfBbeIMSW0wpM9c MIVQMiCmdZ2Vwom+PdaKOPGXYGwNG3mZc40KTYA3oYryD0jxOT5Lk5Tr0Uh6xMbPWMpi IUZA== X-Gm-Message-State: AOUpUlH2kNyBuo44K66QlU5VOFRzg4k3PD4cpSufO43TONkoh6sb6IO4 42vZ9GN9x2WZPFnkBTj3n9eC76I= X-Received: by 2002:a19:d7d1:: with SMTP id q78-v6mr1816597lfi.40.1532104334456; Fri, 20 Jul 2018 09:32:14 -0700 (PDT) Received: from dan-acer.lan (campus.ifmo.ru. [194.85.161.2]) by smtp.gmail.com with ESMTPSA id r1-v6sm429356ljr.81.2018.07.20.09.32.13 (version=TLS1_2 cipher=ECDHE-RSA-AES128-SHA bits=128/128); Fri, 20 Jul 2018 09:32:14 -0700 (PDT) From: Danil Iashchenko To: ffmpeg-devel@ffmpeg.org Date: Fri, 20 Jul 2018 19:31:21 +0300 Message-Id: <1532104281-17163-3-git-send-email-danyaschenko@gmail.com> X-Mailer: git-send-email 2.7.4 In-Reply-To: <1532104281-17163-1-git-send-email-danyaschenko@gmail.com> References: <1532104281-17163-1-git-send-email-danyaschenko@gmail.com> Subject: [FFmpeg-devel] [PATCH 2/2] moved program_opencl and openclsrc filters to OpenCL section 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 Cc: Danil Iashchenko MIME-Version: 1.0 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" moved program_opencl and openclsrc filters to OpenCL section --- doc/filters.texi | 404 +++++++++++++++++++++++++++---------------------------- 1 file changed, 202 insertions(+), 202 deletions(-) diff --git a/doc/filters.texi b/doc/filters.texi index d206972..d9458cd 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -13040,136 +13040,6 @@ 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. @@ -17770,6 +17640,78 @@ Apply emboss: @end example @end itemize +@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. + +@subsection 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 + @section overlay_opencl Overlay one video on top of another. @@ -17831,6 +17773,136 @@ Apply the Prewitt operator with scale set to 2 and delta set to 10. @end example @end itemize +@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 + +@subsection 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 roberts_opencl Apply the Roberts cross operator (@url{https://en.wikipedia.org/wiki/Roberts_cross}) to input video stream. @@ -18686,78 +18758,6 @@ 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