import Oracle_OSS gnome-remote-desktop-47.3-3.el10_1
This commit is contained in:
AlmaLinux RelEng Bot
parent
8da900636f
commit
ffc78ae67d
606
0001-hwaccel-nvidia-Fix-image-corruptions-on-Blackwell-GP.patch
Normal file
606
0001-hwaccel-nvidia-Fix-image-corruptions-on-Blackwell-GP.patch
Normal file
@ -0,0 +1,606 @@
|
||||
From eb4a7aebfd005d3a82b6bd63ee864bef941d9285 Mon Sep 17 00:00:00 2001
|
||||
From: Stephan Seitz <sseitz@nvidia.com>
|
||||
Date: Fri, 29 Aug 2025 11:31:20 +0200
|
||||
Subject: [PATCH] hwaccel-nvidia: Fix image corruptions on Blackwell GPUs
|
||||
|
||||
Nvidia devices never supported `NV_ENC_PARAMS_FRAME_FIELD_MODE_MBAFF`.
|
||||
Though the code worked for past, it causes interlaced line artifacts on
|
||||
Blackwell GPUs.
|
||||
|
||||
Using `NV_ENC_PARAMS_FRAME_FIELD_MODE_FRAME` and changing
|
||||
`grd-cuda-avc-utils` to produce progressive frames works on
|
||||
Blackwell and on older GPU generations.
|
||||
|
||||
Clang was used instead of nvcc, to regenerate the device kernel.
|
||||
`clang --cuda-gpu-arch=sm_30 -S src/grd-cuda-avc-utils.cu -o data/grd-cuda-avc-utils_30.ptx --no-cuda-version-check -O3 --cuda-device-only`
|
||||
|
||||
Closes https://gitlab.gnome.org/GNOME/gnome-remote-desktop/-/issues/263
|
||||
|
||||
Part-of: <https://gitlab.gnome.org/GNOME/gnome-remote-desktop/-/merge_requests/344>
|
||||
---
|
||||
data/grd-cuda-avc-utils_30.ptx | 453 ++++++++++++++++++---------------
|
||||
src/grd-cuda-avc-utils.cu | 55 +---
|
||||
src/grd-hwaccel-nvidia.c | 2 +-
|
||||
3 files changed, 256 insertions(+), 254 deletions(-)
|
||||
|
||||
diff --git a/data/grd-cuda-avc-utils_30.ptx b/data/grd-cuda-avc-utils_30.ptx
|
||||
index b14b193c..62f6d5f2 100644
|
||||
--- a/data/grd-cuda-avc-utils_30.ptx
|
||||
+++ b/data/grd-cuda-avc-utils_30.ptx
|
||||
@@ -1,224 +1,267 @@
|
||||
//
|
||||
-// Generated by NVIDIA NVVM Compiler
|
||||
-//
|
||||
-// Compiler Build ID: CL-27506705
|
||||
-// Cuda compilation tools, release 10.2, V10.2.89
|
||||
-// Based on LLVM 3.4svn
|
||||
+// Generated by LLVM NVPTX Back-End
|
||||
//
|
||||
|
||||
-.version 6.5
|
||||
+.version 8.7
|
||||
.target sm_30
|
||||
.address_size 64
|
||||
|
||||
- // .globl convert_2x2_bgrx_area_to_yuv420_nv12
|
||||
+ // .globl rgb_to_y // -- Begin function rgb_to_y
|
||||
+ // @rgb_to_y
|
||||
+.visible .func (.param .b32 func_retval0) rgb_to_y(
|
||||
+ .param .b32 rgb_to_y_param_0,
|
||||
+ .param .b32 rgb_to_y_param_1,
|
||||
+ .param .b32 rgb_to_y_param_2
|
||||
+)
|
||||
+{
|
||||
+ .reg .b16 %rs<4>;
|
||||
+ .reg .b32 %r<7>;
|
||||
+
|
||||
+// %bb.0:
|
||||
+ ld.param.u8 %rs1, [rgb_to_y_param_0];
|
||||
+ ld.param.u8 %rs2, [rgb_to_y_param_1];
|
||||
+ mul.wide.u16 %r1, %rs1, 54;
|
||||
+ ld.param.u8 %rs3, [rgb_to_y_param_2];
|
||||
+ mul.wide.u16 %r2, %rs2, 183;
|
||||
+ add.s32 %r3, %r2, %r1;
|
||||
+ mul.wide.u16 %r4, %rs3, 18;
|
||||
+ add.s32 %r5, %r3, %r4;
|
||||
+ shr.u32 %r6, %r5, 8;
|
||||
+ st.param.b32 [func_retval0], %r6;
|
||||
+ ret;
|
||||
+ // -- End function
|
||||
+}
|
||||
+ // .globl rgb_to_u // -- Begin function rgb_to_u
|
||||
+.visible .func (.param .b32 func_retval0) rgb_to_u(
|
||||
+ .param .b32 rgb_to_u_param_0,
|
||||
+ .param .b32 rgb_to_u_param_1,
|
||||
+ .param .b32 rgb_to_u_param_2
|
||||
+) // @rgb_to_u
|
||||
+{
|
||||
+ .reg .b16 %rs<7>;
|
||||
+ .reg .b32 %r<7>;
|
||||
+
|
||||
+// %bb.0:
|
||||
+ ld.param.u8 %rs1, [rgb_to_u_param_0];
|
||||
+ ld.param.u8 %rs2, [rgb_to_u_param_1];
|
||||
+ mul.wide.u16 %r1, %rs1, -29;
|
||||
+ ld.param.u8 %rs3, [rgb_to_u_param_2];
|
||||
+ mul.wide.u16 %r2, %rs2, -99;
|
||||
+ add.s32 %r3, %r2, %r1;
|
||||
+ mul.wide.u16 %r4, %rs3, 128;
|
||||
+ add.s32 %r5, %r3, %r4;
|
||||
+ cvt.u16.u32 %rs4, %r5;
|
||||
+ shr.u16 %rs5, %rs4, 8;
|
||||
+ xor.b16 %rs6, %rs5, 128;
|
||||
+ cvt.u32.u16 %r6, %rs6;
|
||||
+ st.param.b32 [func_retval0], %r6;
|
||||
+ ret;
|
||||
+ // -- End function
|
||||
+}
|
||||
+ // .globl rgb_to_v // -- Begin function rgb_to_v
|
||||
+.visible .func (.param .b32 func_retval0) rgb_to_v(
|
||||
+ .param .b32 rgb_to_v_param_0,
|
||||
+ .param .b32 rgb_to_v_param_1,
|
||||
+ .param .b32 rgb_to_v_param_2
|
||||
+) // @rgb_to_v
|
||||
+{
|
||||
+ .reg .b16 %rs<7>;
|
||||
+ .reg .b32 %r<7>;
|
||||
|
||||
+// %bb.0:
|
||||
+ ld.param.u8 %rs1, [rgb_to_v_param_0];
|
||||
+ ld.param.u8 %rs2, [rgb_to_v_param_1];
|
||||
+ mul.wide.u16 %r1, %rs1, 128;
|
||||
+ ld.param.u8 %rs3, [rgb_to_v_param_2];
|
||||
+ mul.wide.u16 %r2, %rs2, -116;
|
||||
+ add.s32 %r3, %r2, %r1;
|
||||
+ mul.wide.u16 %r4, %rs3, -12;
|
||||
+ add.s32 %r5, %r3, %r4;
|
||||
+ cvt.u16.u32 %rs4, %r5;
|
||||
+ shr.u16 %rs5, %rs4, 8;
|
||||
+ xor.b16 %rs6, %rs5, 128;
|
||||
+ cvt.u32.u16 %r6, %rs6;
|
||||
+ st.param.b32 [func_retval0], %r6;
|
||||
+ ret;
|
||||
+ // -- End function
|
||||
+}
|
||||
+ // .globl convert_2x2_bgrx_area_to_yuv420_nv12 // -- Begin function convert_2x2_bgrx_area_to_yuv420_nv12
|
||||
.visible .entry convert_2x2_bgrx_area_to_yuv420_nv12(
|
||||
- .param .u64 convert_2x2_bgrx_area_to_yuv420_nv12_param_0,
|
||||
- .param .u64 convert_2x2_bgrx_area_to_yuv420_nv12_param_1,
|
||||
+ .param .u64 .ptr .align 1 convert_2x2_bgrx_area_to_yuv420_nv12_param_0,
|
||||
+ .param .u64 .ptr .align 1 convert_2x2_bgrx_area_to_yuv420_nv12_param_1,
|
||||
.param .u16 convert_2x2_bgrx_area_to_yuv420_nv12_param_2,
|
||||
.param .u16 convert_2x2_bgrx_area_to_yuv420_nv12_param_3,
|
||||
.param .u16 convert_2x2_bgrx_area_to_yuv420_nv12_param_4,
|
||||
.param .u16 convert_2x2_bgrx_area_to_yuv420_nv12_param_5,
|
||||
.param .u16 convert_2x2_bgrx_area_to_yuv420_nv12_param_6
|
||||
-)
|
||||
+) // @convert_2x2_bgrx_area_to_yuv420_nv12
|
||||
{
|
||||
- .reg .pred %p<17>;
|
||||
- .reg .b16 %rs<18>;
|
||||
- .reg .b32 %r<127>;
|
||||
- .reg .b64 %rd<40>;
|
||||
+ .reg .pred %p<13>;
|
||||
+ .reg .b16 %rs<27>;
|
||||
+ .reg .b32 %r<104>;
|
||||
+ .reg .b64 %rd<25>;
|
||||
|
||||
-
|
||||
- ld.param.u64 %rd6, [convert_2x2_bgrx_area_to_yuv420_nv12_param_0];
|
||||
- ld.param.u64 %rd7, [convert_2x2_bgrx_area_to_yuv420_nv12_param_1];
|
||||
- ld.param.u16 %rs7, [convert_2x2_bgrx_area_to_yuv420_nv12_param_2];
|
||||
- ld.param.u16 %rs8, [convert_2x2_bgrx_area_to_yuv420_nv12_param_3];
|
||||
- ld.param.u16 %rs9, [convert_2x2_bgrx_area_to_yuv420_nv12_param_5];
|
||||
- ld.param.u16 %rs10, [convert_2x2_bgrx_area_to_yuv420_nv12_param_6];
|
||||
- mov.u32 %r30, %ntid.x;
|
||||
- mov.u32 %r31, %ctaid.x;
|
||||
- mov.u32 %r32, %tid.x;
|
||||
- mad.lo.s32 %r1, %r30, %r31, %r32;
|
||||
- mov.u32 %r33, %ntid.y;
|
||||
- mov.u32 %r34, %ctaid.y;
|
||||
- mov.u32 %r35, %tid.y;
|
||||
- mad.lo.s32 %r2, %r33, %r34, %r35;
|
||||
- and.b32 %r36, %r1, 65535;
|
||||
- ld.param.u16 %r37, [convert_2x2_bgrx_area_to_yuv420_nv12_param_4];
|
||||
- shr.u32 %r38, %r37, 1;
|
||||
+// %bb.0:
|
||||
+ ld.param.u64 %rd7, [convert_2x2_bgrx_area_to_yuv420_nv12_param_0];
|
||||
+ ld.param.u64 %rd8, [convert_2x2_bgrx_area_to_yuv420_nv12_param_1];
|
||||
+ cvta.to.global.u64 %rd1, %rd8;
|
||||
+ cvta.to.global.u64 %rd2, %rd7;
|
||||
+ mov.u32 %r32, %ctaid.x;
|
||||
+ ld.param.u16 %rs12, [convert_2x2_bgrx_area_to_yuv420_nv12_param_4];
|
||||
+ mov.u32 %r33, %ntid.x;
|
||||
+ mov.u32 %r34, %tid.x;
|
||||
+ mad.lo.s32 %r1, %r32, %r33, %r34;
|
||||
+ and.b32 %r35, %r1, 65535;
|
||||
+ shr.u16 %rs13, %rs12, 1;
|
||||
+ cvt.u32.u16 %r36, %rs13;
|
||||
+ setp.ge.u32 %p1, %r35, %r36;
|
||||
+ @%p1 bra $L__BB3_11;
|
||||
+// %bb.1:
|
||||
+ ld.param.u16 %rs10, [convert_2x2_bgrx_area_to_yuv420_nv12_param_5];
|
||||
+ mov.u32 %r37, %tid.y;
|
||||
+ mov.u32 %r38, %ntid.y;
|
||||
+ mov.u32 %r39, %ctaid.y;
|
||||
+ mad.lo.s32 %r2, %r39, %r38, %r37;
|
||||
and.b32 %r3, %r2, 65535;
|
||||
- cvt.u32.u16 %r4, %rs9;
|
||||
- shr.u32 %r5, %r4, 1;
|
||||
- setp.ge.u32 %p1, %r3, %r5;
|
||||
- setp.ge.u32 %p2, %r36, %r38;
|
||||
- or.pred %p3, %p1, %p2;
|
||||
- @%p3 bra BB0_10;
|
||||
-
|
||||
- cvta.to.global.u64 %rd8, %rd6;
|
||||
+ cvt.u32.u16 %r4, %rs10;
|
||||
+ shr.u32 %r40, %r4, 1;
|
||||
+ setp.ge.u32 %p2, %r3, %r40;
|
||||
+ @%p2 bra $L__BB3_11;
|
||||
+// %bb.2:
|
||||
+ ld.param.u16 %rs11, [convert_2x2_bgrx_area_to_yuv420_nv12_param_6];
|
||||
+ ld.param.u16 %rs9, [convert_2x2_bgrx_area_to_yuv420_nv12_param_3];
|
||||
+ ld.param.u16 %rs8, [convert_2x2_bgrx_area_to_yuv420_nv12_param_2];
|
||||
+ cvt.u32.u16 %r5, %rs8;
|
||||
shl.b32 %r42, %r1, 1;
|
||||
- and.b32 %r6, %r42, 65534;
|
||||
- mov.u32 %r43, 1;
|
||||
- shl.b32 %r44, %r2, 1;
|
||||
- and.b32 %r7, %r44, 65534;
|
||||
- cvt.u32.u16 %r45, %rs7;
|
||||
- mul.lo.s32 %r46, %r7, %r45;
|
||||
- cvt.u64.u32 %rd9, %r46;
|
||||
- cvt.u64.u32 %rd10, %r6;
|
||||
- add.s64 %rd1, %rd9, %rd10;
|
||||
- cvta.to.global.u64 %rd11, %rd7;
|
||||
- shl.b64 %rd12, %rd1, 2;
|
||||
- add.s64 %rd2, %rd11, %rd12;
|
||||
- setp.lt.u32 %p4, %r7, %r5;
|
||||
- shl.b32 %r47, %r2, 2;
|
||||
- sub.s32 %r48, %r43, %r4;
|
||||
- selp.b32 %r49, 0, %r48, %p4;
|
||||
- mov.u32 %r124, 0;
|
||||
- add.s32 %r50, %r49, %r47;
|
||||
- cvt.u64.u32 %rd13, %r50;
|
||||
- and.b64 %rd14, %rd13, 65535;
|
||||
- cvt.u64.u16 %rd15, %rs10;
|
||||
- mul.lo.s64 %rd16, %rd14, %rd15;
|
||||
- add.s64 %rd17, %rd16, %rd10;
|
||||
- add.s64 %rd3, %rd8, %rd17;
|
||||
- add.s32 %r51, %r7, 1;
|
||||
- and.b32 %r52, %r51, 65535;
|
||||
- setp.lt.u32 %p5, %r52, %r5;
|
||||
- shl.b32 %r53, %r51, 1;
|
||||
- selp.b32 %r54, 0, %r48, %p5;
|
||||
- add.s32 %r55, %r54, %r53;
|
||||
- cvt.u64.u32 %rd18, %r55;
|
||||
- and.b64 %rd19, %rd18, 65535;
|
||||
- mul.lo.s64 %rd20, %rd19, %rd15;
|
||||
- add.s64 %rd21, %rd20, %rd10;
|
||||
- add.s64 %rd4, %rd8, %rd21;
|
||||
- shr.u32 %r56, %r4, 2;
|
||||
- setp.lt.u32 %p6, %r3, %r56;
|
||||
- sub.s32 %r57, %r43, %r5;
|
||||
- selp.b32 %r58, 0, %r57, %p6;
|
||||
- shl.b32 %r59, %r3, 1;
|
||||
- add.s32 %r60, %r59, %r58;
|
||||
- cvt.u64.u32 %rd22, %r60;
|
||||
- and.b64 %rd23, %rd22, 65535;
|
||||
- cvt.u64.u16 %rd24, %rs9;
|
||||
- add.s64 %rd25, %rd23, %rd24;
|
||||
- mul.lo.s64 %rd26, %rd25, %rd15;
|
||||
- add.s64 %rd27, %rd26, %rd10;
|
||||
- add.s64 %rd5, %rd8, %rd27;
|
||||
- setp.ge.u32 %p7, %r6, %r45;
|
||||
- cvt.u32.u16 %r61, %rs8;
|
||||
- setp.ge.u32 %p8, %r7, %r61;
|
||||
- mov.u16 %rs16, 0;
|
||||
- or.pred %p9, %p7, %p8;
|
||||
- mov.u16 %rs15, %rs16;
|
||||
- mov.u32 %r125, %r124;
|
||||
- mov.u32 %r126, %r124;
|
||||
- @%p9 bra BB0_3;
|
||||
-
|
||||
- ld.global.u32 %r62, [%rd2];
|
||||
- and.b32 %r126, %r62, 255;
|
||||
- bfe.u32 %r125, %r62, 8, 8;
|
||||
- bfe.u32 %r124, %r62, 16, 8;
|
||||
- mul.lo.s32 %r63, %r126, 18;
|
||||
- mad.lo.s32 %r64, %r124, 54, %r63;
|
||||
- mad.lo.s32 %r65, %r125, 183, %r64;
|
||||
- shr.u32 %r66, %r65, 8;
|
||||
- cvt.u16.u32 %rs15, %r66;
|
||||
-
|
||||
-BB0_3:
|
||||
- st.global.u8 [%rd3], %rs15;
|
||||
- add.s32 %r14, %r6, 1;
|
||||
- setp.ge.u32 %p11, %r14, %r45;
|
||||
- or.pred %p12, %p11, %p8;
|
||||
- @%p12 bra BB0_5;
|
||||
-
|
||||
- ld.global.u32 %r69, [%rd2+4];
|
||||
- and.b32 %r70, %r69, 255;
|
||||
- add.s32 %r126, %r70, %r126;
|
||||
- bfe.u32 %r71, %r69, 8, 8;
|
||||
- add.s32 %r125, %r71, %r125;
|
||||
- bfe.u32 %r72, %r69, 16, 8;
|
||||
- add.s32 %r124, %r72, %r124;
|
||||
- mul.lo.s32 %r73, %r70, 18;
|
||||
- mad.lo.s32 %r74, %r72, 54, %r73;
|
||||
- mad.lo.s32 %r75, %r71, 183, %r74;
|
||||
- shr.u32 %r76, %r75, 8;
|
||||
- cvt.u16.u32 %rs16, %r76;
|
||||
-
|
||||
-BB0_5:
|
||||
- setp.lt.u32 %p13, %r6, %r45;
|
||||
- st.global.u8 [%rd3+1], %rs16;
|
||||
- and.b32 %r82, %r2, 32767;
|
||||
- shl.b32 %r83, %r82, 1;
|
||||
- add.s32 %r84, %r83, 1;
|
||||
- setp.lt.u32 %p14, %r84, %r61;
|
||||
- and.pred %p15, %p13, %p14;
|
||||
- @%p15 bra BB0_7;
|
||||
- bra.uni BB0_6;
|
||||
-
|
||||
-BB0_7:
|
||||
- cvt.u64.u16 %rd28, %rs7;
|
||||
- add.s64 %rd29, %rd1, %rd28;
|
||||
- shl.b64 %rd31, %rd29, 2;
|
||||
- add.s64 %rd32, %rd11, %rd31;
|
||||
- ld.global.u32 %r87, [%rd32];
|
||||
- and.b32 %r88, %r87, 255;
|
||||
- add.s32 %r126, %r88, %r126;
|
||||
- bfe.u32 %r89, %r87, 8, 8;
|
||||
- add.s32 %r125, %r89, %r125;
|
||||
- bfe.u32 %r90, %r87, 16, 8;
|
||||
- add.s32 %r124, %r90, %r124;
|
||||
- mul.lo.s32 %r91, %r88, 18;
|
||||
- mad.lo.s32 %r92, %r90, 54, %r91;
|
||||
- mad.lo.s32 %r93, %r89, 183, %r92;
|
||||
- shr.u32 %r94, %r93, 8;
|
||||
- st.global.u8 [%rd4], %r94;
|
||||
- mov.u16 %rs17, 0;
|
||||
- @%p11 bra BB0_9;
|
||||
-
|
||||
- add.s64 %rd34, %rd28, 1;
|
||||
- and.b64 %rd35, %rd34, 65535;
|
||||
- add.s64 %rd36, %rd1, %rd35;
|
||||
- shl.b64 %rd38, %rd36, 2;
|
||||
- add.s64 %rd39, %rd11, %rd38;
|
||||
- ld.global.u32 %r95, [%rd39];
|
||||
- and.b32 %r96, %r95, 255;
|
||||
- add.s32 %r126, %r96, %r126;
|
||||
- bfe.u32 %r97, %r95, 8, 8;
|
||||
- add.s32 %r125, %r97, %r125;
|
||||
- bfe.u32 %r98, %r95, 16, 8;
|
||||
- add.s32 %r124, %r98, %r124;
|
||||
- mul.lo.s32 %r99, %r96, 18;
|
||||
- mad.lo.s32 %r100, %r98, 54, %r99;
|
||||
- mad.lo.s32 %r101, %r97, 183, %r100;
|
||||
- shr.u32 %r102, %r101, 8;
|
||||
- cvt.u16.u32 %rs17, %r102;
|
||||
- bra.uni BB0_9;
|
||||
-
|
||||
-BB0_6:
|
||||
- mov.u16 %rs17, 0;
|
||||
- st.global.u8 [%rd4], %rs17;
|
||||
-
|
||||
-BB0_9:
|
||||
- st.global.u8 [%rd4+1], %rs17;
|
||||
- bfe.u32 %r103, %r124, 2, 8;
|
||||
- mul.lo.s32 %r104, %r103, -29;
|
||||
- bfe.u32 %r105, %r125, 2, 8;
|
||||
- mad.lo.s32 %r106, %r105, -99, %r104;
|
||||
- bfe.u32 %r107, %r126, 2, 8;
|
||||
- shl.b32 %r108, %r107, 7;
|
||||
- add.s32 %r109, %r106, %r108;
|
||||
- shr.u32 %r110, %r109, 8;
|
||||
- add.s32 %r111, %r110, 128;
|
||||
- st.global.u8 [%rd5], %r111;
|
||||
- shl.b32 %r112, %r124, 5;
|
||||
- and.b32 %r113, %r112, 32640;
|
||||
- mad.lo.s32 %r114, %r105, -116, %r113;
|
||||
- mad.lo.s32 %r115, %r107, -12, %r114;
|
||||
- shr.u32 %r116, %r115, 8;
|
||||
- add.s32 %r117, %r116, 128;
|
||||
- st.global.u8 [%rd5+1], %r117;
|
||||
-
|
||||
-BB0_10:
|
||||
+ shl.b32 %r43, %r2, 1;
|
||||
+ and.b32 %r6, %r43, 65534;
|
||||
+ mul.lo.s32 %r44, %r6, %r5;
|
||||
+ mul.wide.u32 %rd9, %r44, 4;
|
||||
+ add.s64 %rd10, %rd1, %rd9;
|
||||
+ and.b32 %r7, %r42, 65534;
|
||||
+ cvt.u64.u32 %rd11, %r7;
|
||||
+ mul.wide.u32 %rd12, %r7, 4;
|
||||
+ add.s64 %rd3, %rd10, %rd12;
|
||||
+ cvt.u32.u16 %r45, %rs11;
|
||||
+ mul.lo.s32 %r46, %r6, %r45;
|
||||
+ cvt.u64.u32 %rd13, %r46;
|
||||
+ add.s64 %rd14, %rd2, %rd13;
|
||||
+ add.s64 %rd4, %rd14, %rd11;
|
||||
+ add.s32 %r47, %r46, %r45;
|
||||
+ cvt.u64.u32 %rd15, %r47;
|
||||
+ mul.lo.s32 %r48, %r45, %r4;
|
||||
+ cvt.u64.u32 %rd17, %r48;
|
||||
+ add.s64 %rd18, %rd2, %rd17;
|
||||
+ mul.lo.s32 %r49, %r3, %r45;
|
||||
+ cvt.u64.u32 %rd19, %r49;
|
||||
+ setp.ge.u32 %p3, %r7, %r5;
|
||||
+ cvt.u32.u16 %r9, %rs9;
|
||||
+ setp.ge.u32 %p4, %r6, %r9;
|
||||
+ mov.b32 %r98, 0;
|
||||
+ mov.b16 %rs25, 0;
|
||||
+ or.pred %p5, %p3, %p4;
|
||||
+ mov.u16 %rs24, %rs25;
|
||||
+ mov.u32 %r99, %r98;
|
||||
+ mov.u32 %r100, %r98;
|
||||
+ @%p5 bra $L__BB3_4;
|
||||
+// %bb.3:
|
||||
+ ld.global.u32 %r50, [%rd3];
|
||||
+ and.b32 %r100, %r50, 255;
|
||||
+ bfe.u32 %r99, %r50, 8, 8;
|
||||
+ bfe.u32 %r98, %r50, 16, 8;
|
||||
+ mul.lo.s32 %r51, %r100, 18;
|
||||
+ mad.lo.s32 %r52, %r98, 54, %r51;
|
||||
+ mad.lo.s32 %r53, %r99, 183, %r52;
|
||||
+ shr.u32 %r54, %r53, 8;
|
||||
+ cvt.u16.u32 %rs24, %r54;
|
||||
+$L__BB3_4:
|
||||
+ add.s64 %rd16, %rd2, %rd15;
|
||||
+ add.s64 %rd20, %rd18, %rd19;
|
||||
+ or.b32 %r8, %r6, 1;
|
||||
+ st.global.u8 [%rd4], %rs24;
|
||||
+ or.b32 %r16, %r7, 1;
|
||||
+ setp.ge.u32 %p7, %r16, %r5;
|
||||
+ or.pred %p8, %p7, %p4;
|
||||
+ @%p8 bra $L__BB3_6;
|
||||
+// %bb.5:
|
||||
+ ld.global.u32 %r57, [%rd3+4];
|
||||
+ and.b32 %r58, %r57, 255;
|
||||
+ add.s32 %r100, %r58, %r100;
|
||||
+ bfe.u32 %r59, %r57, 8, 8;
|
||||
+ add.s32 %r99, %r59, %r99;
|
||||
+ bfe.u32 %r60, %r57, 16, 8;
|
||||
+ add.s32 %r98, %r60, %r98;
|
||||
+ mul.lo.s32 %r61, %r58, 18;
|
||||
+ mad.lo.s32 %r62, %r60, 54, %r61;
|
||||
+ mad.lo.s32 %r63, %r59, 183, %r62;
|
||||
+ shr.u32 %r64, %r63, 8;
|
||||
+ cvt.u16.u32 %rs25, %r64;
|
||||
+$L__BB3_6:
|
||||
+ add.s64 %rd5, %rd16, %rd11;
|
||||
+ add.s64 %rd6, %rd20, %rd11;
|
||||
+ st.global.u8 [%rd4+1], %rs25;
|
||||
+ setp.ge.u32 %p10, %r8, %r9;
|
||||
+ or.pred %p11, %p3, %p10;
|
||||
+ @%p11 bra $L__BB3_9;
|
||||
+// %bb.7:
|
||||
+ mul.wide.u32 %rd21, %r5, 4;
|
||||
+ add.s64 %rd22, %rd3, %rd21;
|
||||
+ ld.global.u32 %r68, [%rd22];
|
||||
+ and.b32 %r69, %r68, 255;
|
||||
+ add.s32 %r100, %r69, %r100;
|
||||
+ bfe.u32 %r70, %r68, 8, 8;
|
||||
+ add.s32 %r99, %r70, %r99;
|
||||
+ bfe.u32 %r71, %r68, 16, 8;
|
||||
+ add.s32 %r98, %r71, %r98;
|
||||
+ mul.lo.s32 %r72, %r69, 18;
|
||||
+ mad.lo.s32 %r73, %r71, 54, %r72;
|
||||
+ mad.lo.s32 %r74, %r70, 183, %r73;
|
||||
+ shr.u32 %r75, %r74, 8;
|
||||
+ st.global.u8 [%rd5], %r75;
|
||||
+ mov.b16 %rs26, 0;
|
||||
+ @%p7 bra $L__BB3_10;
|
||||
+// %bb.8:
|
||||
+ add.s16 %rs1, %rs8, 1;
|
||||
+ cvt.u32.u16 %r76, %rs1;
|
||||
+ mul.wide.u32 %rd23, %r76, 4;
|
||||
+ add.s64 %rd24, %rd3, %rd23;
|
||||
+ ld.global.u32 %r77, [%rd24];
|
||||
+ and.b32 %r78, %r77, 255;
|
||||
+ add.s32 %r100, %r78, %r100;
|
||||
+ bfe.u32 %r79, %r77, 8, 8;
|
||||
+ add.s32 %r99, %r79, %r99;
|
||||
+ bfe.u32 %r80, %r77, 16, 8;
|
||||
+ add.s32 %r98, %r80, %r98;
|
||||
+ mul.lo.s32 %r81, %r78, 18;
|
||||
+ mad.lo.s32 %r82, %r80, 54, %r81;
|
||||
+ mad.lo.s32 %r83, %r79, 183, %r82;
|
||||
+ shr.u32 %r84, %r83, 8;
|
||||
+ cvt.u16.u32 %rs26, %r84;
|
||||
+ bra.uni $L__BB3_10;
|
||||
+$L__BB3_9:
|
||||
+ mov.b16 %rs26, 0;
|
||||
+ st.global.u8 [%rd5], %rs26;
|
||||
+$L__BB3_10:
|
||||
+ st.global.u8 [%rd5+1], %rs26;
|
||||
+ shr.u32 %r85, %r100, 2;
|
||||
+ shr.u32 %r86, %r99, 2;
|
||||
+ shr.u32 %r87, %r98, 2;
|
||||
+ mul.lo.s32 %r88, %r87, 65507;
|
||||
+ mad.lo.s32 %r89, %r86, 65437, %r88;
|
||||
+ shl.b32 %r90, %r85, 7;
|
||||
+ add.s32 %r91, %r89, %r90;
|
||||
+ cvt.u16.u32 %rs18, %r91;
|
||||
+ shr.u16 %rs19, %rs18, 8;
|
||||
+ xor.b16 %rs20, %rs19, 128;
|
||||
+ st.global.u8 [%rd6], %rs20;
|
||||
+ shl.b32 %r92, %r87, 7;
|
||||
+ mad.lo.s32 %r93, %r86, 65420, %r92;
|
||||
+ mad.lo.s32 %r94, %r85, 65524, %r93;
|
||||
+ cvt.u16.u32 %rs21, %r94;
|
||||
+ shr.u16 %rs22, %rs21, 8;
|
||||
+ xor.b16 %rs23, %rs22, 128;
|
||||
+ st.global.u8 [%rd6+1], %rs23;
|
||||
+$L__BB3_11:
|
||||
ret;
|
||||
+ // -- End function
|
||||
}
|
||||
-
|
||||
-
|
||||
diff --git a/src/grd-cuda-avc-utils.cu b/src/grd-cuda-avc-utils.cu
|
||||
index 7ce7f295..2f1e197d 100644
|
||||
--- a/src/grd-cuda-avc-utils.cu
|
||||
+++ b/src/grd-cuda-avc-utils.cu
|
||||
@@ -19,6 +19,10 @@
|
||||
|
||||
/*
|
||||
* Generate the PTX instructions with:
|
||||
+ * clang --cuda-gpu-arch=sm_30 -S src/grd-cuda-avc-utils.cu -o data/grd-cuda-avc-utils_30.ptx --no-cuda-version-check -O3 --cuda-device-only -Wall -Wextra
|
||||
+ *
|
||||
+ * or
|
||||
+ *
|
||||
* nvcc -arch=compute_30 -ptx grd-cuda-avc-utils.cu -o grd-cuda-avc-utils_30.ptx
|
||||
*
|
||||
* Note: This requires CUDA < 11, since the generation of Kepler capable
|
||||
@@ -29,24 +33,6 @@
|
||||
|
||||
extern "C"
|
||||
{
|
||||
- __device__ uint16_t
|
||||
- nv12_get_interlaced_y_1x1 (uint16_t y_1x1,
|
||||
- uint16_t aligned_height)
|
||||
- {
|
||||
- if (y_1x1 < aligned_height >> 1)
|
||||
- return y_1x1 << 1;
|
||||
- return (y_1x1 << 1) - aligned_height + 1;
|
||||
- }
|
||||
-
|
||||
- __device__ uint16_t
|
||||
- nv12_get_interlaced_y_2x2 (uint16_t y_2x2,
|
||||
- uint16_t aligned_height)
|
||||
- {
|
||||
- if (y_2x2 < aligned_height >> 2)
|
||||
- return y_2x2 << 1;
|
||||
- return (y_2x2 << 1) - (aligned_height >> 1) + 1;
|
||||
- }
|
||||
-
|
||||
__device__ uint8_t
|
||||
rgb_to_y (uint8_t r,
|
||||
uint8_t g,
|
||||
@@ -106,44 +92,17 @@ extern "C"
|
||||
s1 = 1;
|
||||
s2 = src_width;
|
||||
s3 = src_width + 1;
|
||||
- /*
|
||||
- * Technically, the correct positions for the Y data in the resulting NV12
|
||||
- * image would be the following:
|
||||
- *
|
||||
- * d0 = 0;
|
||||
- * d1 = 1;
|
||||
- * d2 = aligned_stride;
|
||||
- * d3 = aligned_stride + 1;
|
||||
- *
|
||||
- * However, since MBAFF is used as frame field mode, NVENC requires the input
|
||||
- * frame to be interlaced.
|
||||
- * If the frame is not interlaced, then even lines end up in the position
|
||||
- * y / 2, instead of y and odd lines end up in the position y / 2 +
|
||||
- * aligned_height / 2, instead of y.
|
||||
- * So, calculate the interlaced y position via a dedicated function, which
|
||||
- * ensures that the lines in the input frame end up in the resulting frame to
|
||||
- * be at the correct position.
|
||||
- * Doing this now in the kernel here, instead of after the BGRX -> YUV420
|
||||
- * conversion, saves a huge amount of time, since each thread only has a
|
||||
- * super tiny overhead to perform this action, while a normal
|
||||
- * device-to-device copy operation can take at least several milliseconds.
|
||||
- */
|
||||
|
||||
x_1x1 = x_2x2 << 1;
|
||||
y_1x1 = y_2x2 << 1;
|
||||
src_u32 = src_data + y_1x1 * src_width + x_1x1;
|
||||
|
||||
- dst_y0 = dst_data +
|
||||
- nv12_get_interlaced_y_1x1 (y_1x1, aligned_height) * aligned_stride +
|
||||
- x_1x1;
|
||||
+ dst_y0 = dst_data + y_1x1 * aligned_stride + x_1x1;
|
||||
dst_y1 = dst_y0 + 1;
|
||||
- dst_y2 = dst_data +
|
||||
- nv12_get_interlaced_y_1x1 (y_1x1 + 1, aligned_height) * aligned_stride +
|
||||
- x_1x1;
|
||||
+ dst_y2 = dst_data + (y_1x1 + 1) * aligned_stride + x_1x1;
|
||||
dst_y3 = dst_y2 + 1;
|
||||
dst_u = dst_data + aligned_height * aligned_stride +
|
||||
- nv12_get_interlaced_y_2x2 (y_2x2, aligned_height) * aligned_stride +
|
||||
- x_1x1;
|
||||
+ y_2x2 * aligned_stride + x_1x1;
|
||||
dst_v = dst_u + 1;
|
||||
|
||||
/* d_0 */
|
||||
diff --git a/src/grd-hwaccel-nvidia.c b/src/grd-hwaccel-nvidia.c
|
||||
index 76387635..b4640e02 100644
|
||||
--- a/src/grd-hwaccel-nvidia.c
|
||||
+++ b/src/grd-hwaccel-nvidia.c
|
||||
@@ -406,7 +406,7 @@ grd_hwaccel_nvidia_create_nvenc_session (GrdHwAccelNvidia *hwaccel_nvidia,
|
||||
encode_config.profileGUID = NV_ENC_H264_PROFILE_PROGRESSIVE_HIGH_GUID;
|
||||
encode_config.gopLength = NVENC_INFINITE_GOPLENGTH;
|
||||
encode_config.frameIntervalP = 1;
|
||||
- encode_config.frameFieldMode = NV_ENC_PARAMS_FRAME_FIELD_MODE_MBAFF;
|
||||
+ encode_config.frameFieldMode = NV_ENC_PARAMS_FRAME_FIELD_MODE_FRAME;
|
||||
encode_config.mvPrecision = NV_ENC_MV_PRECISION_QUARTER_PEL;
|
||||
encode_config.rcParams.version = NV_ENC_RC_PARAMS_VER;
|
||||
encode_config.rcParams.rateControlMode = NV_ENC_PARAMS_RC_VBR;
|
||||
--
|
||||
2.51.0
|
||||
|
||||
@ -13,7 +13,7 @@
|
||||
|
||||
Name: gnome-remote-desktop
|
||||
Version: 47.3
|
||||
Release: 2%{?dist}
|
||||
Release: 3%{?dist}
|
||||
Summary: GNOME Remote Desktop screen share service
|
||||
|
||||
License: GPL-2.0-or-later
|
||||
@ -24,6 +24,9 @@ Source0: https://download.gnome.org/sources/%{name}/47/%{name}-%{tarball_
|
||||
Patch0: gnutls-anontls.patch
|
||||
Patch1: connection-throttling.patch
|
||||
|
||||
# https://issues.redhat.com/browse/RHEL-146620
|
||||
Patch2: 0001-hwaccel-nvidia-Fix-image-corruptions-on-Blackwell-GP.patch
|
||||
|
||||
BuildRequires: asciidoc
|
||||
BuildRequires: gcc
|
||||
BuildRequires: meson >= 0.47.0
|
||||
@ -148,6 +151,10 @@ GNOME desktop environment.
|
||||
|
||||
|
||||
%changelog
|
||||
* Tue Feb 10 2026 Jonas Ådahl <jadahl@redhat.com> - 47.3-3
|
||||
- Fix image corruptions on NVIDIA Blackwell GPUs
|
||||
Resolves: RHEL-146620
|
||||
|
||||
* Thu Jul 03 2025 Jonas Ådahl <jadahl@redhat.com> - 47.3-2
|
||||
- Backport connection throttling
|
||||
Resolves: RHEL-92787
|
||||
|
||||
Loading…
Reference in New Issue
Block a user