From patchwork Tue Sep 10 18:10:55 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Koushik Dutta X-Patchwork-Id: 51498 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a05:612c:41b1:b0:48e:c0f8:d0de with SMTP id le49csp584252vqb; Tue, 10 Sep 2024 11:11:19 -0700 (PDT) X-Forwarded-Encrypted: i=2; AJvYcCWSI/EYbPm1rCK4Xzgvv+EkYHiCvhL7LAkY+CRZnl2fQ6ro2sqQ/zrz3uqGmolNJrgtt/RtDmvlZ0oCRylij3Au@gmail.com X-Google-Smtp-Source: AGHT+IE1Tbmt2ttaweNTMDS69qGB1P7aK2gK04LOeSeS9t7zo2BKC2BUN6r18Qn5xxokE9t23GiP X-Received: by 2002:a2e:a542:0:b0:2f7:5c23:98fc with SMTP id 38308e7fff4ca-2f75c23991emr22131251fa.11.1725991879369; Tue, 10 Sep 2024 11:11:19 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1725991879; cv=none; d=google.com; s=arc-20240605; b=XxlqRgzgofmo4SCcC9VVLU8+tIq3alFfv3oc8jwayt5BP5hlLl9ym0+1TsPhZmwkIP CaX2yZf4/wRDtUq9SL+LbFGlZLIkvZQAEBaxLFOTFkmIz6LNqrPj4ZAp6imcJ9Ysh7D5 SjbSsKqSbezLISyj1NsP75e7Z1jEeF3TSAg8BwNiNTQ2JFl/i7APhZ0uG6Li3l73KWyd SBDlc80K6rOpT/TpxPSMGTFaj1FfVI8Zd013U0XLX1WLUnMiq9odFt7Dk4BiIF7aNObS JFeweDlryLWafxEKNjaIP7EEAHc2tmFecLiAsQi/7no5/E2Xx8eLECO6q1qM+sD8oJfW TtOg== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20240605; 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=PuFD74xKf9jsesj50wQvutIoK1VmODd62lxjPv9NgSA=; fh=vf1rgm8KpEMxcRKiIfYPC5TR7MMWdHSIxg51vjgwMpE=; b=I4vMzN8Ht7q2mS+PFXqlo+X3twCPgucaNdFNG7rvVAC8TCcdozNK2VxEGfVAWebp4T YZMsBawou14+Dne3fzuZjiq5lCJ0IsW3KT0jW7o+QSYmNkTkaLiWEyjgN/kDv6XO7SLk Yf4wA4732sCdPDG7vej+VWMhRt2nF8YzuAzDOQU5a25ntHPnf29/g+ITVCHWOejapeEQ l26BhtEqbWOxRi3TFbSC9x9ch89eVElTUL2xCrgAf0strtryZ6rwlcJJcm7PYJfB41Cp NpF2B8QwMFquqsaFK8PTJ4r3sP5vMS9vQarpaYjgPHsvlpJzHQpWkDdB2/ycos/7vzei +meg==; dara=google.com ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20230601 header.b=YXPDub0G; 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; dara=fail header.i=@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 38308e7fff4ca-2f75c0b897esi23426181fa.638.2024.09.10.11.11.18; Tue, 10 Sep 2024 11:11:19 -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=20230601 header.b=YXPDub0G; 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; dara=fail header.i=@gmail.com Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id CC44D68E0AF; Tue, 10 Sep 2024 21:11:13 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pl1-f175.google.com (mail-pl1-f175.google.com [209.85.214.175]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 8E68B68DE63 for ; Tue, 10 Sep 2024 21:11:07 +0300 (EEST) Received: by mail-pl1-f175.google.com with SMTP id d9443c01a7336-20688fbaeafso59644255ad.0 for ; Tue, 10 Sep 2024 11:11:07 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20230601; t=1725991865; x=1726596665; darn=ffmpeg.org; h=content-transfer-encoding:mime-version:message-id:date:subject:cc :to:from:from:to:cc:subject:date:message-id:reply-to; bh=WkdzRWa7iClx/3SfHY+PO/C9USqFSrwXVyTKzq4vcgk=; b=YXPDub0GV/phBUu3Y453m+tfkKhkvSxKIhzNbVuLJvM+lvapmV8M94kviQXle9LeIv aca/x5op3xfMmjuSewJoPNDBNKz8BUBZ2j49xy0d9SHoFX0Z75V9aAMv8IYYkcrCcamN Nvkx86wsAEfx5cXZnWHj+Kx7O0chdSEB3HMG/ncAmWi7UkmziQPKsFqAuywaYIHP18fg QvvgL3+f4cGHiNB8LwDWGAffreDrmrZpazc0ZtElk3dVD5UW5E51bjYLzgh9mXu9GUDM UmkIx44jMNCOQVtLKG3MKi+epusyf8MchGXxBr6HDURlNPwwS9bgSPaR+Rbi+luvJggW b78Q== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1725991865; x=1726596665; h=content-transfer-encoding:mime-version:message-id:date:subject:cc :to:from:x-gm-message-state:from:to:cc:subject:date:message-id :reply-to; bh=WkdzRWa7iClx/3SfHY+PO/C9USqFSrwXVyTKzq4vcgk=; b=C9swcwvzo/c4rIlJfN+eszn+cn9qPMuSX1MugFACPTsb/bnbqJwQ84JYeimfWH24TJ 8TzOVTVbS/Kd9NOu6HghGteJHzJBUb0o6MySQKRNnKx+/2AQaGDpN/vQ6LYVihB0d22j psFxZXWzdOoDl0yQm8WktsRLSNoBnnOLtQqaEA/DQJP1vju80d/P/18pf4RZpuCbGaNF 0LnzasB/O1795MqVV14PDEW5hGLgm1RsJx1X3ViALH5nFkrENt96E5yDRXn38YQf/Fuu i4uSnkOnGirP08n90CTU9ktSibDmpnoxyONDaBDMTvd76RR65cvM3qXQZJzVeuMTI70U OI4g== X-Gm-Message-State: AOJu0YzVaGxNEdQSxs5i9Z4YzeOmSKnGsMh3pB5OcJA4n3vLmg7PK/ID iN9cn7H5mjUupGWFPebv/UCfLQfiJL5SiyWkRydaUh5kjgu/jheGBZ3nY+BQ X-Received: by 2002:a17:902:f68e:b0:207:182c:8a52 with SMTP id d9443c01a7336-2074c79fa6amr18690735ad.58.1725991865236; Tue, 10 Sep 2024 11:11:05 -0700 (PDT) Received: from Koushik-MacStudio.tail05204.ts.net ([2001:559:76c:0:ed2e:ef7a:2bf7:c169]) by smtp.gmail.com with ESMTPSA id d9443c01a7336-20710e1d5c3sm51265025ad.59.2024.09.10.11.11.04 (version=TLS1_3 cipher=TLS_CHACHA20_POLY1305_SHA256 bits=256/256); Tue, 10 Sep 2024 11:11:04 -0700 (PDT) From: Koushik Dutta To: ffmpeg-devel@ffmpeg.org Date: Tue, 10 Sep 2024 11:10:55 -0700 Message-Id: <20240910181057.43453-1-koushd@gmail.com> X-Mailer: git-send-email 2.39.3 (Apple Git-146) MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH 1/3] scale_vt frame crop support 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: Koushik Dutta Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: 2JAY4AVPh6gY The crop filter has no effect on scale_vt: -vf crop=100:100,scale_vt=300x300 Hardware frames (AV_PIX_FMT_FLAG_HWACCEL) are expected to use the crop_* properties, as seen in the implementation vf_crop.c. The current workaround is to hwdownload the full frame and perform the crop on CPU. --- libavfilter/vf_scale_vt.c | 50 +++++++++++++++++++++++++++++++++++++++ 1 file changed, 50 insertions(+) diff --git a/libavfilter/vf_scale_vt.c b/libavfilter/vf_scale_vt.c index 05f4e7b797..3da46a6cd5 100644 --- a/libavfilter/vf_scale_vt.c +++ b/libavfilter/vf_scale_vt.c @@ -109,6 +109,8 @@ static av_cold int scale_vt_init(AVFilterContext *avctx) VTSessionSetProperty(s->transfer, kVTPixelTransferPropertyKey_DestinationYCbCrMatrix, value); } + VTSessionSetProperty(s->transfer, kVTPixelTransferPropertyKey_ScalingMode, kVTScalingMode_CropSourceToCleanAperture); + return 0; } @@ -132,6 +134,18 @@ static int scale_vt_filter_frame(AVFilterLink *link, AVFrame *in) CVPixelBufferRef src; CVPixelBufferRef dst; + int left; + int top; + int width; + int height; + CFNumberRef crop_width_num; + CFNumberRef crop_height_num; + CFNumberRef crop_offset_left_num; + CFNumberRef crop_offset_top_num; + const void *clean_aperture_keys[4]; + const void *source_clean_aperture_values[4]; + CFDictionaryRef source_clean_aperture; + AVFrame *out = ff_get_video_buffer(outlink, outlink->w, outlink->h); if (!out) { ret = AVERROR(ENOMEM); @@ -153,8 +167,43 @@ static int scale_vt_filter_frame(AVFilterLink *link, AVFrame *in) if (s->colour_matrix != AVCOL_SPC_UNSPECIFIED) out->colorspace = s->colour_matrix; + width = (in->width - in->crop_right) - in->crop_left; + height = (in->height - in->crop_bottom) - in->crop_top; + // The crop offsets are relative to the center of the frame. + // the crop width and crop height are relative to the center of the crop rect, not top left as normal. + left = in->crop_left - in->width / 2 + width / 2; + top = in->crop_top - in->height / 2 + height / 2; + crop_width_num = CFNumberCreate(kCFAllocatorDefault, kCFNumberIntType, &width); + crop_height_num = CFNumberCreate(kCFAllocatorDefault, kCFNumberIntType, &height); + crop_offset_left_num = CFNumberCreate(kCFAllocatorDefault, kCFNumberIntType, &left); + crop_offset_top_num = CFNumberCreate(kCFAllocatorDefault, kCFNumberIntType, &top); + + clean_aperture_keys[0] = kCVImageBufferCleanApertureWidthKey; + clean_aperture_keys[1] = kCVImageBufferCleanApertureHeightKey; + clean_aperture_keys[2] = kCVImageBufferCleanApertureHorizontalOffsetKey; + clean_aperture_keys[3] = kCVImageBufferCleanApertureVerticalOffsetKey; + + source_clean_aperture_values[0] = crop_width_num; + source_clean_aperture_values[1] = crop_height_num; + source_clean_aperture_values[2] = crop_offset_left_num; + source_clean_aperture_values[3] = crop_offset_top_num; + + source_clean_aperture = CFDictionaryCreate(kCFAllocatorDefault, + clean_aperture_keys, + source_clean_aperture_values, + 4, + &kCFTypeDictionaryKeyCallBacks, + &kCFTypeDictionaryValueCallBacks); + + CFRelease(crop_width_num); + CFRelease(crop_height_num); + CFRelease(crop_offset_left_num); + CFRelease(crop_offset_top_num); + src = (CVPixelBufferRef)in->data[3]; dst = (CVPixelBufferRef)out->data[3]; + CVBufferSetAttachment(src, kCVImageBufferCleanApertureKey, + source_clean_aperture, kCVAttachmentMode_ShouldPropagate); ret = VTPixelTransferSessionTransferImage(s->transfer, src, dst); if (ret != noErr) { av_log(ctx, AV_LOG_ERROR, "transfer image failed, %d\n", ret); @@ -162,6 +211,7 @@ static int scale_vt_filter_frame(AVFilterLink *link, AVFrame *in) goto fail; } + CFRelease(source_clean_aperture); av_frame_free(&in); return ff_filter_frame(outlink, out); From patchwork Tue Sep 10 18:10:56 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Koushik Dutta X-Patchwork-Id: 51499 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a05:612c:41b1:b0:48e:c0f8:d0de with SMTP id le49csp584329vqb; Tue, 10 Sep 2024 11:11:29 -0700 (PDT) X-Forwarded-Encrypted: i=2; AJvYcCXqaH+rnUWwrusab0EzCop/ct4HTDiILcvmearlAJ+ircjLSIYJiqrUuTj3KFRDTXajtywsUYMtPaDo3kxYziOJ@gmail.com X-Google-Smtp-Source: AGHT+IEgDvKTkAI6xs60PD+XxkwPaWZCBxGf74hKj0ooLHihHXv8/wyC4XjglkrUnAr8itr11Q9e X-Received: by 2002:a17:907:da9:b0:a8d:3705:4115 with SMTP id a640c23a62f3a-a8ffab88e1amr131836966b.32.1725991889276; Tue, 10 Sep 2024 11:11:29 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1725991889; cv=none; d=google.com; s=arc-20240605; b=BFnPAdjHuYucV5TkuTAHI7WqlTuoukLTI/a/sMvg68Z9RE7NVVxrRLJo6BAkMmIxYu hPNmgSojs8/YwvfyrUvKM6xZ6ViNqOLomueTC7823Za7BJcGp+kqYtuwCm3yLasReKWI MM3vsqaUyu2Zf0JrPj2LL2Bqqq2C1lLox9j9jfddQcHViei6+d2fKgxV+YsVwywiaRJC 6U5yAWk/1JS6yy0PkoxguOgi95d2HhSsQWzVTzlykTdxnbSgRZNJ0TAtTbmyqCRYxhQD v0Z1aUhJZZNYwIfRBMk+DPYJF6voOmMY9ilHeDPsBS1X90qg1Xv2RA0UBlXmrp8P2e5o JcvA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20240605; 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:references:in-reply-to :message-id:date:to:from:dkim-signature:delivered-to; bh=1vH/lZ5BAyZJO3B/pHlRF2B7TzfA+z7GL54tFcsMsWw=; fh=vf1rgm8KpEMxcRKiIfYPC5TR7MMWdHSIxg51vjgwMpE=; b=carUk/BdtJap0xe3nbPA2viZ5AKIcyreu3nZmnOuscGgTMWF29IpXXtVSDcZOlpig2 ea+oYvMHfKrUoqIwRFpL97FCiE8mFCoZeo8Aw6r/vmm5keU6uyS3BmqwZ05VIE9d3TdF zXi8QpRQ0npvB/QdcGRBCAxbJbT6nsIzdBKeu/qcfLQAMKALA9CZryNu3IHvs7jasZfF I4wLqbvzwtz2XwI29c3muaIxAKFtdkmDAWGvF/TQQ4C/emVxb9HXGNRM2XUuz0m/TBXj ZI18VrU3dxS1GaTegFX2j8vEUyW0ZGFFOtm9u9IoBbizrNACG0AwubPCI9EohmwxyDiw e6HQ==; dara=google.com ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20230601 header.b=nsYWAo64; 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; dara=fail header.i=@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 a640c23a62f3a-a8d25856a1bsi600285666b.134.2024.09.10.11.11.27; Tue, 10 Sep 2024 11:11:29 -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=20230601 header.b=nsYWAo64; 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; dara=fail header.i=@gmail.com Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id D678A68E103; Tue, 10 Sep 2024 21:11:15 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pl1-f179.google.com (mail-pl1-f179.google.com [209.85.214.179]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id E895C68DE63 for ; Tue, 10 Sep 2024 21:11:07 +0300 (EEST) Received: by mail-pl1-f179.google.com with SMTP id d9443c01a7336-206aee40676so48816925ad.0 for ; Tue, 10 Sep 2024 11:11:07 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20230601; t=1725991866; x=1726596666; darn=ffmpeg.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to; bh=/a+xezM4SQBRsfa3hqwtSLDM6yfR+1CgNY0KmfUqP1I=; b=nsYWAo64vqpwWKp4LQy0bNtjvJdnQkfgNWVgXnrQnveaB/kjordZS/P+lPXh9FbQRU pAMgk5OvG6rv/exl5y5E1S3TRWVy7exom2Z9cM1t7LLbCIt+k3DIQfpH7+RPJfV01T/X ih7pk21B96tGfuF88v22clx0WWw2X/T204ruaBCV/bPnziOAG60pZqDlRSSOOm4/xcau UgeGITBUiCidixviFARAEv3Kvk/Os11m1DKQ104LyWF9lbbkesrT+jAzigL4BWfANASe VVd9jDsD6r4BwHo9oPIcnxxfaAW6n8nRLPlA78UFz9gbZRXTXm2ziGT4lZ1wTEo8Zc1H Znyw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1725991866; x=1726596666; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:x-gm-message-state:from:to:cc :subject:date:message-id:reply-to; bh=/a+xezM4SQBRsfa3hqwtSLDM6yfR+1CgNY0KmfUqP1I=; b=PPSWiW2HZreDgCJn36G/6WY19ss/VRGjeK+x5fT5raBBovHHagnu9MA8Rx3GBHFI97 rE517tSnoxVemlGDoeM8OAJoEtKHi9jTbrHwj/N15AvWn5wIHRS0iRdY+6CIpqaVgH+R KMILEbD9WvfR23TaRg/ZHDvt0auGM/JE5MQlxrEHoegxIEL4wagawYKyAIRMll5yoWPB w00D/V4LgU+oFWWU/AWbi19B1EjXe83wOWyMjUcoJ+U4+JvrDw/vvbKYdMdbD8EuEJCw 1bCC9s/8tKjx2JCi7WZEiNLXm2x7Ztt5Al2R/9+0dDEEN0mhT1U/u4CBy9vaSkEKqBXe KGpQ== X-Gm-Message-State: AOJu0YyrjuQR5xTF9pGy3uKYxLxAwGo88Il01Tn76hW7mdzmwKAD9ViO jHOE+Jpnvho/yFNK7u89E/M2jMbcKoYDXWfaxcg5MRzZOjEBNDkPEXkp/swe X-Received: by 2002:a17:902:e846:b0:206:c8dc:e334 with SMTP id d9443c01a7336-2074c69b72emr21799775ad.39.1725991865858; Tue, 10 Sep 2024 11:11:05 -0700 (PDT) Received: from Koushik-MacStudio.tail05204.ts.net ([2001:559:76c:0:ed2e:ef7a:2bf7:c169]) by smtp.gmail.com with ESMTPSA id d9443c01a7336-20710e1d5c3sm51265025ad.59.2024.09.10.11.11.05 (version=TLS1_3 cipher=TLS_CHACHA20_POLY1305_SHA256 bits=256/256); Tue, 10 Sep 2024 11:11:05 -0700 (PDT) From: Koushik Dutta To: ffmpeg-devel@ffmpeg.org Date: Tue, 10 Sep 2024 11:10:56 -0700 Message-Id: <20240910181057.43453-2-koushd@gmail.com> X-Mailer: git-send-email 2.39.3 (Apple Git-146) In-Reply-To: <20240910181057.43453-1-koushd@gmail.com> References: <20240910181057.43453-1-koushd@gmail.com> MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH 2/3] scale_cuda frame crop support 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: Koushik Dutta Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: azowR2aNFA+H The crop filter has no effect on scale_cuda: -vf crop=100:100,scale_cuda=300x300 Hardware frames (AV_PIX_FMT_FLAG_HWACCEL) are expected to use the crop_* properties, as seen in the implementation vf_crop.c. The current workaround is to hwdownload the full frame and perform the crop on CPU. --- libavfilter/vf_scale_cuda.c | 15 ++++++++++----- libavfilter/vf_scale_cuda.cu | 22 ++++++++++++++-------- 2 files changed, 24 insertions(+), 13 deletions(-) diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index 54a340949d..eb8beee771 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -407,7 +407,7 @@ fail: } static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, - CUtexObject src_tex[4], int src_width, int src_height, + CUtexObject src_tex[4], int src_left, int src_top, int src_width, int src_height, AVFrame *out_frame, int dst_width, int dst_height, int dst_pitch) { CUDAScaleContext *s = ctx->priv; @@ -422,7 +422,7 @@ static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, &src_tex[0], &src_tex[1], &src_tex[2], &src_tex[3], &dst_devptr[0], &dst_devptr[1], &dst_devptr[2], &dst_devptr[3], &dst_width, &dst_height, &dst_pitch, - &src_width, &src_height, &s->param + &src_left, &src_top, &src_width, &src_height, &s->param }; return CHECK_CU(cu->cuLaunchKernel(func, @@ -440,6 +440,9 @@ static int scalecuda_resize(AVFilterContext *ctx, CUtexObject tex[4] = { 0, 0, 0, 0 }; + int crop_width = (in->width - in->crop_right) - in->crop_left; + int crop_height = (in->height - in->crop_bottom) - in->crop_top; + ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx)); if (ret < 0) return ret; @@ -477,7 +480,7 @@ static int scalecuda_resize(AVFilterContext *ctx, // scale primary plane(s). Usually Y (and A), or single plane of RGB frames. ret = call_resize_kernel(ctx, s->cu_func, - tex, in->width, in->height, + tex, in->crop_left, in->crop_top, crop_width, crop_height, out, out->width, out->height, out->linesize[0]); if (ret < 0) goto exit; @@ -485,8 +488,10 @@ static int scalecuda_resize(AVFilterContext *ctx, if (s->out_planes > 1) { // scale UV plane. Scale function sets both U and V plane, or singular interleaved plane. ret = call_resize_kernel(ctx, s->cu_func_uv, tex, - AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w), - AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h), + AV_CEIL_RSHIFT(in->crop_left, s->in_desc->log2_chroma_w), + AV_CEIL_RSHIFT(in->crop_top, s->in_desc->log2_chroma_h), + AV_CEIL_RSHIFT(crop_width, s->in_desc->log2_chroma_w), + AV_CEIL_RSHIFT(crop_height, s->in_desc->log2_chroma_h), out, AV_CEIL_RSHIFT(out->width, s->out_desc->log2_chroma_w), AV_CEIL_RSHIFT(out->height, s->out_desc->log2_chroma_h), diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index de06ba9433..271b55cd5d 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -26,6 +26,7 @@ template using subsample_function_t = T (*)(cudaTextureObject_t tex, int xo, int yo, int dst_width, int dst_height, + int src_left, int src_top, int src_width, int src_height, int bit_depth, float param); @@ -64,11 +65,12 @@ static inline __device__ ushort conv_16to10(ushort in) subsample_function_t subsample_func_uv> \ __device__ static inline void N(cudaTextureObject_t src_tex[4], T *dst[4], int xo, int yo, \ int dst_width, int dst_height, int dst_pitch, \ - int src_width, int src_height, float param) + int src_left, int src_top, int src_width, int src_height, float param) #define SUB_F(m, plane) \ subsample_func_##m(src_tex[plane], xo, yo, \ dst_width, dst_height, \ + src_left, src_top, \ src_width, src_height, \ in_bit_depth, param) @@ -1063,13 +1065,14 @@ template __device__ static inline T Subsample_Nearest(cudaTextureObject_t tex, int xo, int yo, int dst_width, int dst_height, + int src_left, int src_top, int src_width, int src_height, int bit_depth, float param) { float hscale = (float)src_width / (float)dst_width; float vscale = (float)src_height / (float)dst_height; - float xi = (xo + 0.5f) * hscale; - float yi = (yo + 0.5f) * vscale; + float xi = (xo + 0.5f) * hscale + src_left; + float yi = (yo + 0.5f) * vscale + src_top; return tex2D(tex, xi, yi); } @@ -1078,13 +1081,14 @@ template __device__ static inline T Subsample_Bilinear(cudaTextureObject_t tex, int xo, int yo, int dst_width, int dst_height, + int src_left, int src_top, int src_width, int src_height, int bit_depth, float param) { float hscale = (float)src_width / (float)dst_width; float vscale = (float)src_height / (float)dst_height; - float xi = (xo + 0.5f) * hscale; - float yi = (yo + 0.5f) * vscale; + float xi = (xo + 0.5f) * hscale + src_left; + float yi = (yo + 0.5f) * vscale + src_top; // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv} float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f); float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f); @@ -1109,13 +1113,14 @@ template __device__ static inline T Subsample_Bicubic(cudaTextureObject_t tex, int xo, int yo, int dst_width, int dst_height, + int src_left, int src_top, int src_width, int src_height, int bit_depth, float param) { float hscale = (float)src_width / (float)dst_width; float vscale = (float)src_height / (float)dst_height; - float xi = (xo + 0.5f) * hscale - 0.5f; - float yi = (yo + 0.5f) * vscale - 0.5f; + float xi = (xo + 0.5f) * hscale - 0.5f + src_left; + float yi = (yo + 0.5f) * vscale - 0.5f + src_top; float px = floor(xi); float py = floor(yi); float fx = xi - px; @@ -1147,7 +1152,7 @@ __device__ static inline T Subsample_Bicubic(cudaTextureObject_t tex, cudaTextureObject_t src_tex_2, cudaTextureObject_t src_tex_3, \ T *dst_0, T *dst_1, T *dst_2, T *dst_3, \ int dst_width, int dst_height, int dst_pitch, \ - int src_width, int src_height, float param + int src_left, int src_top, int src_width, int src_height, float param #define SUBSAMPLE(Convert, T) \ cudaTextureObject_t src_tex[4] = \ @@ -1159,6 +1164,7 @@ __device__ static inline T Subsample_Bicubic(cudaTextureObject_t tex, Convert( \ src_tex, dst, xo, yo, \ dst_width, dst_height, dst_pitch, \ + src_left, src_top, \ src_width, src_height, param); extern "C" { From patchwork Tue Sep 10 18:10:57 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Koushik Dutta X-Patchwork-Id: 51500 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a05:612c:41b1:b0:48e:c0f8:d0de with SMTP id le49csp600069vqb; Tue, 10 Sep 2024 11:49:11 -0700 (PDT) X-Forwarded-Encrypted: i=2; AJvYcCUoCd6vT9cdky6wyJ4M5R0yRs49lq8o0OI9PRPQFxWuomhREp5DsB1hcIV8fnz+CSn6W8Ytbdwfq3p9nQPuzKDo@gmail.com X-Google-Smtp-Source: AGHT+IFpB7cQ5KTDudQcKYgbdiza4jmWt/UsvVWuhcRtV3T4JOWi3L6Uk4zt3/K9EC1GEjR30i+u X-Received: by 2002:a2e:bea8:0:b0:2f6:5fa7:2626 with SMTP id 38308e7fff4ca-2f77b824201mr4891981fa.39.1725994151151; Tue, 10 Sep 2024 11:49:11 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1725994151; cv=none; d=google.com; s=arc-20240605; b=ZKzHuLQtjET4ZoWhfO8sW0raA99qm/CZHDNYLILaVJvrRCkQ3jBSlmFJu2N10n9n4v RguYwsbLI7TTvYgG8lGdyhdgtQCVHu1iUKlKJPJ7YY6G13vDR4tE8mOMSPwbu+9q+hDs 4pGxGHCFcZkL5sNfPtsbI4m7dg2JHA2ncmyje/VKALI5nG+mOGX8nkRAUDZhLpiavFE3 7AEk4bRu4QtKZd/+B9h82KnU5GOEvfDew5/vyYIc032mv/+gmU9h3zy5BWK7esKo1HYs NPHMvbW3glurvlxJU1wvvJ78XWQcL2jZbxoTraUpmdD9ceZBi9cjJy+niZkBlB/atbC5 7VVg== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20240605; 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:references:in-reply-to :message-id:date:to:from:dkim-signature:delivered-to; bh=e6+dzRTJaUmPyLMMSc59woJyRnbh2A4nnz3M7wO1N9M=; fh=vf1rgm8KpEMxcRKiIfYPC5TR7MMWdHSIxg51vjgwMpE=; b=EvbrSxJz/mjnfKJdSZG4YHqrj0lrq2NFJjS+hqMJOJt4iB2zUU8xsyoyimoUo0W/sS kjvAdYv3gOMxGAl9Vyuh2FmvgJqk2m4yE8dUAPGkvPb8U5LbzcYJgGlJoVCg4cS34Jqf lvlCkjUmzWWudt1znlCzHlJnpW+rD/8DzoU2SPU6f2zGvMvpsfUnigoOGBZeOoidZRmq YStpgf8ghJEZ+ckqFmf1JMKeBP6XQdOxOq9eIpFoDJO7QAWWlWzbDTXIMseGd++Qx60A bL/Q630lULCRmHAqxCbpO+4M3CKYu1RlX+vP9/2mmZOKQfSsuyfVo/F2lD9apZEZAzSk SWBw==; dara=google.com ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20230601 header.b=gI1GvQia; 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; dara=fail header.i=@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 4fb4d7f45d1cf-5c3ebd9fd86si5915293a12.452.2024.09.10.11.49.10; Tue, 10 Sep 2024 11:49:11 -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=20230601 header.b=gI1GvQia; 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; dara=fail header.i=@gmail.com Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id 1069368E0F8; Tue, 10 Sep 2024 21:11:17 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pl1-f181.google.com (mail-pl1-f181.google.com [209.85.214.181]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 556FF68DE63 for ; Tue, 10 Sep 2024 21:11:08 +0300 (EEST) Received: by mail-pl1-f181.google.com with SMTP id d9443c01a7336-2057835395aso11863655ad.3 for ; Tue, 10 Sep 2024 11:11:08 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20230601; t=1725991866; x=1726596666; darn=ffmpeg.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to; bh=K396oibq7K4lsT2yqw+N2PLt5dZ6SBVzqKrdBJ3ij7s=; b=gI1GvQiaMbsfuKmTw/6BNVck+wNG5+/ThBPcMR59FBLtMp1E92Af/Q1YQjQI7SohWX ILDjs8DW5X//I2XOg6sPs8hu4xAnaAzfHIqca5tlyqw+p74nifdjGioPuR+qBolrVCtV xgszTqDaSMiDSAO4RiWydmLjU05sLUsIr7ibQlEWbtxBtKjMk9HKmxT0EEaIhaigGE8p o3yF1NGjAVZD6sysUpUv5QOaFBJHB0aiflAoaDdl2jy5YqFtxMQ7a4QWhwb9j8OAAfq8 OFSo2oCTfJjGaDRp4gf87HPLl6NFX421AVUwA+jeG9RbBqadqNxEgFHdVfTu8V1bE88p 22ig== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1725991866; x=1726596666; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:x-gm-message-state:from:to:cc :subject:date:message-id:reply-to; bh=K396oibq7K4lsT2yqw+N2PLt5dZ6SBVzqKrdBJ3ij7s=; b=rDpsuocrH7emOs9A35yFLaSBuIKIeV/evwgNcElSeyFmuLlx2t9UM37jSsyMN6utd8 A4YS+MtHjIg6y8fZJsxaBsD/X713G51SPiW3ULgmhfomkSuAcMMbYTJD6ZFUmBkIe1rd D0YY6xjD98DJmwioscASZzxckOOzaJo295gz25I3vpIgwFZ5xA6Y1zHnpbvd3KUTAa7b 1gk6KYKKRU3twJIUY88a+gK5iW7f07iHZNzkSg00POIzqyoOI7Z3STIAWDf4WTp24DI9 jIXej8/Lk8/zDy2GBWjLNY83WktPKfN0zlFvuG3vhwFNix2HScRNksH3xHxKNomsH/J2 l+7Q== X-Gm-Message-State: AOJu0YyPOxznK/6IcehuCm/7jeM/OOSYeuiYmJVxJ1lUB0YGHUZsQyPe 8gZhxwxQAIXoSmKzwJ9oujnl7PlMLzSFmE9zmUkNEYe9qBVWjPgNStiekSnX X-Received: by 2002:a17:902:dac8:b0:205:5d12:3f24 with SMTP id d9443c01a7336-20752196ec1mr5125645ad.20.1725991866487; Tue, 10 Sep 2024 11:11:06 -0700 (PDT) Received: from Koushik-MacStudio.tail05204.ts.net ([2001:559:76c:0:ed2e:ef7a:2bf7:c169]) by smtp.gmail.com with ESMTPSA id d9443c01a7336-20710e1d5c3sm51265025ad.59.2024.09.10.11.11.06 (version=TLS1_3 cipher=TLS_CHACHA20_POLY1305_SHA256 bits=256/256); Tue, 10 Sep 2024 11:11:06 -0700 (PDT) From: Koushik Dutta To: ffmpeg-devel@ffmpeg.org Date: Tue, 10 Sep 2024 11:10:57 -0700 Message-Id: <20240910181057.43453-3-koushd@gmail.com> X-Mailer: git-send-email 2.39.3 (Apple Git-146) In-Reply-To: <20240910181057.43453-1-koushd@gmail.com> References: <20240910181057.43453-1-koushd@gmail.com> MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH 3/3] scale_qsv frame crop support 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: Koushik Dutta Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: KETrX1+x+DPC The crop filter has no effect on scale_qsv: -vf crop=100:100,scale_qsv=300x300 Hardware frames (AV_PIX_FMT_FLAG_HWACCEL) are expected to use the crop_* properties, as seen in the implementation vf_crop.c. This patch is slightly different from the previously submitted patches since qsv supports frame crop via the vpp_qsv filter. If no explicit crop is detected via that filter, AVFrame.crop_* will be used instead. Removal of vpp_qsv's crop argument may be warranted. --- libavfilter/qsvvpp.c | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/libavfilter/qsvvpp.c b/libavfilter/qsvvpp.c index 0818ada117..05ce387bc3 100644 --- a/libavfilter/qsvvpp.c +++ b/libavfilter/qsvvpp.c @@ -467,6 +467,16 @@ static QSVFrame *submit_frame(QSVVPPContext *s, AVFilterLink *inlink, AVFrame *p else if (qsv_frame->frame->repeat_pict == 4) qsv_frame->surface.Info.PicStruct |= MFX_PICSTRUCT_FRAME_TRIPLING; + // if crop arguments are not present from the vpp_qsv filter, use the provided AVFrame + // crop_* members instead. + if (!qsv_frame->surface.Info.CropX && !qsv_frame->surface.Info.CropY + && qsv_frame->surface.Info.CropW == picref->width && qsv_frame->surface.Info.CropH == picref->height) { + qsv_frame->surface.Info.CropW = (mfxU16)((picref->width - picref->crop_right) - picref->crop_left); + qsv_frame->surface.Info.CropH = (mfxU16)((picref->height - picref->crop_bottom) - picref->crop_top); + qsv_frame->surface.Info.CropX = (mfxU16)picref->crop_left; + qsv_frame->surface.Info.CropY = (mfxU16)picref->crop_top; + } + return qsv_frame; }