From patchwork Fri Jun 16 16:54:44 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Philip Langdale X-Patchwork-Id: 42156 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a05:6a20:c526:b0:117:ac03:c9de with SMTP id gm38csp1691682pzb; Fri, 16 Jun 2023 09:55:08 -0700 (PDT) X-Google-Smtp-Source: ACHHUZ5as+GIsBUpsP8m5V6sIAOjbvqKwINmmniaDIIeYMtreVT1Rd2Je3zGDhdsixPVC+F0OJET X-Received: by 2002:ac2:4d93:0:b0:4f7:6055:8716 with SMTP id g19-20020ac24d93000000b004f760558716mr1804038lfe.4.1686934508419; Fri, 16 Jun 2023 09:55:08 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1686934508; cv=none; d=google.com; s=arc-20160816; b=UJgDWiokWR0Xfl+63oD6Cr+9yJ4zGAShEA5ZFsl1JWGhqQTjJJJ0Es/QTrc3V577UT NxvCT+ukzzV6idhV0WaHMHcvxLUaN9btbJmmgdJ0rZd1Vwmmi8pPotaT4txRxH9NkBOO r2DiApRkljK4hdgpNv71eOJdndUcT5TpgJ1/s9oG20cGTQsJhEZLhNNNxxPGtV9kfRXM gKxSh9Luk0KA7se6Pj4DatGN2H9H19BgFUBMdYAe1Ajov3jg9DwS9fMRnB+p0ocBTd9i 2zZXKqIQUaGw2foUPRSXTmofNZx5N579LG+aMIiS5ITpMQFaXXUBzNNHVVqPzp02YNV1 MLlg== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:content-transfer-encoding:cc:reply-to :list-subscribe:list-help:list-post:list-archive:list-unsubscribe :list-id:precedence:subject:mime-version:message-id:date:to:from :dkim-signature:delivered-to; bh=zB4i6prA9JfoQb4KuJ/RndxYuAtlpGZRvzCufWa+71k=; b=xY4HTW9DOwmDfsM0mIYVOvebZEXCIhWxJQSCdPIzdXaY5/eY5ZHgaM5F9W0QwucYYQ 7csuk11l2v7zknGMUbQmOtV7c5gTQkzS8Hjif9wSgyO6jH9imWig4yQBmTs0ccPAowz9 rrlT1xtSsOGjVFhCeL1fXPmQPm0AVMTU5ArIhn+I9rLDinCR4P7DpIRB93tK1XCqqIT3 lSihVLVRt4Fq/nIsnY4MTRc8cOxY0GA+D1cn/KJpk5EC+oXbWOpNPns95Frvx7WWDSJ8 //R4xm6tt2Q5qjpNIa24uTUkkH12rUIVxLLHADgJxIKE21t1KV46CmZx4c0T5MhPf/H4 jZWw== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@overt.org header.s=mail header.b=c9cgs1J3; 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=QUARANTINE sp=QUARANTINE dis=NONE) header.from=overt.org Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id i7-20020aa7c9c7000000b00514b2a9ce5csi10847768edt.11.2023.06.16.09.55.07; Fri, 16 Jun 2023 09:55:08 -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=@overt.org header.s=mail header.b=c9cgs1J3; 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=QUARANTINE sp=QUARANTINE dis=NONE) header.from=overt.org Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id 081AC68C176; Fri, 16 Jun 2023 19:55:04 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org 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 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: DX61gXDyzQo/ 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) }