From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org [79.124.17.100]) by master.gitmailbox.com (Postfix) with ESMTP id C8C0445ECD for ; Fri, 16 Jun 2023 16:55:06 +0000 (UTC) Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id B5E0968C150; Fri, 16 Jun 2023 19:55:03 +0300 (EEST) Received: from mail.overt.org (mail.overt.org [72.14.183.176]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 6E74D68BFCE for ; Fri, 16 Jun 2023 19:54:56 +0300 (EEST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/simple; d=overt.org; s=mail; t=1686934494; bh=zie+v7babttLWOc7EWOwYneFSi40ak4KvKjECauTB84=; h=From:To:Cc:Subject:Date:From; b=c9cgs1J3orKQXBaHjvSY7wuqnPbf3qsb12lLAUbDcyrpyGNjPI21v6SnJ3VdoCbnE iI0HhyU75GzJ4xswSgFyL86mNzriRbh9CFQeliPZN8i6Rsa8l0BfUpxoKajKzcNrgQ NhibhxF3Bew+rtpvvvlM4Auy4iduSkiLsc4zX2S2nKsVdh31gAMNfwB2w5/Ud4EwD6 ONcBgCmRUfleX8Nxn29++v9rBaNlIjh6Nteoty/qDyQPZqzeEwEbKozxK7bdHYEmd4 bGwcLjthTU0JKKQHQmppFYybafKhD08ubSjTuNNhW3XTXyIY3w99YrIUyPRE8xmGNM yZnlulIcz5okA== Received: from authenticated-user (mail.overt.org [72.14.183.176]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (2048 bits) server-digest SHA256) (No client certificate requested) by mail.overt.org (Postfix) with ESMTPSA id 4A77D60691; Fri, 16 Jun 2023 11:54:54 -0500 (CDT) From: Philip Langdale To: ffmpeg-devel@ffmpeg.org Date: Fri, 16 Jun 2023 09:54:44 -0700 Message-Id: <20230616165444.990190-1-philipl@overt.org> MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH v2] avfilter/scale_cuda: add support for rgb32/bgr32 conversions X-BeenThere: ffmpeg-devel@ffmpeg.org X-Mailman-Version: 2.1.29 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: Philip Langdale Content-Type: text/plain; charset="us-ascii" Content-Transfer-Encoding: 7bit Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Archived-At: List-Archive: List-Post: As we are introducing two new formats and supporting conversions between them, and also with the existing 0rgb32/0bgr32 formats, we get a combinatorial explosion of kernels. I introduced a few new macros to keep the things mostly managable. The conversions are all simple, following existing patterns, with four specific exceptions. When converting from 0rgb32/0bgr32 to rgb32/bgr32, we need to ensure the alpha value is set to 1. In all other cases, it can just be passed through, either to be used or ignored. --- libavfilter/vf_scale_cuda.c | 2 + libavfilter/vf_scale_cuda.cu | 151 ++++++++++++++++++++++++++++------- 2 files changed, 122 insertions(+), 31 deletions(-) diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index 1c99befec8..370cb1d9cd 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -51,6 +51,8 @@ static const enum AVPixelFormat supported_formats[] = { AV_PIX_FMT_YUV444P16, AV_PIX_FMT_0RGB32, AV_PIX_FMT_0BGR32, + AV_PIX_FMT_RGB32, + AV_PIX_FMT_BGR32, }; #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) ) diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index c9c6cafdb6..de06ba9433 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -853,9 +853,67 @@ struct Convert_yuv444p16le_yuv444p16le } }; -// bgr0->X - -struct Convert_bgr0_bgr0 +#define DEF_CONVERT_IDENTITY(fmt1, fmt2)\ + \ +struct Convert_##fmt1##_##fmt2 \ +{ \ + static const int in_bit_depth = 8; \ + typedef uchar4 in_T; \ + typedef uchar in_T_uv; \ + typedef uchar4 out_T; \ + typedef uchar out_T_uv; \ + \ + DEF_F(Convert, out_T) \ + { \ + DEFAULT_DST(0) = SUB_F(y, 0); \ + } \ + \ + DEF_F(Convert_uv, out_T_uv) \ + { \ + } \ +}; \ + +#define DEF_CONVERT_REORDER(fmt1, fmt2) \ + \ +struct Convert_##fmt1##_##fmt2 \ +{ \ + static const int in_bit_depth = 8; \ + typedef uchar4 in_T; \ + typedef uchar in_T_uv; \ + typedef uchar4 out_T; \ + typedef uchar out_T_uv; \ + \ + DEF_F(Convert, out_T) \ + { \ + uchar4 res = SUB_F(y, 0); \ + DEFAULT_DST(0) = make_uchar4( \ + res.z, \ + res.y, \ + res.x, \ + res.w \ + ); \ + } \ + \ + DEF_F(Convert_uv, out_T_uv) \ + { \ + } \ +}; \ + +#define DEF_CONVERT_RGB(fmt1, fmt2) \ + \ +DEF_CONVERT_IDENTITY(fmt1, fmt1) \ +DEF_CONVERT_REORDER (fmt1, fmt2) \ +DEF_CONVERT_REORDER (fmt2, fmt1) \ +DEF_CONVERT_IDENTITY(fmt2, fmt2) + +DEF_CONVERT_RGB(rgb0, bgr0) +DEF_CONVERT_RGB(rgba, bgra) +DEF_CONVERT_IDENTITY(rgba, rgb0) +DEF_CONVERT_IDENTITY(bgra, bgr0) +DEF_CONVERT_REORDER(rgba, bgr0) +DEF_CONVERT_REORDER(bgra, rgb0) + +struct Convert_bgr0_bgra { static const int in_bit_depth = 8; typedef uchar4 in_T; @@ -865,7 +923,13 @@ struct Convert_bgr0_bgr0 DEF_F(Convert, out_T) { - DEFAULT_DST(0) = SUB_F(y, 0); + uchar4 res = SUB_F(y, 0); + DEFAULT_DST(0) = make_uchar4( + res.x, + res.y, + res.z, + 1 + ); } DEF_F(Convert_uv, out_T_uv) @@ -873,7 +937,7 @@ struct Convert_bgr0_bgr0 } }; -struct Convert_bgr0_rgb0 +struct Convert_bgr0_rgba { static const int in_bit_depth = 8; typedef uchar4 in_T; @@ -888,7 +952,7 @@ struct Convert_bgr0_rgb0 res.z, res.y, res.x, - res.w + 1 ); } @@ -897,9 +961,7 @@ struct Convert_bgr0_rgb0 } }; -// rgb0->X - -struct Convert_rgb0_bgr0 +struct Convert_rgb0_bgra { static const int in_bit_depth = 8; typedef uchar4 in_T; @@ -914,7 +976,7 @@ struct Convert_rgb0_bgr0 res.z, res.y, res.x, - res.w + 1 ); } @@ -923,7 +985,7 @@ struct Convert_rgb0_bgr0 } }; -struct Convert_rgb0_rgb0 +struct Convert_rgb0_rgba { static const int in_bit_depth = 8; typedef uchar4 in_T; @@ -933,7 +995,13 @@ struct Convert_rgb0_rgb0 DEF_F(Convert, out_T) { - DEFAULT_DST(0) = SUB_F(y, 0); + uchar4 res = SUB_F(y, 0); + DEFAULT_DST(0) = make_uchar4( + res.x, + res.y, + res.z, + 1 + ); } DEF_F(Convert_uv, out_T_uv) @@ -1117,6 +1185,12 @@ extern "C" { NEAREST_KERNEL_RAW(p016le_ ## C) \ NEAREST_KERNEL_RAW(yuv444p16le_ ## C) +#define NEAREST_KERNELS_RGB(C) \ + NEAREST_KERNEL_RAW(rgb0_ ## C) \ + NEAREST_KERNEL_RAW(bgr0_ ## C) \ + NEAREST_KERNEL_RAW(rgba_ ## C) \ + NEAREST_KERNEL_RAW(bgra_ ## C) \ + NEAREST_KERNELS(yuv420p) NEAREST_KERNELS(nv12) NEAREST_KERNELS(yuv444p) @@ -1124,11 +1198,10 @@ NEAREST_KERNELS(p010le) NEAREST_KERNELS(p016le) NEAREST_KERNELS(yuv444p16le) -NEAREST_KERNEL_RAW(bgr0_bgr0) -NEAREST_KERNEL_RAW(rgb0_rgb0) -NEAREST_KERNEL_RAW(bgr0_rgb0) -NEAREST_KERNEL_RAW(rgb0_bgr0) - +NEAREST_KERNELS_RGB(rgb0) +NEAREST_KERNELS_RGB(bgr0) +NEAREST_KERNELS_RGB(rgba) +NEAREST_KERNELS_RGB(bgra) #define BILINEAR_KERNEL(C, S) \ __global__ void Subsample_Bilinear_##C##S( \ @@ -1152,6 +1225,12 @@ NEAREST_KERNEL_RAW(rgb0_bgr0) BILINEAR_KERNEL_RAW(p016le_ ## C) \ BILINEAR_KERNEL_RAW(yuv444p16le_ ## C) +#define BILINEAR_KERNELS_RGB(C) \ + BILINEAR_KERNEL_RAW(rgb0_ ## C) \ + BILINEAR_KERNEL_RAW(bgr0_ ## C) \ + BILINEAR_KERNEL_RAW(rgba_ ## C) \ + BILINEAR_KERNEL_RAW(bgra_ ## C) + BILINEAR_KERNELS(yuv420p) BILINEAR_KERNELS(nv12) BILINEAR_KERNELS(yuv444p) @@ -1159,10 +1238,10 @@ BILINEAR_KERNELS(p010le) BILINEAR_KERNELS(p016le) BILINEAR_KERNELS(yuv444p16le) -BILINEAR_KERNEL_RAW(bgr0_bgr0) -BILINEAR_KERNEL_RAW(rgb0_rgb0) -BILINEAR_KERNEL_RAW(bgr0_rgb0) -BILINEAR_KERNEL_RAW(rgb0_bgr0) +BILINEAR_KERNELS_RGB(rgb0) +BILINEAR_KERNELS_RGB(bgr0) +BILINEAR_KERNELS_RGB(rgba) +BILINEAR_KERNELS_RGB(bgra) #define BICUBIC_KERNEL(C, S) \ __global__ void Subsample_Bicubic_##C##S( \ @@ -1186,6 +1265,12 @@ BILINEAR_KERNEL_RAW(rgb0_bgr0) BICUBIC_KERNEL_RAW(p016le_ ## C) \ BICUBIC_KERNEL_RAW(yuv444p16le_ ## C) +#define BICUBIC_KERNELS_RGB(C) \ + BICUBIC_KERNEL_RAW(rgb0_ ## C) \ + BICUBIC_KERNEL_RAW(bgr0_ ## C) \ + BICUBIC_KERNEL_RAW(rgba_ ## C) \ + BICUBIC_KERNEL_RAW(bgra_ ## C) + BICUBIC_KERNELS(yuv420p) BICUBIC_KERNELS(nv12) BICUBIC_KERNELS(yuv444p) @@ -1193,11 +1278,10 @@ BICUBIC_KERNELS(p010le) BICUBIC_KERNELS(p016le) BICUBIC_KERNELS(yuv444p16le) -BICUBIC_KERNEL_RAW(bgr0_bgr0) -BICUBIC_KERNEL_RAW(rgb0_rgb0) -BICUBIC_KERNEL_RAW(bgr0_rgb0) -BICUBIC_KERNEL_RAW(rgb0_bgr0) - +BICUBIC_KERNELS_RGB(rgb0) +BICUBIC_KERNELS_RGB(bgr0) +BICUBIC_KERNELS_RGB(rgba) +BICUBIC_KERNELS_RGB(bgra) #define LANCZOS_KERNEL(C, S) \ __global__ void Subsample_Lanczos_##C##S( \ @@ -1221,6 +1305,12 @@ BICUBIC_KERNEL_RAW(rgb0_bgr0) LANCZOS_KERNEL_RAW(p016le_ ## C) \ LANCZOS_KERNEL_RAW(yuv444p16le_ ## C) +#define LANCZOS_KERNELS_RGB(C) \ + LANCZOS_KERNEL_RAW(rgb0_ ## C) \ + LANCZOS_KERNEL_RAW(bgr0_ ## C) \ + LANCZOS_KERNEL_RAW(rgba_ ## C) \ + LANCZOS_KERNEL_RAW(bgra_ ## C) + LANCZOS_KERNELS(yuv420p) LANCZOS_KERNELS(nv12) LANCZOS_KERNELS(yuv444p) @@ -1228,9 +1318,8 @@ LANCZOS_KERNELS(p010le) LANCZOS_KERNELS(p016le) LANCZOS_KERNELS(yuv444p16le) -LANCZOS_KERNEL_RAW(bgr0_bgr0) -LANCZOS_KERNEL_RAW(rgb0_rgb0) -LANCZOS_KERNEL_RAW(bgr0_rgb0) -LANCZOS_KERNEL_RAW(rgb0_bgr0) - +LANCZOS_KERNELS_RGB(rgb0) +LANCZOS_KERNELS_RGB(bgr0) +LANCZOS_KERNELS_RGB(rgba) +LANCZOS_KERNELS_RGB(bgra) } -- 2.39.2 _______________________________________________ ffmpeg-devel mailing list ffmpeg-devel@ffmpeg.org https://ffmpeg.org/mailman/listinfo/ffmpeg-devel To unsubscribe, visit link above, or email ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".