summaryrefslogtreecommitdiff
path: root/libavfilter/vf_overlay_cuda.cu
diff options
context:
space:
mode:
authorYaroslav Pogrebnyak <yyyaroslav@gmail.com>2020-03-19 17:45:48 +0100
committerTimo Rothenpieler <timo@rothenpieler.org>2020-03-28 18:39:40 +0100
commit4976b102d89787d59a3aaa438b76c62aec86ad5a (patch)
treeb436a28e8857c791b6a13dc2619752172bf30cee /libavfilter/vf_overlay_cuda.cu
parent77d5ea1c7cf85d16da330c60351fee3dce3c9e4c (diff)
avfilter: add vf_overlay_cuda
Signed-off-by: Timo Rothenpieler <timo@rothenpieler.org>
Diffstat (limited to 'libavfilter/vf_overlay_cuda.cu')
-rw-r--r--libavfilter/vf_overlay_cuda.cu54
1 files changed, 54 insertions, 0 deletions
diff --git a/libavfilter/vf_overlay_cuda.cu b/libavfilter/vf_overlay_cuda.cu
new file mode 100644
index 0000000000..43ec36c2ed
--- /dev/null
+++ b/libavfilter/vf_overlay_cuda.cu
@@ -0,0 +1,54 @@
+/*
+ * Copyright (c) 2020 Yaroslav Pogrebnyak <yyyaroslav@gmail.com>
+ *
+ * This file is part of FFmpeg.
+ *
+ * FFmpeg is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * FFmpeg is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with FFmpeg; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+ */
+
+extern "C" {
+
+__global__ void Overlay_Cuda(
+ int x_position, int y_position,
+ unsigned char* main, int main_linesize,
+ unsigned char* overlay, int overlay_linesize,
+ int overlay_w, int overlay_h,
+ unsigned char* overlay_alpha, int alpha_linesize,
+ int alpha_adj_x, int alpha_adj_y)
+{
+ int x = blockIdx.x * blockDim.x + threadIdx.x;
+ int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if (x >= overlay_w + x_position ||
+ y >= overlay_h + y_position ||
+ x < x_position ||
+ y < y_position ) {
+
+ return;
+ }
+
+ int overlay_x = x - x_position;
+ int overlay_y = y - y_position;
+
+ float alpha = 1.0;
+ if (alpha_linesize) {
+ alpha = overlay_alpha[alpha_adj_x * overlay_x + alpha_adj_y * overlay_y * alpha_linesize] / 255.0f;
+ }
+
+ main[x + y*main_linesize] = alpha * overlay[overlay_x + overlay_y * overlay_linesize] + (1.0f - alpha) * main[x + y*main_linesize];
+}
+
+}
+