From ffc78ae67d93c806ba6a29ae7a20320e718b89d6 Mon Sep 17 00:00:00 2001 From: AlmaLinux RelEng Bot Date: Tue, 17 Mar 2026 12:27:59 -0400 Subject: [PATCH] import Oracle_OSS gnome-remote-desktop-47.3-3.el10_1 --- ...ix-image-corruptions-on-Blackwell-GP.patch | 606 ++++++++++++++++++ gnome-remote-desktop.spec | 9 +- 2 files changed, 614 insertions(+), 1 deletion(-) create mode 100644 0001-hwaccel-nvidia-Fix-image-corruptions-on-Blackwell-GP.patch diff --git a/0001-hwaccel-nvidia-Fix-image-corruptions-on-Blackwell-GP.patch b/0001-hwaccel-nvidia-Fix-image-corruptions-on-Blackwell-GP.patch new file mode 100644 index 0000000..c11cbb4 --- /dev/null +++ b/0001-hwaccel-nvidia-Fix-image-corruptions-on-Blackwell-GP.patch @@ -0,0 +1,606 @@ +From eb4a7aebfd005d3a82b6bd63ee864bef941d9285 Mon Sep 17 00:00:00 2001 +From: Stephan Seitz +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: +--- + 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 + diff --git a/gnome-remote-desktop.spec b/gnome-remote-desktop.spec index dbca173..644eb98 100644 --- a/gnome-remote-desktop.spec +++ b/gnome-remote-desktop.spec @@ -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 - 47.3-3 +- Fix image corruptions on NVIDIA Blackwell GPUs + Resolves: RHEL-146620 + * Thu Jul 03 2025 Jonas Ådahl - 47.3-2 - Backport connection throttling Resolves: RHEL-92787