// // Generated by NVIDIA NVVM Compiler // // Compiler Build ID: CL-38244171 // Cuda compilation tools, release 13.3, V13.3.73 // Based on NVVM 7.0.1 // .version 9.3 .target sm_75 .address_size 64 // .globl blend_argb .visible .entry blend_argb( .param .u64 blend_argb_param_0, .param .u32 blend_argb_param_1, .param .u32 blend_argb_param_2, .param .u32 blend_argb_param_3, .param .u64 blend_argb_param_4, .param .u32 blend_argb_param_5, .param .u32 blend_argb_param_6, .param .u32 blend_argb_param_7, .param .u32 blend_argb_param_8 ) { .reg .pred %p<10>; .reg .b16 %rs<2>; .reg .b32 %r<34>; .reg .b64 %rd<17>; ld.param.u64 %rd2, [blend_argb_param_0]; ld.param.u32 %r5, [blend_argb_param_1]; ld.param.u32 %r6, [blend_argb_param_2]; ld.param.u32 %r7, [blend_argb_param_3]; ld.param.u64 %rd3, [blend_argb_param_4]; ld.param.u32 %r8, [blend_argb_param_5]; ld.param.u32 %r11, [blend_argb_param_6]; ld.param.u32 %r9, [blend_argb_param_7]; ld.param.u32 %r10, [blend_argb_param_8]; mov.u32 %r12, %ntid.x; mov.u32 %r13, %ctaid.x; mov.u32 %r14, %tid.x; mad.lo.s32 %r1, %r13, %r12, %r14; mov.u32 %r15, %ntid.y; mov.u32 %r16, %ctaid.y; mov.u32 %r17, %tid.y; mad.lo.s32 %r2, %r16, %r15, %r17; setp.ge.s32 %p1, %r1, %r8; setp.ge.s32 %p2, %r2, %r11; or.pred %p3, %p1, %p2; @%p3 bra $L__BB0_4; add.s32 %r3, %r1, %r9; add.s32 %r4, %r2, %r10; or.b32 %r18, %r4, %r3; setp.lt.s32 %p4, %r18, 0; setp.ge.s32 %p5, %r3, %r6; or.pred %p6, %p5, %p4; setp.ge.s32 %p7, %r4, %r7; or.pred %p8, %p7, %p6; @%p8 bra $L__BB0_4; mad.lo.s32 %r19, %r2, %r8, %r1; mul.wide.s32 %rd4, %r19, 4; cvta.to.global.u64 %rd5, %rd3; add.s64 %rd1, %rd5, %rd4; ld.global.u8 %rs1, [%rd1+3]; setp.eq.s16 %p9, %rs1, 0; @%p9 bra $L__BB0_4; cvt.u32.u16 %r20, %rs1; mul.wide.s32 %rd6, %r4, %r5; mul.wide.s32 %rd7, %r3, 4; add.s64 %rd8, %rd6, %rd7; cvta.to.global.u64 %rd9, %rd2; add.s64 %rd10, %rd9, %rd8; ld.global.u8 %r21, [%rd10]; ld.global.u8 %r22, [%rd1+2]; xor.b32 %r23, %r20, 255; mul.lo.s32 %r24, %r23, %r21; mad.lo.s32 %r25, %r22, %r20, %r24; mul.wide.u32 %rd11, %r25, -2139062143; shr.u64 %rd12, %rd11, 39; st.global.u8 [%rd10], %rd12; ld.global.u8 %r26, [%rd10+1]; ld.global.u8 %r27, [%rd1+1]; mul.lo.s32 %r28, %r23, %r26; mad.lo.s32 %r29, %r27, %r20, %r28; mul.wide.u32 %rd13, %r29, -2139062143; shr.u64 %rd14, %rd13, 39; st.global.u8 [%rd10+1], %rd14; ld.global.u8 %r30, [%rd10+2]; ld.global.u8 %r31, [%rd1]; mul.lo.s32 %r32, %r23, %r30; mad.lo.s32 %r33, %r31, %r20, %r32; mul.wide.u32 %rd15, %r33, -2139062143; shr.u64 %rd16, %rd15, 39; st.global.u8 [%rd10+2], %rd16; $L__BB0_4: ret; } // .globl blend_yuv444 .visible .entry blend_yuv444( .param .u64 blend_yuv444_param_0, .param .u32 blend_yuv444_param_1, .param .u32 blend_yuv444_param_2, .param .u32 blend_yuv444_param_3, .param .u64 blend_yuv444_param_4, .param .u32 blend_yuv444_param_5, .param .u32 blend_yuv444_param_6, .param .u32 blend_yuv444_param_7, .param .u32 blend_yuv444_param_8 ) { .reg .pred %p<10>; .reg .b16 %rs<5>; .reg .f32 %f<16>; .reg .b32 %r<49>; .reg .b64 %rd<14>; ld.param.u64 %rd2, [blend_yuv444_param_0]; ld.param.u32 %r5, [blend_yuv444_param_1]; ld.param.u32 %r6, [blend_yuv444_param_2]; ld.param.u32 %r7, [blend_yuv444_param_3]; ld.param.u64 %rd3, [blend_yuv444_param_4]; ld.param.u32 %r8, [blend_yuv444_param_5]; ld.param.u32 %r11, [blend_yuv444_param_6]; ld.param.u32 %r9, [blend_yuv444_param_7]; ld.param.u32 %r10, [blend_yuv444_param_8]; mov.u32 %r12, %ntid.x; mov.u32 %r13, %ctaid.x; mov.u32 %r14, %tid.x; mad.lo.s32 %r1, %r13, %r12, %r14; mov.u32 %r15, %ntid.y; mov.u32 %r16, %ctaid.y; mov.u32 %r17, %tid.y; mad.lo.s32 %r2, %r16, %r15, %r17; setp.ge.s32 %p1, %r1, %r8; setp.ge.s32 %p2, %r2, %r11; or.pred %p3, %p1, %p2; @%p3 bra $L__BB1_4; add.s32 %r3, %r1, %r9; add.s32 %r4, %r2, %r10; or.b32 %r18, %r4, %r3; setp.lt.s32 %p4, %r18, 0; setp.ge.s32 %p5, %r3, %r6; or.pred %p6, %p5, %p4; setp.ge.s32 %p7, %r4, %r7; or.pred %p8, %p7, %p6; @%p8 bra $L__BB1_4; mad.lo.s32 %r19, %r2, %r8, %r1; mul.wide.s32 %rd4, %r19, 4; cvta.to.global.u64 %rd5, %rd3; add.s64 %rd1, %rd5, %rd4; ld.global.u8 %rs1, [%rd1+3]; setp.eq.s16 %p9, %rs1, 0; @%p9 bra $L__BB1_4; cvt.u32.u16 %r20, %rs1; ld.global.u8 %rs2, [%rd1]; cvt.rn.f32.u16 %f1, %rs2; ld.global.u8 %rs3, [%rd1+1]; cvt.rn.f32.u16 %f2, %rs3; ld.global.u8 %rs4, [%rd1+2]; cvt.rn.f32.u16 %f3, %rs4; fma.rn.f32 %f4, %f1, 0f3E3AFB7F, 0f41800000; fma.rn.f32 %f5, %f2, 0f3F1D3C36, %f4; fma.rn.f32 %f6, %f3, 0f3D7DF3B6, %f5; add.f32 %f7, %f6, 0f3F000000; cvt.rzi.s32.f32 %r21, %f7; fma.rn.f32 %f8, %f1, 0fBDCE075F, 0f43000000; fma.rn.f32 %f9, %f2, 0fBEAD5CFB, %f8; fma.rn.f32 %f10, %f3, 0f3EE0DED3, %f9; add.f32 %f11, %f10, 0f3F000000; cvt.rzi.s32.f32 %r22, %f11; fma.rn.f32 %f12, %f1, 0f3EE0DED3, 0f43000000; fma.rn.f32 %f13, %f2, 0fBECC3C9F, %f12; fma.rn.f32 %f14, %f3, 0fBD25119D, %f13; add.f32 %f15, %f14, 0f3F000000; cvt.rzi.s32.f32 %r23, %f15; mul.wide.s32 %rd6, %r4, %r5; cvt.s64.s32 %rd7, %r3; add.s64 %rd8, %rd6, %rd7; cvta.to.global.u64 %rd9, %rd2; add.s64 %rd10, %rd9, %rd8; ld.global.u8 %r24, [%rd10]; mul.lo.s32 %r25, %r21, %r20; xor.b32 %r26, %r20, 255; mad.lo.s32 %r27, %r26, %r24, %r25; mul.hi.s32 %r28, %r27, -2139062143; add.s32 %r29, %r28, %r27; shr.u32 %r30, %r29, 31; shr.u32 %r31, %r29, 7; add.s32 %r32, %r31, %r30; st.global.u8 [%rd10], %r32; mul.wide.s32 %rd11, %r7, %r5; add.s64 %rd12, %rd10, %rd11; ld.global.u8 %r33, [%rd12]; mul.lo.s32 %r34, %r22, %r20; mad.lo.s32 %r35, %r26, %r33, %r34; mul.hi.s32 %r36, %r35, -2139062143; add.s32 %r37, %r36, %r35; shr.u32 %r38, %r37, 31; shr.u32 %r39, %r37, 7; add.s32 %r40, %r39, %r38; st.global.u8 [%rd12], %r40; add.s64 %rd13, %rd12, %rd11; ld.global.u8 %r41, [%rd13]; mul.lo.s32 %r42, %r23, %r20; mad.lo.s32 %r43, %r26, %r41, %r42; mul.hi.s32 %r44, %r43, -2139062143; add.s32 %r45, %r44, %r43; shr.u32 %r46, %r45, 31; shr.u32 %r47, %r45, 7; add.s32 %r48, %r47, %r46; st.global.u8 [%rd13], %r48; $L__BB1_4: ret; } // .globl blend_nv12 .visible .entry blend_nv12( .param .u64 blend_nv12_param_0, .param .u32 blend_nv12_param_1, .param .u64 blend_nv12_param_2, .param .u32 blend_nv12_param_3, .param .u32 blend_nv12_param_4, .param .u32 blend_nv12_param_5, .param .u64 blend_nv12_param_6, .param .u32 blend_nv12_param_7, .param .u32 blend_nv12_param_8, .param .u32 blend_nv12_param_9, .param .u32 blend_nv12_param_10 ) { .reg .pred %p<43>; .reg .b16 %rs<17>; .reg .f32 %f<108>; .reg .b32 %r<123>; .reg .b64 %rd<35>; ld.param.u64 %rd11, [blend_nv12_param_0]; ld.param.u32 %r21, [blend_nv12_param_1]; ld.param.u64 %rd10, [blend_nv12_param_2]; ld.param.u32 %r22, [blend_nv12_param_3]; ld.param.u32 %r23, [blend_nv12_param_4]; ld.param.u32 %r24, [blend_nv12_param_5]; ld.param.u64 %rd12, [blend_nv12_param_6]; ld.param.u32 %r25, [blend_nv12_param_7]; ld.param.u32 %r26, [blend_nv12_param_8]; ld.param.u32 %r27, [blend_nv12_param_9]; ld.param.u32 %r28, [blend_nv12_param_10]; cvta.to.global.u64 %rd1, %rd11; cvta.to.global.u64 %rd2, %rd12; mov.u32 %r29, %ntid.x; mov.u32 %r30, %ctaid.x; mov.u32 %r31, %tid.x; mad.lo.s32 %r32, %r30, %r29, %r31; mov.u32 %r33, %ntid.y; mov.u32 %r34, %ctaid.y; mov.u32 %r35, %tid.y; mad.lo.s32 %r36, %r34, %r33, %r35; shl.b32 %r1, %r32, 1; shl.b32 %r2, %r36, 1; setp.ge.s32 %p1, %r1, %r25; setp.ge.s32 %p2, %r2, %r26; or.pred %p3, %p1, %p2; mov.f32 %f102, 0f00000000; mov.f32 %f103, 0f00000000; mov.f32 %f104, 0f00000000; @%p3 bra $L__BB2_19; cvt.s64.s32 %rd3, %r21; add.s32 %r3, %r2, %r28; setp.ge.s32 %p4, %r3, %r24; mul.lo.s32 %r4, %r2, %r25; mul.wide.s32 %rd4, %r3, %r21; add.s32 %r5, %r1, %r27; or.b32 %r38, %r5, %r3; setp.lt.s32 %p5, %r38, 0; mov.u32 %r121, 0; setp.ge.s32 %p6, %r5, %r23; or.pred %p7, %p6, %p5; or.pred %p8, %p4, %p7; @%p8 bra $L__BB2_4; add.s32 %r40, %r1, %r4; mul.wide.s32 %rd13, %r40, 4; add.s64 %rd5, %rd2, %rd13; ld.global.u8 %rs1, [%rd5+3]; setp.eq.s16 %p9, %rs1, 0; @%p9 bra $L__BB2_4; cvt.u32.u16 %r42, %rs1; ld.global.u8 %rs5, [%rd5]; cvt.rn.f32.u16 %f31, %rs5; ld.global.u8 %rs6, [%rd5+1]; cvt.rn.f32.u16 %f32, %rs6; ld.global.u8 %rs7, [%rd5+2]; cvt.rn.f32.u16 %f33, %rs7; fma.rn.f32 %f34, %f31, 0f3E3AFB7F, 0f41800000; fma.rn.f32 %f35, %f32, 0f3F1D3C36, %f34; fma.rn.f32 %f36, %f33, 0f3D7DF3B6, %f35; add.f32 %f37, %f36, 0f3F000000; cvt.rzi.s32.f32 %r43, %f37; cvt.s64.s32 %rd14, %r5; add.s64 %rd15, %rd4, %rd14; add.s64 %rd16, %rd1, %rd15; ld.global.u8 %r44, [%rd16]; mul.lo.s32 %r45, %r43, %r42; xor.b32 %r46, %r42, 255; mad.lo.s32 %r47, %r46, %r44, %r45; mul.hi.s32 %r48, %r47, -2139062143; add.s32 %r49, %r48, %r47; shr.u32 %r50, %r49, 31; shr.u32 %r51, %r49, 7; add.s32 %r52, %r51, %r50; st.global.u8 [%rd16], %r52; fma.rn.f32 %f38, %f31, 0fBDCE075F, 0f43000000; fma.rn.f32 %f39, %f32, 0fBEAD5CFB, %f38; fma.rn.f32 %f40, %f33, 0f3EE0DED3, %f39; cvt.rn.f32.u16 %f104, %rs1; fma.rn.f32 %f102, %f40, %f104, 0f00000000; fma.rn.f32 %f41, %f31, 0f3EE0DED3, 0f43000000; fma.rn.f32 %f42, %f32, 0fBECC3C9F, %f41; fma.rn.f32 %f43, %f33, 0fBD25119D, %f42; fma.rn.f32 %f103, %f43, %f104, 0f00000000; mov.u32 %r121, 1; $L__BB2_4: add.s32 %r7, %r1, 1; setp.ge.s32 %p10, %r7, %r25; @%p10 bra $L__BB2_8; add.s32 %r8, %r7, %r27; or.b32 %r53, %r8, %r3; setp.lt.s32 %p12, %r53, 0; setp.ge.s32 %p13, %r8, %r23; or.pred %p14, %p13, %p12; or.pred %p15, %p4, %p14; @%p15 bra $L__BB2_8; add.s32 %r54, %r7, %r4; mul.wide.s32 %rd17, %r54, 4; add.s64 %rd6, %rd2, %rd17; ld.global.u8 %rs2, [%rd6+3]; setp.eq.s16 %p16, %rs2, 0; @%p16 bra $L__BB2_8; cvt.u32.u16 %r55, %rs2; ld.global.u8 %rs8, [%rd6]; cvt.rn.f32.u16 %f44, %rs8; ld.global.u8 %rs9, [%rd6+1]; cvt.rn.f32.u16 %f45, %rs9; ld.global.u8 %rs10, [%rd6+2]; cvt.rn.f32.u16 %f46, %rs10; fma.rn.f32 %f47, %f44, 0f3E3AFB7F, 0f41800000; fma.rn.f32 %f48, %f45, 0f3F1D3C36, %f47; fma.rn.f32 %f49, %f46, 0f3D7DF3B6, %f48; add.f32 %f50, %f49, 0f3F000000; cvt.rzi.s32.f32 %r56, %f50; cvt.s64.s32 %rd18, %r8; add.s64 %rd19, %rd4, %rd18; add.s64 %rd20, %rd1, %rd19; ld.global.u8 %r57, [%rd20]; mul.lo.s32 %r58, %r56, %r55; xor.b32 %r59, %r55, 255; mad.lo.s32 %r60, %r59, %r57, %r58; mul.hi.s32 %r61, %r60, -2139062143; add.s32 %r62, %r61, %r60; shr.u32 %r63, %r62, 31; shr.u32 %r64, %r62, 7; add.s32 %r65, %r64, %r63; st.global.u8 [%rd20], %r65; fma.rn.f32 %f51, %f44, 0fBDCE075F, 0f43000000; fma.rn.f32 %f52, %f45, 0fBEAD5CFB, %f51; fma.rn.f32 %f53, %f46, 0f3EE0DED3, %f52; cvt.rn.f32.u16 %f54, %rs2; fma.rn.f32 %f102, %f53, %f54, %f102; fma.rn.f32 %f55, %f44, 0f3EE0DED3, 0f43000000; fma.rn.f32 %f56, %f45, 0fBECC3C9F, %f55; fma.rn.f32 %f57, %f46, 0fBD25119D, %f56; fma.rn.f32 %f103, %f57, %f54, %f103; add.f32 %f104, %f104, %f54; add.s32 %r121, %r121, 1; $L__BB2_8: add.s32 %r11, %r2, 1; setp.ge.s32 %p17, %r11, %r26; add.s32 %r12, %r11, %r28; add.s32 %r13, %r4, %r25; cvt.s64.s32 %rd21, %r12; mul.lo.s64 %rd7, %rd21, %rd3; @%p17 bra $L__BB2_12; setp.ge.s32 %p18, %r12, %r24; or.b32 %r66, %r5, %r12; setp.lt.s32 %p19, %r66, 0; or.pred %p21, %p6, %p19; or.pred %p22, %p18, %p21; @%p22 bra $L__BB2_12; add.s32 %r67, %r1, %r13; mul.wide.s32 %rd22, %r67, 4; add.s64 %rd8, %rd2, %rd22; ld.global.u8 %rs3, [%rd8+3]; setp.eq.s16 %p23, %rs3, 0; @%p23 bra $L__BB2_12; cvt.u32.u16 %r68, %rs3; ld.global.u8 %rs11, [%rd8]; cvt.rn.f32.u16 %f58, %rs11; ld.global.u8 %rs12, [%rd8+1]; cvt.rn.f32.u16 %f59, %rs12; ld.global.u8 %rs13, [%rd8+2]; cvt.rn.f32.u16 %f60, %rs13; fma.rn.f32 %f61, %f58, 0f3E3AFB7F, 0f41800000; fma.rn.f32 %f62, %f59, 0f3F1D3C36, %f61; fma.rn.f32 %f63, %f60, 0f3D7DF3B6, %f62; add.f32 %f64, %f63, 0f3F000000; cvt.rzi.s32.f32 %r69, %f64; cvt.s64.s32 %rd23, %r5; add.s64 %rd24, %rd7, %rd23; add.s64 %rd25, %rd1, %rd24; ld.global.u8 %r70, [%rd25]; mul.lo.s32 %r71, %r69, %r68; xor.b32 %r72, %r68, 255; mad.lo.s32 %r73, %r72, %r70, %r71; mul.hi.s32 %r74, %r73, -2139062143; add.s32 %r75, %r74, %r73; shr.u32 %r76, %r75, 31; shr.u32 %r77, %r75, 7; add.s32 %r78, %r77, %r76; st.global.u8 [%rd25], %r78; fma.rn.f32 %f65, %f58, 0fBDCE075F, 0f43000000; fma.rn.f32 %f66, %f59, 0fBEAD5CFB, %f65; fma.rn.f32 %f67, %f60, 0f3EE0DED3, %f66; cvt.rn.f32.u16 %f68, %rs3; fma.rn.f32 %f102, %f67, %f68, %f102; fma.rn.f32 %f69, %f58, 0f3EE0DED3, 0f43000000; fma.rn.f32 %f70, %f59, 0fBECC3C9F, %f69; fma.rn.f32 %f71, %f60, 0fBD25119D, %f70; fma.rn.f32 %f103, %f71, %f68, %f103; add.f32 %f104, %f104, %f68; add.s32 %r121, %r121, 1; $L__BB2_12: or.pred %p26, %p17, %p10; @%p26 bra $L__BB2_16; setp.ge.s32 %p27, %r12, %r24; add.s32 %r16, %r7, %r27; or.b32 %r79, %r16, %r12; setp.lt.s32 %p28, %r79, 0; setp.ge.s32 %p29, %r16, %r23; or.pred %p30, %p29, %p28; or.pred %p31, %p27, %p30; @%p31 bra $L__BB2_16; add.s32 %r80, %r7, %r13; mul.wide.s32 %rd26, %r80, 4; add.s64 %rd9, %rd2, %rd26; ld.global.u8 %rs4, [%rd9+3]; setp.eq.s16 %p32, %rs4, 0; @%p32 bra $L__BB2_16; cvt.u32.u16 %r81, %rs4; ld.global.u8 %rs14, [%rd9]; cvt.rn.f32.u16 %f72, %rs14; ld.global.u8 %rs15, [%rd9+1]; cvt.rn.f32.u16 %f73, %rs15; ld.global.u8 %rs16, [%rd9+2]; cvt.rn.f32.u16 %f74, %rs16; fma.rn.f32 %f75, %f72, 0f3E3AFB7F, 0f41800000; fma.rn.f32 %f76, %f73, 0f3F1D3C36, %f75; fma.rn.f32 %f77, %f74, 0f3D7DF3B6, %f76; add.f32 %f78, %f77, 0f3F000000; cvt.rzi.s32.f32 %r82, %f78; cvt.s64.s32 %rd27, %r16; add.s64 %rd28, %rd7, %rd27; add.s64 %rd29, %rd1, %rd28; ld.global.u8 %r83, [%rd29]; mul.lo.s32 %r84, %r82, %r81; xor.b32 %r85, %r81, 255; mad.lo.s32 %r86, %r85, %r83, %r84; mul.hi.s32 %r87, %r86, -2139062143; add.s32 %r88, %r87, %r86; shr.u32 %r89, %r88, 31; shr.u32 %r90, %r88, 7; add.s32 %r91, %r90, %r89; st.global.u8 [%rd29], %r91; fma.rn.f32 %f79, %f72, 0fBDCE075F, 0f43000000; fma.rn.f32 %f80, %f73, 0fBEAD5CFB, %f79; fma.rn.f32 %f81, %f74, 0f3EE0DED3, %f80; cvt.rn.f32.u16 %f82, %rs4; fma.rn.f32 %f102, %f81, %f82, %f102; fma.rn.f32 %f83, %f72, 0f3EE0DED3, 0f43000000; fma.rn.f32 %f84, %f73, 0fBECC3C9F, %f83; fma.rn.f32 %f85, %f74, 0fBD25119D, %f84; fma.rn.f32 %f103, %f85, %f82, %f103; add.f32 %f104, %f104, %f82; add.s32 %r121, %r121, 1; $L__BB2_16: setp.eq.s32 %p33, %r121, 0; setp.le.f32 %p34, %f104, 0f00000000; or.pred %p35, %p34, %p33; @%p35 bra $L__BB2_19; shr.u32 %r92, %r5, 31; add.s32 %r93, %r5, %r92; shr.s32 %r19, %r93, 1; setp.lt.s32 %p36, %r3, -1; setp.lt.s32 %p37, %r5, -1; or.pred %p38, %p37, %p36; and.b32 %r94, %r93, -2; setp.ge.s32 %p39, %r94, %r23; or.pred %p40, %p38, %p39; shr.u32 %r95, %r3, 31; add.s32 %r96, %r3, %r95; shr.s32 %r20, %r96, 1; and.b32 %r97, %r96, -2; setp.ge.s32 %p41, %r97, %r24; or.pred %p42, %p40, %p41; @%p42 bra $L__BB2_19; div.rn.f32 %f86, %f102, %f104; add.f32 %f87, %f86, 0f3F000000; cvt.rzi.s32.f32 %r98, %f87; div.rn.f32 %f88, %f103, %f104; add.f32 %f89, %f88, 0f3F000000; cvt.rzi.s32.f32 %r99, %f89; cvt.rn.f32.s32 %f90, %r121; div.rn.f32 %f91, %f104, %f90; add.f32 %f92, %f91, 0f3F000000; cvt.rzi.s32.f32 %r100, %f92; mul.wide.s32 %rd30, %r20, %r22; mul.wide.s32 %rd31, %r19, 2; add.s64 %rd32, %rd30, %rd31; cvta.to.global.u64 %rd33, %rd10; add.s64 %rd34, %rd33, %rd32; ld.global.u8 %r101, [%rd34]; mul.lo.s32 %r102, %r100, %r98; mov.u32 %r103, 255; sub.s32 %r104, %r103, %r100; mad.lo.s32 %r105, %r104, %r101, %r102; mul.hi.s32 %r106, %r105, -2139062143; add.s32 %r107, %r106, %r105; shr.u32 %r108, %r107, 31; shr.u32 %r109, %r107, 7; add.s32 %r110, %r109, %r108; st.global.u8 [%rd34], %r110; ld.global.u8 %r111, [%rd34+1]; mul.lo.s32 %r112, %r100, %r99; mad.lo.s32 %r113, %r104, %r111, %r112; mul.hi.s32 %r114, %r113, -2139062143; add.s32 %r115, %r114, %r113; shr.u32 %r116, %r115, 31; shr.u32 %r117, %r115, 7; add.s32 %r118, %r117, %r116; st.global.u8 [%rd34+1], %r118; $L__BB2_19: ret; }