// // Generated by NVIDIA NVVM Compiler // // Compiler Build ID: CL-32688072 // Cuda compilation tools, release 12.1, V12.1.105 // Based on NVVM 7.0.1 // .version 8.1 .target sm_80 .address_size 64 // _ZZ22sgemm_128x128x8_kernelPKfS0_PfjjjjjE4smem has been demoted .func _Z9C_tile_wb7StgFragPfPKfjjjjj( .param .align 4 .b8 _Z9C_tile_wb7StgFragPfPKfjjjjj_param_0[64], .param .b64 _Z9C_tile_wb7StgFragPfPKfjjjjj_param_1, .param .b64 _Z9C_tile_wb7StgFragPfPKfjjjjj_param_2, .param .b32 _Z9C_tile_wb7StgFragPfPKfjjjjj_param_3, .param .b32 _Z9C_tile_wb7StgFragPfPKfjjjjj_param_4, .param .b32 _Z9C_tile_wb7StgFragPfPKfjjjjj_param_5, .param .b32 _Z9C_tile_wb7StgFragPfPKfjjjjj_param_6, .param .b32 _Z9C_tile_wb7StgFragPfPKfjjjjj_param_7 ) { .reg .pred %p<34>; .reg .f32 %f<33>; .reg .b32 %r<42>; .reg .b64 %rd<33>; ld.param.u64 %rd1, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_1]; ld.param.u64 %rd17, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_2]; ld.param.u32 %r1, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_3]; ld.param.u32 %r21, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_4]; ld.param.u32 %r22, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_5]; ld.param.u32 %r23, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_6]; ld.param.u32 %r24, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_7]; { .reg .b64 %tmp; cvta.to.shared.u64 %tmp, %rd17; cvt.u32.u64 %r25, %tmp; } ld.param.f32 %f13, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+48]; ld.param.f32 %f14, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+52]; ld.param.f32 %f15, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+56]; ld.param.f32 %f16, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+60]; ld.param.f32 %f9, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+32]; ld.param.f32 %f10, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+36]; ld.param.f32 %f11, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+40]; ld.param.f32 %f12, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+44]; ld.param.f32 %f5, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+16]; ld.param.f32 %f6, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+20]; ld.param.f32 %f7, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+24]; ld.param.f32 %f8, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+28]; ld.param.f32 %f1, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0]; ld.param.f32 %f2, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+4]; ld.param.f32 %f3, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+8]; ld.param.f32 %f4, [_Z9C_tile_wb7StgFragPfPKfjjjjj_param_0+12]; bar.sync 0; // begin inline asm st.shared.v4.f32 [%r1], {%f1, %f2, %f3, %f4}; // end inline asm add.s32 %r2, %r1, 128; // begin inline asm st.shared.v4.f32 [%r2], {%f5, %f6, %f7, %f8}; // end inline asm add.s32 %r3, %r1, 256; // begin inline asm st.shared.v4.f32 [%r3], {%f9, %f10, %f11, %f12}; // end inline asm add.s32 %r4, %r1, 384; // begin inline asm st.shared.v4.f32 [%r4], {%f13, %f14, %f15, %f16}; // end inline asm bar.sync 0; max.u32 %r26, %r21, %r23; sub.s32 %r27, %r26, %r23; setp.lt.u32 %p1, %r24, %r22; setp.ne.s32 %p2, %r27, 0; and.pred %p3, %p1, %p2; ld.shared.f32 %f17, [%r25]; selp.u32 %r5, 1, 0, %p3; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r5, 0; @p st.global.f32 [%rd1], %f17;} // end inline asm mul.wide.u32 %rd18, %r22, 4; add.s64 %rd2, %rd1, %rd18; setp.gt.u32 %p4, %r27, 1; and.pred %p5, %p1, %p4; ld.shared.f32 %f18, [%r25+128]; selp.u32 %r6, 1, 0, %p5; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r6, 0; @p st.global.f32 [%rd2], %f18;} // end inline asm shl.b32 %r28, %r22, 1; mul.wide.u32 %rd19, %r28, 4; add.s64 %rd3, %rd1, %rd19; setp.gt.u32 %p6, %r27, 2; and.pred %p7, %p1, %p6; ld.shared.f32 %f19, [%r25+256]; selp.u32 %r7, 1, 0, %p7; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r7, 0; @p st.global.f32 [%rd3], %f19;} // end inline asm mul.lo.s32 %r29, %r22, 3; mul.wide.u32 %rd20, %r29, 4; add.s64 %rd4, %rd1, %rd20; setp.gt.u32 %p8, %r27, 3; and.pred %p9, %p1, %p8; ld.shared.f32 %f20, [%r25+384]; selp.u32 %r8, 1, 0, %p9; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r8, 0; @p st.global.f32 [%rd4], %f20;} // end inline asm shl.b32 %r30, %r22, 2; mul.wide.u32 %rd21, %r30, 4; add.s64 %rd5, %rd1, %rd21; setp.gt.u32 %p10, %r27, 4; and.pred %p11, %p1, %p10; ld.shared.f32 %f21, [%r25+512]; selp.u32 %r9, 1, 0, %p11; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r9, 0; @p st.global.f32 [%rd5], %f21;} // end inline asm mul.lo.s32 %r31, %r22, 5; mul.wide.u32 %rd22, %r31, 4; add.s64 %rd6, %rd1, %rd22; setp.gt.u32 %p12, %r27, 5; and.pred %p13, %p1, %p12; ld.shared.f32 %f22, [%r25+640]; selp.u32 %r10, 1, 0, %p13; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r10, 0; @p st.global.f32 [%rd6], %f22;} // end inline asm mul.lo.s32 %r32, %r22, 6; mul.wide.u32 %rd23, %r32, 4; add.s64 %rd7, %rd1, %rd23; setp.gt.u32 %p14, %r27, 6; and.pred %p15, %p1, %p14; ld.shared.f32 %f23, [%r25+768]; selp.u32 %r11, 1, 0, %p15; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r11, 0; @p st.global.f32 [%rd7], %f23;} // end inline asm mul.lo.s32 %r33, %r22, 7; mul.wide.u32 %rd24, %r33, 4; add.s64 %rd8, %rd1, %rd24; setp.gt.u32 %p16, %r27, 7; and.pred %p17, %p1, %p16; ld.shared.f32 %f24, [%r25+896]; selp.u32 %r12, 1, 0, %p17; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r12, 0; @p st.global.f32 [%rd8], %f24;} // end inline asm shl.b32 %r34, %r22, 3; mul.wide.u32 %rd25, %r34, 4; add.s64 %rd9, %rd1, %rd25; setp.gt.u32 %p18, %r27, 8; and.pred %p19, %p1, %p18; ld.shared.f32 %f25, [%r25+1024]; selp.u32 %r13, 1, 0, %p19; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r13, 0; @p st.global.f32 [%rd9], %f25;} // end inline asm mul.lo.s32 %r35, %r22, 9; mul.wide.u32 %rd26, %r35, 4; add.s64 %rd10, %rd1, %rd26; setp.gt.u32 %p20, %r27, 9; and.pred %p21, %p1, %p20; ld.shared.f32 %f26, [%r25+1152]; selp.u32 %r14, 1, 0, %p21; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r14, 0; @p st.global.f32 [%rd10], %f26;} // end inline asm mul.lo.s32 %r36, %r22, 10; mul.wide.u32 %rd27, %r36, 4; add.s64 %rd11, %rd1, %rd27; setp.gt.u32 %p22, %r27, 10; and.pred %p23, %p1, %p22; ld.shared.f32 %f27, [%r25+1280]; selp.u32 %r15, 1, 0, %p23; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r15, 0; @p st.global.f32 [%rd11], %f27;} // end inline asm mul.lo.s32 %r37, %r22, 11; mul.wide.u32 %rd28, %r37, 4; add.s64 %rd12, %rd1, %rd28; setp.gt.u32 %p24, %r27, 11; and.pred %p25, %p1, %p24; ld.shared.f32 %f28, [%r25+1408]; selp.u32 %r16, 1, 0, %p25; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r16, 0; @p st.global.f32 [%rd12], %f28;} // end inline asm mul.lo.s32 %r38, %r22, 12; mul.wide.u32 %rd29, %r38, 4; add.s64 %rd13, %rd1, %rd29; setp.gt.u32 %p26, %r27, 12; and.pred %p27, %p1, %p26; ld.shared.f32 %f29, [%r25+1536]; selp.u32 %r17, 1, 0, %p27; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r17, 0; @p st.global.f32 [%rd13], %f29;} // end inline asm mul.lo.s32 %r39, %r22, 13; mul.wide.u32 %rd30, %r39, 4; add.s64 %rd14, %rd1, %rd30; setp.gt.u32 %p28, %r27, 13; and.pred %p29, %p1, %p28; ld.shared.f32 %f30, [%r25+1664]; selp.u32 %r18, 1, 0, %p29; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r18, 0; @p st.global.f32 [%rd14], %f30;} // end inline asm mul.lo.s32 %r40, %r22, 14; mul.wide.u32 %rd31, %r40, 4; add.s64 %rd15, %rd1, %rd31; setp.gt.u32 %p30, %r27, 14; and.pred %p31, %p1, %p30; ld.shared.f32 %f31, [%r25+1792]; selp.u32 %r19, 1, 0, %p31; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r19, 0; @p st.global.f32 [%rd15], %f31;} // end inline asm mul.lo.s32 %r41, %r22, 15; mul.wide.u32 %rd32, %r41, 4; add.s64 %rd16, %rd1, %rd32; setp.gt.u32 %p32, %r27, 15; and.pred %p33, %p1, %p32; ld.shared.f32 %f32, [%r25+1920]; selp.u32 %r20, 1, 0, %p33; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r20, 0; @p st.global.f32 [%rd16], %f32;} // end inline asm ret; } // .globl _Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj .visible .entry _Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj( .param .u64 _Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_0, .param .u64 _Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_1, .param .u64 _Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_2, .param .u32 _Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_3, .param .u32 _Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_4, .param .u32 _Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_5, .param .u32 _Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_6, .param .u32 _Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_7 ) .maxntid 256, 1, 1 .minnctapersm 2 { .reg .pred %p<33>; .reg .f32 %f<1921>; .reg .b32 %r<360>; .reg .b64 %rd<176>; // demoted variable .shared .align 16384 .b8 _ZZ22sgemm_128x128x8_kernelPKfS0_PfjjjjjE4smem[24576]; ld.param.u64 %rd26, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_0]; ld.param.u64 %rd27, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_1]; ld.param.u32 %r34, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_3]; ld.param.u32 %r35, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_4]; ld.param.u32 %r58, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_5]; ld.param.u32 %r59, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_6]; ld.param.u32 %r36, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_7]; mov.u32 %r60, %tid.x; and.b32 %r61, %r60, 31; shr.u32 %r1, %r60, 5; shr.u32 %r62, %r60, 1; and.b32 %r2, %r62, 7; shr.u32 %r63, %r61, 4; and.b32 %r64, %r60, 1; bfi.b32 %r3, %r63, %r64, 1, 31; mov.u32 %r65, %ctaid.y; shl.b32 %r4, %r65, 7; and.b32 %r66, %r62, 2147483644; add.s32 %r5, %r66, %r4; mul.lo.s32 %r67, %r5, %r58; cvt.u64.u32 %rd28, %r67; and.b32 %r68, %r60, 7; cvt.u64.u32 %rd29, %r68; add.s64 %rd30, %rd28, %rd29; shl.b64 %rd31, %rd30, 2; add.s64 %rd18, %rd26, %rd31; mul.lo.s32 %r69, %r1, %r35; mov.u32 %r70, %ctaid.x; bfi.b32 %r11, %r70, %r61, 7, 25; cvt.u64.u32 %rd32, %r11; mad.lo.s32 %r71, %r68, 132, %r66; mul.wide.u32 %rd33, %r71, 4; mov.u32 %r72, _ZZ22sgemm_128x128x8_kernelPKfS0_PfjjjjjE4smem; { .reg .b64 %tmp; cvt.u64.u32 %tmp, %r72; cvta.shared.u64 %rd34, %tmp; } add.s64 %rd14, %rd34, %rd33; // begin inline asm {.reg .u64 u64addr; cvta.to.shared.u64 u64addr, %rd14; cvt.u32.u64 %r355, u64addr;} // end inline asm bfi.b32 %r73, %r1, %r61, 7, 25; mul.wide.u32 %rd35, %r73, 4; add.s64 %rd36, %rd34, 16384; add.s64 %rd15, %rd36, %rd35; // begin inline asm {.reg .u64 u64addr; cvta.to.shared.u64 u64addr, %rd15; cvt.u32.u64 %r354, u64addr;} // end inline asm and.b32 %r8, %r62, 2147483616; shl.b32 %r74, %r3, 2; or.b32 %r75, %r74, %r8; mul.wide.u32 %rd37, %r75, 4; add.s64 %rd16, %rd34, %rd37; // begin inline asm {.reg .u64 u64addr; cvta.to.shared.u64 u64addr, %rd16; cvt.u32.u64 %r359, u64addr;} // end inline asm shl.b32 %r76, %r60, 1; and.b32 %r77, %r76, 92; mul.wide.u32 %rd38, %r77, 4; add.s64 %rd17, %rd36, %rd38; // begin inline asm {.reg .u64 u64addr; cvta.to.shared.u64 u64addr, %rd17; cvt.u32.u64 %r358, u64addr;} // end inline asm cvt.u64.u32 %rd39, %r69; setp.lt.u32 %p1, %r5, %r34; or.b32 %r78, %r5, 1; setp.lt.u32 %p2, %r78, %r34; or.b32 %r79, %r5, 2; setp.lt.u32 %p3, %r79, %r34; or.b32 %r80, %r5, 3; setp.lt.u32 %p4, %r80, %r34; add.s64 %rd40, %rd32, %rd39; setp.lt.u32 %p5, %r11, %r35; or.b32 %r81, %r11, 32; setp.lt.u32 %p6, %r81, %r35; or.b32 %r82, %r11, 64; setp.lt.u32 %p7, %r82, %r35; or.b32 %r83, %r11, 96; setp.lt.u32 %p8, %r83, %r35; add.s32 %r84, %r58, 7; shr.u32 %r85, %r84, 3; add.s32 %r353, %r85, -1; shl.b32 %r86, %r353, 3; sub.s32 %r87, %r58, %r86; setp.lt.u32 %p9, %r68, %r87; and.pred %p10, %p9, %p1; selp.u32 %r41, 1, 0, %p10; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r41, 0; @!p mov.b32 %f321, 0; @p ld.global.nc.f32 %f321, [%rd18];} // end inline asm and.pred %p11, %p9, %p2; cvt.u64.u32 %rd1, %r59; add.s64 %rd19, %rd18, %rd1; selp.u32 %r42, 1, 0, %p11; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r42, 0; @!p mov.b32 %f322, 0; @p ld.global.nc.f32 %f322, [%rd19];} // end inline asm and.pred %p12, %p9, %p3; shl.b32 %r88, %r59, 1; cvt.u64.u32 %rd2, %r88; add.s64 %rd20, %rd18, %rd2; selp.u32 %r43, 1, 0, %p12; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r43, 0; @!p mov.b32 %f323, 0; @p ld.global.nc.f32 %f323, [%rd20];} // end inline asm and.pred %p13, %p9, %p4; mul.lo.s32 %r89, %r59, 3; cvt.u64.u32 %rd3, %r89; add.s64 %rd21, %rd18, %rd3; selp.u32 %r44, 1, 0, %p13; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r44, 0; @!p mov.b32 %f324, 0; @p ld.global.nc.f32 %f324, [%rd21];} // end inline asm // begin inline asm st.shared.v4.f32 [%r355], {%f321, %f322, %f323, %f324}; // end inline asm setp.lt.u32 %p14, %r1, %r87; and.pred %p15, %p14, %p5; shl.b64 %rd41, %rd40, 2; add.s64 %rd22, %rd27, %rd41; selp.u32 %r46, 1, 0, %p15; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r46, 0; @!p mov.b32 %f329, 0; @p ld.global.nc.f32 %f329, [%rd22];} // end inline asm and.pred %p16, %p14, %p6; add.s64 %rd23, %rd22, 128; selp.u32 %r47, 1, 0, %p16; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r47, 0; @!p mov.b32 %f330, 0; @p ld.global.nc.f32 %f330, [%rd23];} // end inline asm and.pred %p17, %p14, %p7; add.s64 %rd24, %rd22, 256; selp.u32 %r48, 1, 0, %p17; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r48, 0; @!p mov.b32 %f331, 0; @p ld.global.nc.f32 %f331, [%rd24];} // end inline asm and.pred %p18, %p14, %p8; add.s64 %rd25, %rd22, 384; selp.u32 %r49, 1, 0, %p18; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r49, 0; @!p mov.b32 %f332, 0; @p ld.global.nc.f32 %f332, [%rd25];} // end inline asm // begin inline asm st.shared.f32 [%r354], %f329; // end inline asm add.s32 %r51, %r354, 128; // begin inline asm st.shared.f32 [%r51], %f330; // end inline asm add.s32 %r52, %r354, 256; // begin inline asm st.shared.f32 [%r52], %f331; // end inline asm add.s32 %r53, %r354, 384; // begin inline asm st.shared.f32 [%r53], %f332; // end inline asm bar.sync 0; mul.wide.u32 %rd42, %r87, 4; add.s64 %rd175, %rd18, %rd42; mul.lo.s32 %r90, %r87, %r35; mul.wide.u32 %rd43, %r90, 4; add.s64 %rd174, %rd22, %rd43; // begin inline asm ld.shared.v4.f32 {%f1913, %f1914, %f1915, %f1916}, [%r359]; // end inline asm add.s32 %r55, %r359, 64; // begin inline asm ld.shared.v4.f32 {%f1917, %f1918, %f1919, %f1920}, [%r55]; // end inline asm // begin inline asm ld.shared.v4.f32 {%f1848, %f1847, %f1846, %f1845}, [%r358]; // end inline asm add.s32 %r57, %r358, 128; // begin inline asm ld.shared.v4.f32 {%f1844, %f1843, %f1842, %f1841}, [%r57]; // end inline asm setp.eq.s32 %p19, %r353, 0; mov.f32 %f1849, 0f00000000; mov.f32 %f1850, %f1849; mov.f32 %f1851, %f1849; mov.f32 %f1852, %f1849; mov.f32 %f1853, %f1849; mov.f32 %f1854, %f1849; mov.f32 %f1855, %f1849; mov.f32 %f1856, %f1849; mov.f32 %f1857, %f1849; mov.f32 %f1858, %f1849; mov.f32 %f1859, %f1849; mov.f32 %f1860, %f1849; mov.f32 %f1861, %f1849; mov.f32 %f1862, %f1849; mov.f32 %f1863, %f1849; mov.f32 %f1864, %f1849; mov.f32 %f1865, %f1849; mov.f32 %f1866, %f1849; mov.f32 %f1867, %f1849; mov.f32 %f1868, %f1849; mov.f32 %f1869, %f1849; mov.f32 %f1870, %f1849; mov.f32 %f1871, %f1849; mov.f32 %f1872, %f1849; mov.f32 %f1873, %f1849; mov.f32 %f1874, %f1849; mov.f32 %f1875, %f1849; mov.f32 %f1876, %f1849; mov.f32 %f1877, %f1849; mov.f32 %f1878, %f1849; mov.f32 %f1879, %f1849; mov.f32 %f1880, %f1849; mov.f32 %f1881, %f1849; mov.f32 %f1882, %f1849; mov.f32 %f1883, %f1849; mov.f32 %f1884, %f1849; mov.f32 %f1885, %f1849; mov.f32 %f1886, %f1849; mov.f32 %f1887, %f1849; mov.f32 %f1888, %f1849; mov.f32 %f1889, %f1849; mov.f32 %f1890, %f1849; mov.f32 %f1891, %f1849; mov.f32 %f1892, %f1849; mov.f32 %f1893, %f1849; mov.f32 %f1894, %f1849; mov.f32 %f1895, %f1849; mov.f32 %f1896, %f1849; mov.f32 %f1897, %f1849; mov.f32 %f1898, %f1849; mov.f32 %f1899, %f1849; mov.f32 %f1900, %f1849; mov.f32 %f1901, %f1849; mov.f32 %f1902, %f1849; mov.f32 %f1903, %f1849; mov.f32 %f1904, %f1849; mov.f32 %f1905, %f1849; mov.f32 %f1906, %f1849; mov.f32 %f1907, %f1849; mov.f32 %f1908, %f1849; mov.f32 %f1909, %f1849; mov.f32 %f1910, %f1849; mov.f32 %f1911, %f1849; mov.f32 %f1912, %f1849; @%p19 bra $L__BB1_3; add.s32 %r91, %r11, 96; setp.lt.u32 %p20, %r91, %r35; add.s32 %r92, %r11, 64; setp.lt.u32 %p21, %r92, %r35; add.s32 %r93, %r11, 32; setp.lt.u32 %p22, %r93, %r35; add.s32 %r94, %r5, 3; setp.lt.u32 %p23, %r94, %r34; add.s32 %r95, %r5, 2; setp.lt.u32 %p24, %r95, %r34; add.s32 %r96, %r5, 1; setp.lt.u32 %p25, %r96, %r34; selp.u32 %r13, 1, 0, %p1; selp.u32 %r14, 1, 0, %p25; selp.u32 %r15, 1, 0, %p24; selp.u32 %r16, 1, 0, %p23; selp.u32 %r17, 1, 0, %p5; selp.u32 %r18, 1, 0, %p22; selp.u32 %r19, 1, 0, %p21; selp.u32 %r20, 1, 0, %p20; cvt.u64.u32 %rd8, %r36; $L__BB1_2: add.s32 %r97, %r359, 528; // begin inline asm ld.shared.v4.f32 {%f481, %f482, %f483, %f484}, [%r97]; // end inline asm add.s32 %r98, %r359, 592; // begin inline asm ld.shared.v4.f32 {%f485, %f486, %f487, %f488}, [%r98]; // end inline asm add.s32 %r99, %r358, 512; // begin inline asm ld.shared.v4.f32 {%f489, %f490, %f491, %f492}, [%r99]; // end inline asm add.s32 %r100, %r358, 640; // begin inline asm ld.shared.v4.f32 {%f493, %f494, %f495, %f496}, [%r100]; // end inline asm // begin inline asm {.reg .pred p; setp.ne.b32 p, %r13, 0; @p ld.global.nc.f32 %f497, [%rd175];} // end inline asm add.s64 %rd45, %rd175, %rd1; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r14, 0; @p ld.global.nc.f32 %f498, [%rd45];} // end inline asm add.s64 %rd46, %rd175, %rd2; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r15, 0; @p ld.global.nc.f32 %f499, [%rd46];} // end inline asm add.s64 %rd47, %rd175, %rd3; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r16, 0; @p ld.global.nc.f32 %f500, [%rd47];} // end inline asm // begin inline asm {.reg .pred p; setp.ne.b32 p, %r17, 0; @p ld.global.nc.f32 %f501, [%rd174];} // end inline asm add.s64 %rd49, %rd174, 128; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r18, 0; @p ld.global.nc.f32 %f502, [%rd49];} // end inline asm add.s64 %rd50, %rd174, 256; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r19, 0; @p ld.global.nc.f32 %f503, [%rd50];} // end inline asm add.s64 %rd51, %rd174, 384; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r20, 0; @p ld.global.nc.f32 %f504, [%rd51];} // end inline asm fma.rn.f32 %f625, %f1913, %f1848, %f1912; fma.rn.f32 %f626, %f1913, %f1847, %f1911; fma.rn.f32 %f627, %f1913, %f1846, %f1910; fma.rn.f32 %f628, %f1913, %f1845, %f1909; fma.rn.f32 %f629, %f1913, %f1844, %f1908; fma.rn.f32 %f630, %f1913, %f1843, %f1907; fma.rn.f32 %f631, %f1913, %f1842, %f1906; fma.rn.f32 %f632, %f1913, %f1841, %f1905; fma.rn.f32 %f633, %f1914, %f1848, %f1904; fma.rn.f32 %f634, %f1914, %f1847, %f1903; fma.rn.f32 %f635, %f1914, %f1846, %f1902; fma.rn.f32 %f636, %f1914, %f1845, %f1901; fma.rn.f32 %f637, %f1914, %f1844, %f1900; fma.rn.f32 %f638, %f1914, %f1843, %f1899; fma.rn.f32 %f639, %f1914, %f1842, %f1898; fma.rn.f32 %f640, %f1914, %f1841, %f1897; fma.rn.f32 %f641, %f1915, %f1848, %f1896; fma.rn.f32 %f642, %f1915, %f1847, %f1895; fma.rn.f32 %f643, %f1915, %f1846, %f1894; fma.rn.f32 %f644, %f1915, %f1845, %f1893; fma.rn.f32 %f645, %f1915, %f1844, %f1892; fma.rn.f32 %f646, %f1915, %f1843, %f1891; fma.rn.f32 %f647, %f1915, %f1842, %f1890; fma.rn.f32 %f648, %f1915, %f1841, %f1889; fma.rn.f32 %f649, %f1916, %f1848, %f1888; fma.rn.f32 %f650, %f1916, %f1847, %f1887; fma.rn.f32 %f651, %f1916, %f1846, %f1886; fma.rn.f32 %f652, %f1916, %f1845, %f1885; fma.rn.f32 %f653, %f1916, %f1844, %f1884; fma.rn.f32 %f654, %f1916, %f1843, %f1883; fma.rn.f32 %f655, %f1916, %f1842, %f1882; fma.rn.f32 %f656, %f1916, %f1841, %f1881; fma.rn.f32 %f657, %f1917, %f1848, %f1880; fma.rn.f32 %f658, %f1917, %f1847, %f1879; fma.rn.f32 %f659, %f1917, %f1846, %f1878; fma.rn.f32 %f660, %f1917, %f1845, %f1877; fma.rn.f32 %f661, %f1917, %f1844, %f1876; fma.rn.f32 %f662, %f1917, %f1843, %f1875; fma.rn.f32 %f663, %f1917, %f1842, %f1874; fma.rn.f32 %f664, %f1917, %f1841, %f1873; fma.rn.f32 %f665, %f1918, %f1848, %f1872; fma.rn.f32 %f666, %f1918, %f1847, %f1871; fma.rn.f32 %f667, %f1918, %f1846, %f1870; fma.rn.f32 %f668, %f1918, %f1845, %f1869; fma.rn.f32 %f669, %f1918, %f1844, %f1868; fma.rn.f32 %f670, %f1918, %f1843, %f1867; fma.rn.f32 %f671, %f1918, %f1842, %f1866; fma.rn.f32 %f672, %f1918, %f1841, %f1865; fma.rn.f32 %f673, %f1919, %f1848, %f1864; fma.rn.f32 %f674, %f1919, %f1847, %f1863; fma.rn.f32 %f675, %f1919, %f1846, %f1862; fma.rn.f32 %f676, %f1919, %f1845, %f1861; fma.rn.f32 %f677, %f1919, %f1844, %f1860; fma.rn.f32 %f678, %f1919, %f1843, %f1859; fma.rn.f32 %f679, %f1919, %f1842, %f1858; fma.rn.f32 %f680, %f1919, %f1841, %f1857; fma.rn.f32 %f681, %f1920, %f1848, %f1856; fma.rn.f32 %f682, %f1920, %f1847, %f1855; fma.rn.f32 %f683, %f1920, %f1846, %f1854; fma.rn.f32 %f684, %f1920, %f1845, %f1853; fma.rn.f32 %f685, %f1920, %f1844, %f1852; fma.rn.f32 %f686, %f1920, %f1843, %f1851; fma.rn.f32 %f687, %f1920, %f1842, %f1850; fma.rn.f32 %f688, %f1920, %f1841, %f1849; add.s32 %r109, %r359, 1056; // begin inline asm ld.shared.v4.f32 {%f505, %f506, %f507, %f508}, [%r109]; // end inline asm add.s32 %r110, %r359, 1120; // begin inline asm ld.shared.v4.f32 {%f509, %f510, %f511, %f512}, [%r110]; // end inline asm add.s32 %r111, %r358, 1024; // begin inline asm ld.shared.v4.f32 {%f513, %f514, %f515, %f516}, [%r111]; // end inline asm add.s32 %r112, %r358, 1152; // begin inline asm ld.shared.v4.f32 {%f517, %f518, %f519, %f520}, [%r112]; // end inline asm fma.rn.f32 %f689, %f481, %f489, %f625; fma.rn.f32 %f690, %f481, %f490, %f626; fma.rn.f32 %f691, %f481, %f491, %f627; fma.rn.f32 %f692, %f481, %f492, %f628; fma.rn.f32 %f693, %f481, %f493, %f629; fma.rn.f32 %f694, %f481, %f494, %f630; fma.rn.f32 %f695, %f481, %f495, %f631; fma.rn.f32 %f696, %f481, %f496, %f632; fma.rn.f32 %f697, %f482, %f489, %f633; fma.rn.f32 %f698, %f482, %f490, %f634; fma.rn.f32 %f699, %f482, %f491, %f635; fma.rn.f32 %f700, %f482, %f492, %f636; fma.rn.f32 %f701, %f482, %f493, %f637; fma.rn.f32 %f702, %f482, %f494, %f638; fma.rn.f32 %f703, %f482, %f495, %f639; fma.rn.f32 %f704, %f482, %f496, %f640; fma.rn.f32 %f705, %f483, %f489, %f641; fma.rn.f32 %f706, %f483, %f490, %f642; fma.rn.f32 %f707, %f483, %f491, %f643; fma.rn.f32 %f708, %f483, %f492, %f644; fma.rn.f32 %f709, %f483, %f493, %f645; fma.rn.f32 %f710, %f483, %f494, %f646; fma.rn.f32 %f711, %f483, %f495, %f647; fma.rn.f32 %f712, %f483, %f496, %f648; fma.rn.f32 %f713, %f484, %f489, %f649; fma.rn.f32 %f714, %f484, %f490, %f650; fma.rn.f32 %f715, %f484, %f491, %f651; fma.rn.f32 %f716, %f484, %f492, %f652; fma.rn.f32 %f717, %f484, %f493, %f653; fma.rn.f32 %f718, %f484, %f494, %f654; fma.rn.f32 %f719, %f484, %f495, %f655; fma.rn.f32 %f720, %f484, %f496, %f656; fma.rn.f32 %f721, %f485, %f489, %f657; fma.rn.f32 %f722, %f485, %f490, %f658; fma.rn.f32 %f723, %f485, %f491, %f659; fma.rn.f32 %f724, %f485, %f492, %f660; fma.rn.f32 %f725, %f485, %f493, %f661; fma.rn.f32 %f726, %f485, %f494, %f662; fma.rn.f32 %f727, %f485, %f495, %f663; fma.rn.f32 %f728, %f485, %f496, %f664; fma.rn.f32 %f729, %f486, %f489, %f665; fma.rn.f32 %f730, %f486, %f490, %f666; fma.rn.f32 %f731, %f486, %f491, %f667; fma.rn.f32 %f732, %f486, %f492, %f668; fma.rn.f32 %f733, %f486, %f493, %f669; fma.rn.f32 %f734, %f486, %f494, %f670; fma.rn.f32 %f735, %f486, %f495, %f671; fma.rn.f32 %f736, %f486, %f496, %f672; fma.rn.f32 %f737, %f487, %f489, %f673; fma.rn.f32 %f738, %f487, %f490, %f674; fma.rn.f32 %f739, %f487, %f491, %f675; fma.rn.f32 %f740, %f487, %f492, %f676; fma.rn.f32 %f741, %f487, %f493, %f677; fma.rn.f32 %f742, %f487, %f494, %f678; fma.rn.f32 %f743, %f487, %f495, %f679; fma.rn.f32 %f744, %f487, %f496, %f680; fma.rn.f32 %f745, %f488, %f489, %f681; fma.rn.f32 %f746, %f488, %f490, %f682; fma.rn.f32 %f747, %f488, %f491, %f683; fma.rn.f32 %f748, %f488, %f492, %f684; fma.rn.f32 %f749, %f488, %f493, %f685; fma.rn.f32 %f750, %f488, %f494, %f686; fma.rn.f32 %f751, %f488, %f495, %f687; fma.rn.f32 %f752, %f488, %f496, %f688; add.s32 %r113, %r359, 1584; // begin inline asm ld.shared.v4.f32 {%f521, %f522, %f523, %f524}, [%r113]; // end inline asm add.s32 %r114, %r359, 1648; // begin inline asm ld.shared.v4.f32 {%f525, %f526, %f527, %f528}, [%r114]; // end inline asm add.s32 %r115, %r358, 1536; // begin inline asm ld.shared.v4.f32 {%f529, %f530, %f531, %f532}, [%r115]; // end inline asm add.s32 %r116, %r358, 1664; // begin inline asm ld.shared.v4.f32 {%f533, %f534, %f535, %f536}, [%r116]; // end inline asm fma.rn.f32 %f753, %f505, %f513, %f689; fma.rn.f32 %f754, %f505, %f514, %f690; fma.rn.f32 %f755, %f505, %f515, %f691; fma.rn.f32 %f756, %f505, %f516, %f692; fma.rn.f32 %f757, %f505, %f517, %f693; fma.rn.f32 %f758, %f505, %f518, %f694; fma.rn.f32 %f759, %f505, %f519, %f695; fma.rn.f32 %f760, %f505, %f520, %f696; fma.rn.f32 %f761, %f506, %f513, %f697; fma.rn.f32 %f762, %f506, %f514, %f698; fma.rn.f32 %f763, %f506, %f515, %f699; fma.rn.f32 %f764, %f506, %f516, %f700; fma.rn.f32 %f765, %f506, %f517, %f701; fma.rn.f32 %f766, %f506, %f518, %f702; fma.rn.f32 %f767, %f506, %f519, %f703; fma.rn.f32 %f768, %f506, %f520, %f704; fma.rn.f32 %f769, %f507, %f513, %f705; fma.rn.f32 %f770, %f507, %f514, %f706; fma.rn.f32 %f771, %f507, %f515, %f707; fma.rn.f32 %f772, %f507, %f516, %f708; fma.rn.f32 %f773, %f507, %f517, %f709; fma.rn.f32 %f774, %f507, %f518, %f710; fma.rn.f32 %f775, %f507, %f519, %f711; fma.rn.f32 %f776, %f507, %f520, %f712; fma.rn.f32 %f777, %f508, %f513, %f713; fma.rn.f32 %f778, %f508, %f514, %f714; fma.rn.f32 %f779, %f508, %f515, %f715; fma.rn.f32 %f780, %f508, %f516, %f716; fma.rn.f32 %f781, %f508, %f517, %f717; fma.rn.f32 %f782, %f508, %f518, %f718; fma.rn.f32 %f783, %f508, %f519, %f719; fma.rn.f32 %f784, %f508, %f520, %f720; fma.rn.f32 %f785, %f509, %f513, %f721; fma.rn.f32 %f786, %f509, %f514, %f722; fma.rn.f32 %f787, %f509, %f515, %f723; fma.rn.f32 %f788, %f509, %f516, %f724; fma.rn.f32 %f789, %f509, %f517, %f725; fma.rn.f32 %f790, %f509, %f518, %f726; fma.rn.f32 %f791, %f509, %f519, %f727; fma.rn.f32 %f792, %f509, %f520, %f728; fma.rn.f32 %f793, %f510, %f513, %f729; fma.rn.f32 %f794, %f510, %f514, %f730; fma.rn.f32 %f795, %f510, %f515, %f731; fma.rn.f32 %f796, %f510, %f516, %f732; fma.rn.f32 %f797, %f510, %f517, %f733; fma.rn.f32 %f798, %f510, %f518, %f734; fma.rn.f32 %f799, %f510, %f519, %f735; fma.rn.f32 %f800, %f510, %f520, %f736; fma.rn.f32 %f801, %f511, %f513, %f737; fma.rn.f32 %f802, %f511, %f514, %f738; fma.rn.f32 %f803, %f511, %f515, %f739; fma.rn.f32 %f804, %f511, %f516, %f740; fma.rn.f32 %f805, %f511, %f517, %f741; fma.rn.f32 %f806, %f511, %f518, %f742; fma.rn.f32 %f807, %f511, %f519, %f743; fma.rn.f32 %f808, %f511, %f520, %f744; fma.rn.f32 %f809, %f512, %f513, %f745; fma.rn.f32 %f810, %f512, %f514, %f746; fma.rn.f32 %f811, %f512, %f515, %f747; fma.rn.f32 %f812, %f512, %f516, %f748; fma.rn.f32 %f813, %f512, %f517, %f749; fma.rn.f32 %f814, %f512, %f518, %f750; fma.rn.f32 %f815, %f512, %f519, %f751; fma.rn.f32 %f816, %f512, %f520, %f752; add.s32 %r117, %r359, 2112; // begin inline asm ld.shared.v4.f32 {%f537, %f538, %f539, %f540}, [%r117]; // end inline asm add.s32 %r118, %r359, 2176; // begin inline asm ld.shared.v4.f32 {%f541, %f542, %f543, %f544}, [%r118]; // end inline asm add.s32 %r119, %r358, 2048; // begin inline asm ld.shared.v4.f32 {%f545, %f546, %f547, %f548}, [%r119]; // end inline asm add.s32 %r120, %r358, 2176; // begin inline asm ld.shared.v4.f32 {%f549, %f550, %f551, %f552}, [%r120]; // end inline asm fma.rn.f32 %f817, %f521, %f529, %f753; fma.rn.f32 %f818, %f521, %f530, %f754; fma.rn.f32 %f819, %f521, %f531, %f755; fma.rn.f32 %f820, %f521, %f532, %f756; fma.rn.f32 %f821, %f521, %f533, %f757; fma.rn.f32 %f822, %f521, %f534, %f758; fma.rn.f32 %f823, %f521, %f535, %f759; fma.rn.f32 %f824, %f521, %f536, %f760; fma.rn.f32 %f825, %f522, %f529, %f761; fma.rn.f32 %f826, %f522, %f530, %f762; fma.rn.f32 %f827, %f522, %f531, %f763; fma.rn.f32 %f828, %f522, %f532, %f764; fma.rn.f32 %f829, %f522, %f533, %f765; fma.rn.f32 %f830, %f522, %f534, %f766; fma.rn.f32 %f831, %f522, %f535, %f767; fma.rn.f32 %f832, %f522, %f536, %f768; fma.rn.f32 %f833, %f523, %f529, %f769; fma.rn.f32 %f834, %f523, %f530, %f770; fma.rn.f32 %f835, %f523, %f531, %f771; fma.rn.f32 %f836, %f523, %f532, %f772; fma.rn.f32 %f837, %f523, %f533, %f773; fma.rn.f32 %f838, %f523, %f534, %f774; fma.rn.f32 %f839, %f523, %f535, %f775; fma.rn.f32 %f840, %f523, %f536, %f776; fma.rn.f32 %f841, %f524, %f529, %f777; fma.rn.f32 %f842, %f524, %f530, %f778; fma.rn.f32 %f843, %f524, %f531, %f779; fma.rn.f32 %f844, %f524, %f532, %f780; fma.rn.f32 %f845, %f524, %f533, %f781; fma.rn.f32 %f846, %f524, %f534, %f782; fma.rn.f32 %f847, %f524, %f535, %f783; fma.rn.f32 %f848, %f524, %f536, %f784; fma.rn.f32 %f849, %f525, %f529, %f785; fma.rn.f32 %f850, %f525, %f530, %f786; fma.rn.f32 %f851, %f525, %f531, %f787; fma.rn.f32 %f852, %f525, %f532, %f788; fma.rn.f32 %f853, %f525, %f533, %f789; fma.rn.f32 %f854, %f525, %f534, %f790; fma.rn.f32 %f855, %f525, %f535, %f791; fma.rn.f32 %f856, %f525, %f536, %f792; fma.rn.f32 %f857, %f526, %f529, %f793; fma.rn.f32 %f858, %f526, %f530, %f794; fma.rn.f32 %f859, %f526, %f531, %f795; fma.rn.f32 %f860, %f526, %f532, %f796; fma.rn.f32 %f861, %f526, %f533, %f797; fma.rn.f32 %f862, %f526, %f534, %f798; fma.rn.f32 %f863, %f526, %f535, %f799; fma.rn.f32 %f864, %f526, %f536, %f800; fma.rn.f32 %f865, %f527, %f529, %f801; fma.rn.f32 %f866, %f527, %f530, %f802; fma.rn.f32 %f867, %f527, %f531, %f803; fma.rn.f32 %f868, %f527, %f532, %f804; fma.rn.f32 %f869, %f527, %f533, %f805; fma.rn.f32 %f870, %f527, %f534, %f806; fma.rn.f32 %f871, %f527, %f535, %f807; fma.rn.f32 %f872, %f527, %f536, %f808; fma.rn.f32 %f873, %f528, %f529, %f809; fma.rn.f32 %f874, %f528, %f530, %f810; fma.rn.f32 %f875, %f528, %f531, %f811; fma.rn.f32 %f876, %f528, %f532, %f812; fma.rn.f32 %f877, %f528, %f533, %f813; fma.rn.f32 %f878, %f528, %f534, %f814; fma.rn.f32 %f879, %f528, %f535, %f815; fma.rn.f32 %f880, %f528, %f536, %f816; add.s32 %r121, %r359, 2640; // begin inline asm ld.shared.v4.f32 {%f553, %f554, %f555, %f556}, [%r121]; // end inline asm add.s32 %r122, %r359, 2704; // begin inline asm ld.shared.v4.f32 {%f557, %f558, %f559, %f560}, [%r122]; // end inline asm add.s32 %r123, %r358, 2560; // begin inline asm ld.shared.v4.f32 {%f561, %f562, %f563, %f564}, [%r123]; // end inline asm add.s32 %r124, %r358, 2688; // begin inline asm ld.shared.v4.f32 {%f565, %f566, %f567, %f568}, [%r124]; // end inline asm fma.rn.f32 %f881, %f537, %f545, %f817; fma.rn.f32 %f882, %f537, %f546, %f818; fma.rn.f32 %f883, %f537, %f547, %f819; fma.rn.f32 %f884, %f537, %f548, %f820; fma.rn.f32 %f885, %f537, %f549, %f821; fma.rn.f32 %f886, %f537, %f550, %f822; fma.rn.f32 %f887, %f537, %f551, %f823; fma.rn.f32 %f888, %f537, %f552, %f824; fma.rn.f32 %f889, %f538, %f545, %f825; fma.rn.f32 %f890, %f538, %f546, %f826; fma.rn.f32 %f891, %f538, %f547, %f827; fma.rn.f32 %f892, %f538, %f548, %f828; fma.rn.f32 %f893, %f538, %f549, %f829; fma.rn.f32 %f894, %f538, %f550, %f830; fma.rn.f32 %f895, %f538, %f551, %f831; fma.rn.f32 %f896, %f538, %f552, %f832; fma.rn.f32 %f897, %f539, %f545, %f833; fma.rn.f32 %f898, %f539, %f546, %f834; fma.rn.f32 %f899, %f539, %f547, %f835; fma.rn.f32 %f900, %f539, %f548, %f836; fma.rn.f32 %f901, %f539, %f549, %f837; fma.rn.f32 %f902, %f539, %f550, %f838; fma.rn.f32 %f903, %f539, %f551, %f839; fma.rn.f32 %f904, %f539, %f552, %f840; fma.rn.f32 %f905, %f540, %f545, %f841; fma.rn.f32 %f906, %f540, %f546, %f842; fma.rn.f32 %f907, %f540, %f547, %f843; fma.rn.f32 %f908, %f540, %f548, %f844; fma.rn.f32 %f909, %f540, %f549, %f845; fma.rn.f32 %f910, %f540, %f550, %f846; fma.rn.f32 %f911, %f540, %f551, %f847; fma.rn.f32 %f912, %f540, %f552, %f848; fma.rn.f32 %f913, %f541, %f545, %f849; fma.rn.f32 %f914, %f541, %f546, %f850; fma.rn.f32 %f915, %f541, %f547, %f851; fma.rn.f32 %f916, %f541, %f548, %f852; fma.rn.f32 %f917, %f541, %f549, %f853; fma.rn.f32 %f918, %f541, %f550, %f854; fma.rn.f32 %f919, %f541, %f551, %f855; fma.rn.f32 %f920, %f541, %f552, %f856; fma.rn.f32 %f921, %f542, %f545, %f857; fma.rn.f32 %f922, %f542, %f546, %f858; fma.rn.f32 %f923, %f542, %f547, %f859; fma.rn.f32 %f924, %f542, %f548, %f860; fma.rn.f32 %f925, %f542, %f549, %f861; fma.rn.f32 %f926, %f542, %f550, %f862; fma.rn.f32 %f927, %f542, %f551, %f863; fma.rn.f32 %f928, %f542, %f552, %f864; fma.rn.f32 %f929, %f543, %f545, %f865; fma.rn.f32 %f930, %f543, %f546, %f866; fma.rn.f32 %f931, %f543, %f547, %f867; fma.rn.f32 %f932, %f543, %f548, %f868; fma.rn.f32 %f933, %f543, %f549, %f869; fma.rn.f32 %f934, %f543, %f550, %f870; fma.rn.f32 %f935, %f543, %f551, %f871; fma.rn.f32 %f936, %f543, %f552, %f872; fma.rn.f32 %f937, %f544, %f545, %f873; fma.rn.f32 %f938, %f544, %f546, %f874; fma.rn.f32 %f939, %f544, %f547, %f875; fma.rn.f32 %f940, %f544, %f548, %f876; fma.rn.f32 %f941, %f544, %f549, %f877; fma.rn.f32 %f942, %f544, %f550, %f878; fma.rn.f32 %f943, %f544, %f551, %f879; fma.rn.f32 %f944, %f544, %f552, %f880; add.s32 %r125, %r359, 3168; // begin inline asm ld.shared.v4.f32 {%f569, %f570, %f571, %f572}, [%r125]; // end inline asm add.s32 %r126, %r359, 3232; // begin inline asm ld.shared.v4.f32 {%f573, %f574, %f575, %f576}, [%r126]; // end inline asm add.s32 %r127, %r358, 3072; // begin inline asm ld.shared.v4.f32 {%f577, %f578, %f579, %f580}, [%r127]; // end inline asm add.s32 %r128, %r358, 3200; // begin inline asm ld.shared.v4.f32 {%f581, %f582, %f583, %f584}, [%r128]; // end inline asm fma.rn.f32 %f945, %f553, %f561, %f881; fma.rn.f32 %f946, %f553, %f562, %f882; fma.rn.f32 %f947, %f553, %f563, %f883; fma.rn.f32 %f948, %f553, %f564, %f884; fma.rn.f32 %f949, %f553, %f565, %f885; fma.rn.f32 %f950, %f553, %f566, %f886; fma.rn.f32 %f951, %f553, %f567, %f887; fma.rn.f32 %f952, %f553, %f568, %f888; fma.rn.f32 %f953, %f554, %f561, %f889; fma.rn.f32 %f954, %f554, %f562, %f890; fma.rn.f32 %f955, %f554, %f563, %f891; fma.rn.f32 %f956, %f554, %f564, %f892; fma.rn.f32 %f957, %f554, %f565, %f893; fma.rn.f32 %f958, %f554, %f566, %f894; fma.rn.f32 %f959, %f554, %f567, %f895; fma.rn.f32 %f960, %f554, %f568, %f896; fma.rn.f32 %f961, %f555, %f561, %f897; fma.rn.f32 %f962, %f555, %f562, %f898; fma.rn.f32 %f963, %f555, %f563, %f899; fma.rn.f32 %f964, %f555, %f564, %f900; fma.rn.f32 %f965, %f555, %f565, %f901; fma.rn.f32 %f966, %f555, %f566, %f902; fma.rn.f32 %f967, %f555, %f567, %f903; fma.rn.f32 %f968, %f555, %f568, %f904; fma.rn.f32 %f969, %f556, %f561, %f905; fma.rn.f32 %f970, %f556, %f562, %f906; fma.rn.f32 %f971, %f556, %f563, %f907; fma.rn.f32 %f972, %f556, %f564, %f908; fma.rn.f32 %f973, %f556, %f565, %f909; fma.rn.f32 %f974, %f556, %f566, %f910; fma.rn.f32 %f975, %f556, %f567, %f911; fma.rn.f32 %f976, %f556, %f568, %f912; fma.rn.f32 %f977, %f557, %f561, %f913; fma.rn.f32 %f978, %f557, %f562, %f914; fma.rn.f32 %f979, %f557, %f563, %f915; fma.rn.f32 %f980, %f557, %f564, %f916; fma.rn.f32 %f981, %f557, %f565, %f917; fma.rn.f32 %f982, %f557, %f566, %f918; fma.rn.f32 %f983, %f557, %f567, %f919; fma.rn.f32 %f984, %f557, %f568, %f920; fma.rn.f32 %f985, %f558, %f561, %f921; fma.rn.f32 %f986, %f558, %f562, %f922; fma.rn.f32 %f987, %f558, %f563, %f923; fma.rn.f32 %f988, %f558, %f564, %f924; fma.rn.f32 %f989, %f558, %f565, %f925; fma.rn.f32 %f990, %f558, %f566, %f926; fma.rn.f32 %f991, %f558, %f567, %f927; fma.rn.f32 %f992, %f558, %f568, %f928; fma.rn.f32 %f993, %f559, %f561, %f929; fma.rn.f32 %f994, %f559, %f562, %f930; fma.rn.f32 %f995, %f559, %f563, %f931; fma.rn.f32 %f996, %f559, %f564, %f932; fma.rn.f32 %f997, %f559, %f565, %f933; fma.rn.f32 %f998, %f559, %f566, %f934; fma.rn.f32 %f999, %f559, %f567, %f935; fma.rn.f32 %f1000, %f559, %f568, %f936; fma.rn.f32 %f1001, %f560, %f561, %f937; fma.rn.f32 %f1002, %f560, %f562, %f938; fma.rn.f32 %f1003, %f560, %f563, %f939; fma.rn.f32 %f1004, %f560, %f564, %f940; fma.rn.f32 %f1005, %f560, %f565, %f941; fma.rn.f32 %f1006, %f560, %f566, %f942; fma.rn.f32 %f1007, %f560, %f567, %f943; fma.rn.f32 %f1008, %f560, %f568, %f944; add.s32 %r129, %r359, 3696; // begin inline asm ld.shared.v4.f32 {%f585, %f586, %f587, %f588}, [%r129]; // end inline asm add.s32 %r130, %r359, 3760; // begin inline asm ld.shared.v4.f32 {%f589, %f590, %f591, %f592}, [%r130]; // end inline asm add.s32 %r131, %r358, 3584; // begin inline asm ld.shared.v4.f32 {%f593, %f594, %f595, %f596}, [%r131]; // end inline asm add.s32 %r132, %r358, 3712; // begin inline asm ld.shared.v4.f32 {%f597, %f598, %f599, %f600}, [%r132]; // end inline asm fma.rn.f32 %f1009, %f569, %f577, %f945; fma.rn.f32 %f1010, %f569, %f578, %f946; fma.rn.f32 %f1011, %f569, %f579, %f947; fma.rn.f32 %f1012, %f569, %f580, %f948; fma.rn.f32 %f1013, %f569, %f581, %f949; fma.rn.f32 %f1014, %f569, %f582, %f950; fma.rn.f32 %f1015, %f569, %f583, %f951; fma.rn.f32 %f1016, %f569, %f584, %f952; fma.rn.f32 %f1017, %f570, %f577, %f953; fma.rn.f32 %f1018, %f570, %f578, %f954; fma.rn.f32 %f1019, %f570, %f579, %f955; fma.rn.f32 %f1020, %f570, %f580, %f956; fma.rn.f32 %f1021, %f570, %f581, %f957; fma.rn.f32 %f1022, %f570, %f582, %f958; fma.rn.f32 %f1023, %f570, %f583, %f959; fma.rn.f32 %f1024, %f570, %f584, %f960; fma.rn.f32 %f1025, %f571, %f577, %f961; fma.rn.f32 %f1026, %f571, %f578, %f962; fma.rn.f32 %f1027, %f571, %f579, %f963; fma.rn.f32 %f1028, %f571, %f580, %f964; fma.rn.f32 %f1029, %f571, %f581, %f965; fma.rn.f32 %f1030, %f571, %f582, %f966; fma.rn.f32 %f1031, %f571, %f583, %f967; fma.rn.f32 %f1032, %f571, %f584, %f968; fma.rn.f32 %f1033, %f572, %f577, %f969; fma.rn.f32 %f1034, %f572, %f578, %f970; fma.rn.f32 %f1035, %f572, %f579, %f971; fma.rn.f32 %f1036, %f572, %f580, %f972; fma.rn.f32 %f1037, %f572, %f581, %f973; fma.rn.f32 %f1038, %f572, %f582, %f974; fma.rn.f32 %f1039, %f572, %f583, %f975; fma.rn.f32 %f1040, %f572, %f584, %f976; fma.rn.f32 %f1041, %f573, %f577, %f977; fma.rn.f32 %f1042, %f573, %f578, %f978; fma.rn.f32 %f1043, %f573, %f579, %f979; fma.rn.f32 %f1044, %f573, %f580, %f980; fma.rn.f32 %f1045, %f573, %f581, %f981; fma.rn.f32 %f1046, %f573, %f582, %f982; fma.rn.f32 %f1047, %f573, %f583, %f983; fma.rn.f32 %f1048, %f573, %f584, %f984; fma.rn.f32 %f1049, %f574, %f577, %f985; fma.rn.f32 %f1050, %f574, %f578, %f986; fma.rn.f32 %f1051, %f574, %f579, %f987; fma.rn.f32 %f1052, %f574, %f580, %f988; fma.rn.f32 %f1053, %f574, %f581, %f989; fma.rn.f32 %f1054, %f574, %f582, %f990; fma.rn.f32 %f1055, %f574, %f583, %f991; fma.rn.f32 %f1056, %f574, %f584, %f992; fma.rn.f32 %f1057, %f575, %f577, %f993; fma.rn.f32 %f1058, %f575, %f578, %f994; fma.rn.f32 %f1059, %f575, %f579, %f995; fma.rn.f32 %f1060, %f575, %f580, %f996; fma.rn.f32 %f1061, %f575, %f581, %f997; fma.rn.f32 %f1062, %f575, %f582, %f998; fma.rn.f32 %f1063, %f575, %f583, %f999; fma.rn.f32 %f1064, %f575, %f584, %f1000; fma.rn.f32 %f1065, %f576, %f577, %f1001; fma.rn.f32 %f1066, %f576, %f578, %f1002; fma.rn.f32 %f1067, %f576, %f579, %f1003; fma.rn.f32 %f1068, %f576, %f580, %f1004; fma.rn.f32 %f1069, %f576, %f581, %f1005; fma.rn.f32 %f1070, %f576, %f582, %f1006; fma.rn.f32 %f1071, %f576, %f583, %f1007; fma.rn.f32 %f1072, %f576, %f584, %f1008; xor.b32 %r355, %r355, 8192; // begin inline asm st.shared.v4.f32 [%r355], {%f497, %f498, %f499, %f500}; // end inline asm xor.b32 %r354, %r354, 4096; // begin inline asm st.shared.f32 [%r354], %f501; // end inline asm add.s32 %r135, %r354, 128; // begin inline asm st.shared.f32 [%r135], %f502; // end inline asm add.s32 %r136, %r354, 256; // begin inline asm st.shared.f32 [%r136], %f503; // end inline asm add.s32 %r137, %r354, 384; // begin inline asm st.shared.f32 [%r137], %f504; // end inline asm bar.sync 0; add.s64 %rd175, %rd175, 32; add.s64 %rd174, %rd174, %rd8; xor.b32 %r359, %r359, 8192; // begin inline asm ld.shared.v4.f32 {%f1913, %f1914, %f1915, %f1916}, [%r359]; // end inline asm add.s32 %r139, %r359, 64; // begin inline asm ld.shared.v4.f32 {%f1917, %f1918, %f1919, %f1920}, [%r139]; // end inline asm xor.b32 %r358, %r358, 4096; // begin inline asm ld.shared.v4.f32 {%f1848, %f1847, %f1846, %f1845}, [%r358]; // end inline asm add.s32 %r141, %r358, 128; // begin inline asm ld.shared.v4.f32 {%f1844, %f1843, %f1842, %f1841}, [%r141]; // end inline asm fma.rn.f32 %f1912, %f585, %f593, %f1009; fma.rn.f32 %f1911, %f585, %f594, %f1010; fma.rn.f32 %f1910, %f585, %f595, %f1011; fma.rn.f32 %f1909, %f585, %f596, %f1012; fma.rn.f32 %f1908, %f585, %f597, %f1013; fma.rn.f32 %f1907, %f585, %f598, %f1014; fma.rn.f32 %f1906, %f585, %f599, %f1015; fma.rn.f32 %f1905, %f585, %f600, %f1016; fma.rn.f32 %f1904, %f586, %f593, %f1017; fma.rn.f32 %f1903, %f586, %f594, %f1018; fma.rn.f32 %f1902, %f586, %f595, %f1019; fma.rn.f32 %f1901, %f586, %f596, %f1020; fma.rn.f32 %f1900, %f586, %f597, %f1021; fma.rn.f32 %f1899, %f586, %f598, %f1022; fma.rn.f32 %f1898, %f586, %f599, %f1023; fma.rn.f32 %f1897, %f586, %f600, %f1024; fma.rn.f32 %f1896, %f587, %f593, %f1025; fma.rn.f32 %f1895, %f587, %f594, %f1026; fma.rn.f32 %f1894, %f587, %f595, %f1027; fma.rn.f32 %f1893, %f587, %f596, %f1028; fma.rn.f32 %f1892, %f587, %f597, %f1029; fma.rn.f32 %f1891, %f587, %f598, %f1030; fma.rn.f32 %f1890, %f587, %f599, %f1031; fma.rn.f32 %f1889, %f587, %f600, %f1032; fma.rn.f32 %f1888, %f588, %f593, %f1033; fma.rn.f32 %f1887, %f588, %f594, %f1034; fma.rn.f32 %f1886, %f588, %f595, %f1035; fma.rn.f32 %f1885, %f588, %f596, %f1036; fma.rn.f32 %f1884, %f588, %f597, %f1037; fma.rn.f32 %f1883, %f588, %f598, %f1038; fma.rn.f32 %f1882, %f588, %f599, %f1039; fma.rn.f32 %f1881, %f588, %f600, %f1040; fma.rn.f32 %f1880, %f589, %f593, %f1041; fma.rn.f32 %f1879, %f589, %f594, %f1042; fma.rn.f32 %f1878, %f589, %f595, %f1043; fma.rn.f32 %f1877, %f589, %f596, %f1044; fma.rn.f32 %f1876, %f589, %f597, %f1045; fma.rn.f32 %f1875, %f589, %f598, %f1046; fma.rn.f32 %f1874, %f589, %f599, %f1047; fma.rn.f32 %f1873, %f589, %f600, %f1048; fma.rn.f32 %f1872, %f590, %f593, %f1049; fma.rn.f32 %f1871, %f590, %f594, %f1050; fma.rn.f32 %f1870, %f590, %f595, %f1051; fma.rn.f32 %f1869, %f590, %f596, %f1052; fma.rn.f32 %f1868, %f590, %f597, %f1053; fma.rn.f32 %f1867, %f590, %f598, %f1054; fma.rn.f32 %f1866, %f590, %f599, %f1055; fma.rn.f32 %f1865, %f590, %f600, %f1056; fma.rn.f32 %f1864, %f591, %f593, %f1057; fma.rn.f32 %f1863, %f591, %f594, %f1058; fma.rn.f32 %f1862, %f591, %f595, %f1059; fma.rn.f32 %f1861, %f591, %f596, %f1060; fma.rn.f32 %f1860, %f591, %f597, %f1061; fma.rn.f32 %f1859, %f591, %f598, %f1062; fma.rn.f32 %f1858, %f591, %f599, %f1063; fma.rn.f32 %f1857, %f591, %f600, %f1064; fma.rn.f32 %f1856, %f592, %f593, %f1065; fma.rn.f32 %f1855, %f592, %f594, %f1066; fma.rn.f32 %f1854, %f592, %f595, %f1067; fma.rn.f32 %f1853, %f592, %f596, %f1068; fma.rn.f32 %f1852, %f592, %f597, %f1069; fma.rn.f32 %f1851, %f592, %f598, %f1070; fma.rn.f32 %f1850, %f592, %f599, %f1071; fma.rn.f32 %f1849, %f592, %f600, %f1072; add.s32 %r353, %r353, -1; setp.ne.s32 %p28, %r353, 0; @%p28 bra $L__BB1_2; $L__BB1_3: add.s32 %r142, %r359, 528; // begin inline asm ld.shared.v4.f32 {%f1073, %f1074, %f1075, %f1076}, [%r142]; // end inline asm add.s32 %r143, %r359, 592; // begin inline asm ld.shared.v4.f32 {%f1077, %f1078, %f1079, %f1080}, [%r143]; // end inline asm add.s32 %r144, %r358, 512; // begin inline asm ld.shared.v4.f32 {%f1081, %f1082, %f1083, %f1084}, [%r144]; // end inline asm add.s32 %r145, %r358, 640; // begin inline asm ld.shared.v4.f32 {%f1085, %f1086, %f1087, %f1088}, [%r145]; // end inline asm fma.rn.f32 %f1185, %f1913, %f1848, %f1912; fma.rn.f32 %f1186, %f1913, %f1847, %f1911; fma.rn.f32 %f1187, %f1913, %f1846, %f1910; fma.rn.f32 %f1188, %f1913, %f1845, %f1909; fma.rn.f32 %f1189, %f1913, %f1844, %f1908; fma.rn.f32 %f1190, %f1913, %f1843, %f1907; fma.rn.f32 %f1191, %f1913, %f1842, %f1906; fma.rn.f32 %f1192, %f1913, %f1841, %f1905; fma.rn.f32 %f1193, %f1914, %f1848, %f1904; fma.rn.f32 %f1194, %f1914, %f1847, %f1903; fma.rn.f32 %f1195, %f1914, %f1846, %f1902; fma.rn.f32 %f1196, %f1914, %f1845, %f1901; fma.rn.f32 %f1197, %f1914, %f1844, %f1900; fma.rn.f32 %f1198, %f1914, %f1843, %f1899; fma.rn.f32 %f1199, %f1914, %f1842, %f1898; fma.rn.f32 %f1200, %f1914, %f1841, %f1897; fma.rn.f32 %f1201, %f1915, %f1848, %f1896; fma.rn.f32 %f1202, %f1915, %f1847, %f1895; fma.rn.f32 %f1203, %f1915, %f1846, %f1894; fma.rn.f32 %f1204, %f1915, %f1845, %f1893; fma.rn.f32 %f1205, %f1915, %f1844, %f1892; fma.rn.f32 %f1206, %f1915, %f1843, %f1891; fma.rn.f32 %f1207, %f1915, %f1842, %f1890; fma.rn.f32 %f1208, %f1915, %f1841, %f1889; fma.rn.f32 %f1209, %f1916, %f1848, %f1888; fma.rn.f32 %f1210, %f1916, %f1847, %f1887; fma.rn.f32 %f1211, %f1916, %f1846, %f1886; fma.rn.f32 %f1212, %f1916, %f1845, %f1885; fma.rn.f32 %f1213, %f1916, %f1844, %f1884; fma.rn.f32 %f1214, %f1916, %f1843, %f1883; fma.rn.f32 %f1215, %f1916, %f1842, %f1882; fma.rn.f32 %f1216, %f1916, %f1841, %f1881; fma.rn.f32 %f1217, %f1917, %f1848, %f1880; fma.rn.f32 %f1218, %f1917, %f1847, %f1879; fma.rn.f32 %f1219, %f1917, %f1846, %f1878; fma.rn.f32 %f1220, %f1917, %f1845, %f1877; fma.rn.f32 %f1221, %f1917, %f1844, %f1876; fma.rn.f32 %f1222, %f1917, %f1843, %f1875; fma.rn.f32 %f1223, %f1917, %f1842, %f1874; fma.rn.f32 %f1224, %f1917, %f1841, %f1873; fma.rn.f32 %f1225, %f1918, %f1848, %f1872; fma.rn.f32 %f1226, %f1918, %f1847, %f1871; fma.rn.f32 %f1227, %f1918, %f1846, %f1870; fma.rn.f32 %f1228, %f1918, %f1845, %f1869; fma.rn.f32 %f1229, %f1918, %f1844, %f1868; fma.rn.f32 %f1230, %f1918, %f1843, %f1867; fma.rn.f32 %f1231, %f1918, %f1842, %f1866; fma.rn.f32 %f1232, %f1918, %f1841, %f1865; fma.rn.f32 %f1233, %f1919, %f1848, %f1864; fma.rn.f32 %f1234, %f1919, %f1847, %f1863; fma.rn.f32 %f1235, %f1919, %f1846, %f1862; fma.rn.f32 %f1236, %f1919, %f1845, %f1861; fma.rn.f32 %f1237, %f1919, %f1844, %f1860; fma.rn.f32 %f1238, %f1919, %f1843, %f1859; fma.rn.f32 %f1239, %f1919, %f1842, %f1858; fma.rn.f32 %f1240, %f1919, %f1841, %f1857; fma.rn.f32 %f1241, %f1920, %f1848, %f1856; fma.rn.f32 %f1242, %f1920, %f1847, %f1855; fma.rn.f32 %f1243, %f1920, %f1846, %f1854; fma.rn.f32 %f1244, %f1920, %f1845, %f1853; fma.rn.f32 %f1245, %f1920, %f1844, %f1852; fma.rn.f32 %f1246, %f1920, %f1843, %f1851; fma.rn.f32 %f1247, %f1920, %f1842, %f1850; fma.rn.f32 %f1248, %f1920, %f1841, %f1849; add.s32 %r146, %r359, 1056; // begin inline asm ld.shared.v4.f32 {%f1089, %f1090, %f1091, %f1092}, [%r146]; // end inline asm add.s32 %r147, %r359, 1120; // begin inline asm ld.shared.v4.f32 {%f1093, %f1094, %f1095, %f1096}, [%r147]; // end inline asm add.s32 %r148, %r358, 1024; // begin inline asm ld.shared.v4.f32 {%f1097, %f1098, %f1099, %f1100}, [%r148]; // end inline asm add.s32 %r149, %r358, 1152; // begin inline asm ld.shared.v4.f32 {%f1101, %f1102, %f1103, %f1104}, [%r149]; // end inline asm fma.rn.f32 %f1249, %f1073, %f1081, %f1185; fma.rn.f32 %f1250, %f1073, %f1082, %f1186; fma.rn.f32 %f1251, %f1073, %f1083, %f1187; fma.rn.f32 %f1252, %f1073, %f1084, %f1188; fma.rn.f32 %f1253, %f1073, %f1085, %f1189; fma.rn.f32 %f1254, %f1073, %f1086, %f1190; fma.rn.f32 %f1255, %f1073, %f1087, %f1191; fma.rn.f32 %f1256, %f1073, %f1088, %f1192; fma.rn.f32 %f1257, %f1074, %f1081, %f1193; fma.rn.f32 %f1258, %f1074, %f1082, %f1194; fma.rn.f32 %f1259, %f1074, %f1083, %f1195; fma.rn.f32 %f1260, %f1074, %f1084, %f1196; fma.rn.f32 %f1261, %f1074, %f1085, %f1197; fma.rn.f32 %f1262, %f1074, %f1086, %f1198; fma.rn.f32 %f1263, %f1074, %f1087, %f1199; fma.rn.f32 %f1264, %f1074, %f1088, %f1200; fma.rn.f32 %f1265, %f1075, %f1081, %f1201; fma.rn.f32 %f1266, %f1075, %f1082, %f1202; fma.rn.f32 %f1267, %f1075, %f1083, %f1203; fma.rn.f32 %f1268, %f1075, %f1084, %f1204; fma.rn.f32 %f1269, %f1075, %f1085, %f1205; fma.rn.f32 %f1270, %f1075, %f1086, %f1206; fma.rn.f32 %f1271, %f1075, %f1087, %f1207; fma.rn.f32 %f1272, %f1075, %f1088, %f1208; fma.rn.f32 %f1273, %f1076, %f1081, %f1209; fma.rn.f32 %f1274, %f1076, %f1082, %f1210; fma.rn.f32 %f1275, %f1076, %f1083, %f1211; fma.rn.f32 %f1276, %f1076, %f1084, %f1212; fma.rn.f32 %f1277, %f1076, %f1085, %f1213; fma.rn.f32 %f1278, %f1076, %f1086, %f1214; fma.rn.f32 %f1279, %f1076, %f1087, %f1215; fma.rn.f32 %f1280, %f1076, %f1088, %f1216; fma.rn.f32 %f1281, %f1077, %f1081, %f1217; fma.rn.f32 %f1282, %f1077, %f1082, %f1218; fma.rn.f32 %f1283, %f1077, %f1083, %f1219; fma.rn.f32 %f1284, %f1077, %f1084, %f1220; fma.rn.f32 %f1285, %f1077, %f1085, %f1221; fma.rn.f32 %f1286, %f1077, %f1086, %f1222; fma.rn.f32 %f1287, %f1077, %f1087, %f1223; fma.rn.f32 %f1288, %f1077, %f1088, %f1224; fma.rn.f32 %f1289, %f1078, %f1081, %f1225; fma.rn.f32 %f1290, %f1078, %f1082, %f1226; fma.rn.f32 %f1291, %f1078, %f1083, %f1227; fma.rn.f32 %f1292, %f1078, %f1084, %f1228; fma.rn.f32 %f1293, %f1078, %f1085, %f1229; fma.rn.f32 %f1294, %f1078, %f1086, %f1230; fma.rn.f32 %f1295, %f1078, %f1087, %f1231; fma.rn.f32 %f1296, %f1078, %f1088, %f1232; fma.rn.f32 %f1297, %f1079, %f1081, %f1233; fma.rn.f32 %f1298, %f1079, %f1082, %f1234; fma.rn.f32 %f1299, %f1079, %f1083, %f1235; fma.rn.f32 %f1300, %f1079, %f1084, %f1236; fma.rn.f32 %f1301, %f1079, %f1085, %f1237; fma.rn.f32 %f1302, %f1079, %f1086, %f1238; fma.rn.f32 %f1303, %f1079, %f1087, %f1239; fma.rn.f32 %f1304, %f1079, %f1088, %f1240; fma.rn.f32 %f1305, %f1080, %f1081, %f1241; fma.rn.f32 %f1306, %f1080, %f1082, %f1242; fma.rn.f32 %f1307, %f1080, %f1083, %f1243; fma.rn.f32 %f1308, %f1080, %f1084, %f1244; fma.rn.f32 %f1309, %f1080, %f1085, %f1245; fma.rn.f32 %f1310, %f1080, %f1086, %f1246; fma.rn.f32 %f1311, %f1080, %f1087, %f1247; fma.rn.f32 %f1312, %f1080, %f1088, %f1248; add.s32 %r150, %r359, 1584; // begin inline asm ld.shared.v4.f32 {%f1105, %f1106, %f1107, %f1108}, [%r150]; // end inline asm add.s32 %r151, %r359, 1648; // begin inline asm ld.shared.v4.f32 {%f1109, %f1110, %f1111, %f1112}, [%r151]; // end inline asm add.s32 %r152, %r358, 1536; // begin inline asm ld.shared.v4.f32 {%f1113, %f1114, %f1115, %f1116}, [%r152]; // end inline asm add.s32 %r153, %r358, 1664; // begin inline asm ld.shared.v4.f32 {%f1117, %f1118, %f1119, %f1120}, [%r153]; // end inline asm fma.rn.f32 %f1313, %f1089, %f1097, %f1249; fma.rn.f32 %f1314, %f1089, %f1098, %f1250; fma.rn.f32 %f1315, %f1089, %f1099, %f1251; fma.rn.f32 %f1316, %f1089, %f1100, %f1252; fma.rn.f32 %f1317, %f1089, %f1101, %f1253; fma.rn.f32 %f1318, %f1089, %f1102, %f1254; fma.rn.f32 %f1319, %f1089, %f1103, %f1255; fma.rn.f32 %f1320, %f1089, %f1104, %f1256; fma.rn.f32 %f1321, %f1090, %f1097, %f1257; fma.rn.f32 %f1322, %f1090, %f1098, %f1258; fma.rn.f32 %f1323, %f1090, %f1099, %f1259; fma.rn.f32 %f1324, %f1090, %f1100, %f1260; fma.rn.f32 %f1325, %f1090, %f1101, %f1261; fma.rn.f32 %f1326, %f1090, %f1102, %f1262; fma.rn.f32 %f1327, %f1090, %f1103, %f1263; fma.rn.f32 %f1328, %f1090, %f1104, %f1264; fma.rn.f32 %f1329, %f1091, %f1097, %f1265; fma.rn.f32 %f1330, %f1091, %f1098, %f1266; fma.rn.f32 %f1331, %f1091, %f1099, %f1267; fma.rn.f32 %f1332, %f1091, %f1100, %f1268; fma.rn.f32 %f1333, %f1091, %f1101, %f1269; fma.rn.f32 %f1334, %f1091, %f1102, %f1270; fma.rn.f32 %f1335, %f1091, %f1103, %f1271; fma.rn.f32 %f1336, %f1091, %f1104, %f1272; fma.rn.f32 %f1337, %f1092, %f1097, %f1273; fma.rn.f32 %f1338, %f1092, %f1098, %f1274; fma.rn.f32 %f1339, %f1092, %f1099, %f1275; fma.rn.f32 %f1340, %f1092, %f1100, %f1276; fma.rn.f32 %f1341, %f1092, %f1101, %f1277; fma.rn.f32 %f1342, %f1092, %f1102, %f1278; fma.rn.f32 %f1343, %f1092, %f1103, %f1279; fma.rn.f32 %f1344, %f1092, %f1104, %f1280; fma.rn.f32 %f1345, %f1093, %f1097, %f1281; fma.rn.f32 %f1346, %f1093, %f1098, %f1282; fma.rn.f32 %f1347, %f1093, %f1099, %f1283; fma.rn.f32 %f1348, %f1093, %f1100, %f1284; fma.rn.f32 %f1349, %f1093, %f1101, %f1285; fma.rn.f32 %f1350, %f1093, %f1102, %f1286; fma.rn.f32 %f1351, %f1093, %f1103, %f1287; fma.rn.f32 %f1352, %f1093, %f1104, %f1288; fma.rn.f32 %f1353, %f1094, %f1097, %f1289; fma.rn.f32 %f1354, %f1094, %f1098, %f1290; fma.rn.f32 %f1355, %f1094, %f1099, %f1291; fma.rn.f32 %f1356, %f1094, %f1100, %f1292; fma.rn.f32 %f1357, %f1094, %f1101, %f1293; fma.rn.f32 %f1358, %f1094, %f1102, %f1294; fma.rn.f32 %f1359, %f1094, %f1103, %f1295; fma.rn.f32 %f1360, %f1094, %f1104, %f1296; fma.rn.f32 %f1361, %f1095, %f1097, %f1297; fma.rn.f32 %f1362, %f1095, %f1098, %f1298; fma.rn.f32 %f1363, %f1095, %f1099, %f1299; fma.rn.f32 %f1364, %f1095, %f1100, %f1300; fma.rn.f32 %f1365, %f1095, %f1101, %f1301; fma.rn.f32 %f1366, %f1095, %f1102, %f1302; fma.rn.f32 %f1367, %f1095, %f1103, %f1303; fma.rn.f32 %f1368, %f1095, %f1104, %f1304; fma.rn.f32 %f1369, %f1096, %f1097, %f1305; fma.rn.f32 %f1370, %f1096, %f1098, %f1306; fma.rn.f32 %f1371, %f1096, %f1099, %f1307; fma.rn.f32 %f1372, %f1096, %f1100, %f1308; fma.rn.f32 %f1373, %f1096, %f1101, %f1309; fma.rn.f32 %f1374, %f1096, %f1102, %f1310; fma.rn.f32 %f1375, %f1096, %f1103, %f1311; fma.rn.f32 %f1376, %f1096, %f1104, %f1312; add.s32 %r154, %r359, 2112; // begin inline asm ld.shared.v4.f32 {%f1121, %f1122, %f1123, %f1124}, [%r154]; // end inline asm add.s32 %r155, %r359, 2176; // begin inline asm ld.shared.v4.f32 {%f1125, %f1126, %f1127, %f1128}, [%r155]; // end inline asm add.s32 %r156, %r358, 2048; // begin inline asm ld.shared.v4.f32 {%f1129, %f1130, %f1131, %f1132}, [%r156]; // end inline asm add.s32 %r157, %r358, 2176; // begin inline asm ld.shared.v4.f32 {%f1133, %f1134, %f1135, %f1136}, [%r157]; // end inline asm fma.rn.f32 %f1377, %f1105, %f1113, %f1313; fma.rn.f32 %f1378, %f1105, %f1114, %f1314; fma.rn.f32 %f1379, %f1105, %f1115, %f1315; fma.rn.f32 %f1380, %f1105, %f1116, %f1316; fma.rn.f32 %f1381, %f1105, %f1117, %f1317; fma.rn.f32 %f1382, %f1105, %f1118, %f1318; fma.rn.f32 %f1383, %f1105, %f1119, %f1319; fma.rn.f32 %f1384, %f1105, %f1120, %f1320; fma.rn.f32 %f1385, %f1106, %f1113, %f1321; fma.rn.f32 %f1386, %f1106, %f1114, %f1322; fma.rn.f32 %f1387, %f1106, %f1115, %f1323; fma.rn.f32 %f1388, %f1106, %f1116, %f1324; fma.rn.f32 %f1389, %f1106, %f1117, %f1325; fma.rn.f32 %f1390, %f1106, %f1118, %f1326; fma.rn.f32 %f1391, %f1106, %f1119, %f1327; fma.rn.f32 %f1392, %f1106, %f1120, %f1328; fma.rn.f32 %f1393, %f1107, %f1113, %f1329; fma.rn.f32 %f1394, %f1107, %f1114, %f1330; fma.rn.f32 %f1395, %f1107, %f1115, %f1331; fma.rn.f32 %f1396, %f1107, %f1116, %f1332; fma.rn.f32 %f1397, %f1107, %f1117, %f1333; fma.rn.f32 %f1398, %f1107, %f1118, %f1334; fma.rn.f32 %f1399, %f1107, %f1119, %f1335; fma.rn.f32 %f1400, %f1107, %f1120, %f1336; fma.rn.f32 %f1401, %f1108, %f1113, %f1337; fma.rn.f32 %f1402, %f1108, %f1114, %f1338; fma.rn.f32 %f1403, %f1108, %f1115, %f1339; fma.rn.f32 %f1404, %f1108, %f1116, %f1340; fma.rn.f32 %f1405, %f1108, %f1117, %f1341; fma.rn.f32 %f1406, %f1108, %f1118, %f1342; fma.rn.f32 %f1407, %f1108, %f1119, %f1343; fma.rn.f32 %f1408, %f1108, %f1120, %f1344; fma.rn.f32 %f1409, %f1109, %f1113, %f1345; fma.rn.f32 %f1410, %f1109, %f1114, %f1346; fma.rn.f32 %f1411, %f1109, %f1115, %f1347; fma.rn.f32 %f1412, %f1109, %f1116, %f1348; fma.rn.f32 %f1413, %f1109, %f1117, %f1349; fma.rn.f32 %f1414, %f1109, %f1118, %f1350; fma.rn.f32 %f1415, %f1109, %f1119, %f1351; fma.rn.f32 %f1416, %f1109, %f1120, %f1352; fma.rn.f32 %f1417, %f1110, %f1113, %f1353; fma.rn.f32 %f1418, %f1110, %f1114, %f1354; fma.rn.f32 %f1419, %f1110, %f1115, %f1355; fma.rn.f32 %f1420, %f1110, %f1116, %f1356; fma.rn.f32 %f1421, %f1110, %f1117, %f1357; fma.rn.f32 %f1422, %f1110, %f1118, %f1358; fma.rn.f32 %f1423, %f1110, %f1119, %f1359; fma.rn.f32 %f1424, %f1110, %f1120, %f1360; fma.rn.f32 %f1425, %f1111, %f1113, %f1361; fma.rn.f32 %f1426, %f1111, %f1114, %f1362; fma.rn.f32 %f1427, %f1111, %f1115, %f1363; fma.rn.f32 %f1428, %f1111, %f1116, %f1364; fma.rn.f32 %f1429, %f1111, %f1117, %f1365; fma.rn.f32 %f1430, %f1111, %f1118, %f1366; fma.rn.f32 %f1431, %f1111, %f1119, %f1367; fma.rn.f32 %f1432, %f1111, %f1120, %f1368; fma.rn.f32 %f1433, %f1112, %f1113, %f1369; fma.rn.f32 %f1434, %f1112, %f1114, %f1370; fma.rn.f32 %f1435, %f1112, %f1115, %f1371; fma.rn.f32 %f1436, %f1112, %f1116, %f1372; fma.rn.f32 %f1437, %f1112, %f1117, %f1373; fma.rn.f32 %f1438, %f1112, %f1118, %f1374; fma.rn.f32 %f1439, %f1112, %f1119, %f1375; fma.rn.f32 %f1440, %f1112, %f1120, %f1376; add.s32 %r158, %r359, 2640; // begin inline asm ld.shared.v4.f32 {%f1137, %f1138, %f1139, %f1140}, [%r158]; // end inline asm add.s32 %r159, %r359, 2704; // begin inline asm ld.shared.v4.f32 {%f1141, %f1142, %f1143, %f1144}, [%r159]; // end inline asm add.s32 %r160, %r358, 2560; // begin inline asm ld.shared.v4.f32 {%f1145, %f1146, %f1147, %f1148}, [%r160]; // end inline asm add.s32 %r161, %r358, 2688; // begin inline asm ld.shared.v4.f32 {%f1149, %f1150, %f1151, %f1152}, [%r161]; // end inline asm fma.rn.f32 %f1441, %f1121, %f1129, %f1377; fma.rn.f32 %f1442, %f1121, %f1130, %f1378; fma.rn.f32 %f1443, %f1121, %f1131, %f1379; fma.rn.f32 %f1444, %f1121, %f1132, %f1380; fma.rn.f32 %f1445, %f1121, %f1133, %f1381; fma.rn.f32 %f1446, %f1121, %f1134, %f1382; fma.rn.f32 %f1447, %f1121, %f1135, %f1383; fma.rn.f32 %f1448, %f1121, %f1136, %f1384; fma.rn.f32 %f1449, %f1122, %f1129, %f1385; fma.rn.f32 %f1450, %f1122, %f1130, %f1386; fma.rn.f32 %f1451, %f1122, %f1131, %f1387; fma.rn.f32 %f1452, %f1122, %f1132, %f1388; fma.rn.f32 %f1453, %f1122, %f1133, %f1389; fma.rn.f32 %f1454, %f1122, %f1134, %f1390; fma.rn.f32 %f1455, %f1122, %f1135, %f1391; fma.rn.f32 %f1456, %f1122, %f1136, %f1392; fma.rn.f32 %f1457, %f1123, %f1129, %f1393; fma.rn.f32 %f1458, %f1123, %f1130, %f1394; fma.rn.f32 %f1459, %f1123, %f1131, %f1395; fma.rn.f32 %f1460, %f1123, %f1132, %f1396; fma.rn.f32 %f1461, %f1123, %f1133, %f1397; fma.rn.f32 %f1462, %f1123, %f1134, %f1398; fma.rn.f32 %f1463, %f1123, %f1135, %f1399; fma.rn.f32 %f1464, %f1123, %f1136, %f1400; fma.rn.f32 %f1465, %f1124, %f1129, %f1401; fma.rn.f32 %f1466, %f1124, %f1130, %f1402; fma.rn.f32 %f1467, %f1124, %f1131, %f1403; fma.rn.f32 %f1468, %f1124, %f1132, %f1404; fma.rn.f32 %f1469, %f1124, %f1133, %f1405; fma.rn.f32 %f1470, %f1124, %f1134, %f1406; fma.rn.f32 %f1471, %f1124, %f1135, %f1407; fma.rn.f32 %f1472, %f1124, %f1136, %f1408; fma.rn.f32 %f1473, %f1125, %f1129, %f1409; fma.rn.f32 %f1474, %f1125, %f1130, %f1410; fma.rn.f32 %f1475, %f1125, %f1131, %f1411; fma.rn.f32 %f1476, %f1125, %f1132, %f1412; fma.rn.f32 %f1477, %f1125, %f1133, %f1413; fma.rn.f32 %f1478, %f1125, %f1134, %f1414; fma.rn.f32 %f1479, %f1125, %f1135, %f1415; fma.rn.f32 %f1480, %f1125, %f1136, %f1416; fma.rn.f32 %f1481, %f1126, %f1129, %f1417; fma.rn.f32 %f1482, %f1126, %f1130, %f1418; fma.rn.f32 %f1483, %f1126, %f1131, %f1419; fma.rn.f32 %f1484, %f1126, %f1132, %f1420; fma.rn.f32 %f1485, %f1126, %f1133, %f1421; fma.rn.f32 %f1486, %f1126, %f1134, %f1422; fma.rn.f32 %f1487, %f1126, %f1135, %f1423; fma.rn.f32 %f1488, %f1126, %f1136, %f1424; fma.rn.f32 %f1489, %f1127, %f1129, %f1425; fma.rn.f32 %f1490, %f1127, %f1130, %f1426; fma.rn.f32 %f1491, %f1127, %f1131, %f1427; fma.rn.f32 %f1492, %f1127, %f1132, %f1428; fma.rn.f32 %f1493, %f1127, %f1133, %f1429; fma.rn.f32 %f1494, %f1127, %f1134, %f1430; fma.rn.f32 %f1495, %f1127, %f1135, %f1431; fma.rn.f32 %f1496, %f1127, %f1136, %f1432; fma.rn.f32 %f1497, %f1128, %f1129, %f1433; fma.rn.f32 %f1498, %f1128, %f1130, %f1434; fma.rn.f32 %f1499, %f1128, %f1131, %f1435; fma.rn.f32 %f1500, %f1128, %f1132, %f1436; fma.rn.f32 %f1501, %f1128, %f1133, %f1437; fma.rn.f32 %f1502, %f1128, %f1134, %f1438; fma.rn.f32 %f1503, %f1128, %f1135, %f1439; fma.rn.f32 %f1504, %f1128, %f1136, %f1440; add.s32 %r162, %r359, 3168; // begin inline asm ld.shared.v4.f32 {%f1153, %f1154, %f1155, %f1156}, [%r162]; // end inline asm add.s32 %r163, %r359, 3232; // begin inline asm ld.shared.v4.f32 {%f1157, %f1158, %f1159, %f1160}, [%r163]; // end inline asm add.s32 %r164, %r358, 3072; // begin inline asm ld.shared.v4.f32 {%f1161, %f1162, %f1163, %f1164}, [%r164]; // end inline asm add.s32 %r165, %r358, 3200; // begin inline asm ld.shared.v4.f32 {%f1165, %f1166, %f1167, %f1168}, [%r165]; // end inline asm fma.rn.f32 %f1505, %f1137, %f1145, %f1441; fma.rn.f32 %f1506, %f1137, %f1146, %f1442; fma.rn.f32 %f1507, %f1137, %f1147, %f1443; fma.rn.f32 %f1508, %f1137, %f1148, %f1444; fma.rn.f32 %f1509, %f1137, %f1149, %f1445; fma.rn.f32 %f1510, %f1137, %f1150, %f1446; fma.rn.f32 %f1511, %f1137, %f1151, %f1447; fma.rn.f32 %f1512, %f1137, %f1152, %f1448; fma.rn.f32 %f1513, %f1138, %f1145, %f1449; fma.rn.f32 %f1514, %f1138, %f1146, %f1450; fma.rn.f32 %f1515, %f1138, %f1147, %f1451; fma.rn.f32 %f1516, %f1138, %f1148, %f1452; fma.rn.f32 %f1517, %f1138, %f1149, %f1453; fma.rn.f32 %f1518, %f1138, %f1150, %f1454; fma.rn.f32 %f1519, %f1138, %f1151, %f1455; fma.rn.f32 %f1520, %f1138, %f1152, %f1456; fma.rn.f32 %f1521, %f1139, %f1145, %f1457; fma.rn.f32 %f1522, %f1139, %f1146, %f1458; fma.rn.f32 %f1523, %f1139, %f1147, %f1459; fma.rn.f32 %f1524, %f1139, %f1148, %f1460; fma.rn.f32 %f1525, %f1139, %f1149, %f1461; fma.rn.f32 %f1526, %f1139, %f1150, %f1462; fma.rn.f32 %f1527, %f1139, %f1151, %f1463; fma.rn.f32 %f1528, %f1139, %f1152, %f1464; fma.rn.f32 %f1529, %f1140, %f1145, %f1465; fma.rn.f32 %f1530, %f1140, %f1146, %f1466; fma.rn.f32 %f1531, %f1140, %f1147, %f1467; fma.rn.f32 %f1532, %f1140, %f1148, %f1468; fma.rn.f32 %f1533, %f1140, %f1149, %f1469; fma.rn.f32 %f1534, %f1140, %f1150, %f1470; fma.rn.f32 %f1535, %f1140, %f1151, %f1471; fma.rn.f32 %f1536, %f1140, %f1152, %f1472; fma.rn.f32 %f1537, %f1141, %f1145, %f1473; fma.rn.f32 %f1538, %f1141, %f1146, %f1474; fma.rn.f32 %f1539, %f1141, %f1147, %f1475; fma.rn.f32 %f1540, %f1141, %f1148, %f1476; fma.rn.f32 %f1541, %f1141, %f1149, %f1477; fma.rn.f32 %f1542, %f1141, %f1150, %f1478; fma.rn.f32 %f1543, %f1141, %f1151, %f1479; fma.rn.f32 %f1544, %f1141, %f1152, %f1480; fma.rn.f32 %f1545, %f1142, %f1145, %f1481; fma.rn.f32 %f1546, %f1142, %f1146, %f1482; fma.rn.f32 %f1547, %f1142, %f1147, %f1483; fma.rn.f32 %f1548, %f1142, %f1148, %f1484; fma.rn.f32 %f1549, %f1142, %f1149, %f1485; fma.rn.f32 %f1550, %f1142, %f1150, %f1486; fma.rn.f32 %f1551, %f1142, %f1151, %f1487; fma.rn.f32 %f1552, %f1142, %f1152, %f1488; fma.rn.f32 %f1553, %f1143, %f1145, %f1489; fma.rn.f32 %f1554, %f1143, %f1146, %f1490; fma.rn.f32 %f1555, %f1143, %f1147, %f1491; fma.rn.f32 %f1556, %f1143, %f1148, %f1492; fma.rn.f32 %f1557, %f1143, %f1149, %f1493; fma.rn.f32 %f1558, %f1143, %f1150, %f1494; fma.rn.f32 %f1559, %f1143, %f1151, %f1495; fma.rn.f32 %f1560, %f1143, %f1152, %f1496; fma.rn.f32 %f1561, %f1144, %f1145, %f1497; fma.rn.f32 %f1562, %f1144, %f1146, %f1498; fma.rn.f32 %f1563, %f1144, %f1147, %f1499; fma.rn.f32 %f1564, %f1144, %f1148, %f1500; fma.rn.f32 %f1565, %f1144, %f1149, %f1501; fma.rn.f32 %f1566, %f1144, %f1150, %f1502; fma.rn.f32 %f1567, %f1144, %f1151, %f1503; fma.rn.f32 %f1568, %f1144, %f1152, %f1504; add.s32 %r166, %r359, 3696; // begin inline asm ld.shared.v4.f32 {%f1169, %f1170, %f1171, %f1172}, [%r166]; // end inline asm add.s32 %r167, %r359, 3760; // begin inline asm ld.shared.v4.f32 {%f1173, %f1174, %f1175, %f1176}, [%r167]; // end inline asm add.s32 %r168, %r358, 3584; // begin inline asm ld.shared.v4.f32 {%f1177, %f1178, %f1179, %f1180}, [%r168]; // end inline asm add.s32 %r169, %r358, 3712; // begin inline asm ld.shared.v4.f32 {%f1181, %f1182, %f1183, %f1184}, [%r169]; // end inline asm fma.rn.f32 %f1569, %f1153, %f1161, %f1505; fma.rn.f32 %f1570, %f1153, %f1162, %f1506; fma.rn.f32 %f1571, %f1153, %f1163, %f1507; fma.rn.f32 %f1572, %f1153, %f1164, %f1508; fma.rn.f32 %f1573, %f1153, %f1165, %f1509; fma.rn.f32 %f1574, %f1153, %f1166, %f1510; fma.rn.f32 %f1575, %f1153, %f1167, %f1511; fma.rn.f32 %f1576, %f1153, %f1168, %f1512; fma.rn.f32 %f1577, %f1154, %f1161, %f1513; fma.rn.f32 %f1578, %f1154, %f1162, %f1514; fma.rn.f32 %f1579, %f1154, %f1163, %f1515; fma.rn.f32 %f1580, %f1154, %f1164, %f1516; fma.rn.f32 %f1581, %f1154, %f1165, %f1517; fma.rn.f32 %f1582, %f1154, %f1166, %f1518; fma.rn.f32 %f1583, %f1154, %f1167, %f1519; fma.rn.f32 %f1584, %f1154, %f1168, %f1520; fma.rn.f32 %f1585, %f1155, %f1161, %f1521; fma.rn.f32 %f1586, %f1155, %f1162, %f1522; fma.rn.f32 %f1587, %f1155, %f1163, %f1523; fma.rn.f32 %f1588, %f1155, %f1164, %f1524; fma.rn.f32 %f1589, %f1155, %f1165, %f1525; fma.rn.f32 %f1590, %f1155, %f1166, %f1526; fma.rn.f32 %f1591, %f1155, %f1167, %f1527; fma.rn.f32 %f1592, %f1155, %f1168, %f1528; fma.rn.f32 %f1593, %f1156, %f1161, %f1529; fma.rn.f32 %f1594, %f1156, %f1162, %f1530; fma.rn.f32 %f1595, %f1156, %f1163, %f1531; fma.rn.f32 %f1596, %f1156, %f1164, %f1532; fma.rn.f32 %f1597, %f1156, %f1165, %f1533; fma.rn.f32 %f1598, %f1156, %f1166, %f1534; fma.rn.f32 %f1599, %f1156, %f1167, %f1535; fma.rn.f32 %f1600, %f1156, %f1168, %f1536; fma.rn.f32 %f1601, %f1157, %f1161, %f1537; fma.rn.f32 %f1602, %f1157, %f1162, %f1538; fma.rn.f32 %f1603, %f1157, %f1163, %f1539; fma.rn.f32 %f1604, %f1157, %f1164, %f1540; fma.rn.f32 %f1605, %f1157, %f1165, %f1541; fma.rn.f32 %f1606, %f1157, %f1166, %f1542; fma.rn.f32 %f1607, %f1157, %f1167, %f1543; fma.rn.f32 %f1608, %f1157, %f1168, %f1544; fma.rn.f32 %f1609, %f1158, %f1161, %f1545; fma.rn.f32 %f1610, %f1158, %f1162, %f1546; fma.rn.f32 %f1611, %f1158, %f1163, %f1547; fma.rn.f32 %f1612, %f1158, %f1164, %f1548; fma.rn.f32 %f1613, %f1158, %f1165, %f1549; fma.rn.f32 %f1614, %f1158, %f1166, %f1550; fma.rn.f32 %f1615, %f1158, %f1167, %f1551; fma.rn.f32 %f1616, %f1158, %f1168, %f1552; fma.rn.f32 %f1617, %f1159, %f1161, %f1553; fma.rn.f32 %f1618, %f1159, %f1162, %f1554; fma.rn.f32 %f1619, %f1159, %f1163, %f1555; fma.rn.f32 %f1620, %f1159, %f1164, %f1556; fma.rn.f32 %f1621, %f1159, %f1165, %f1557; fma.rn.f32 %f1622, %f1159, %f1166, %f1558; fma.rn.f32 %f1623, %f1159, %f1167, %f1559; fma.rn.f32 %f1624, %f1159, %f1168, %f1560; fma.rn.f32 %f1625, %f1160, %f1161, %f1561; fma.rn.f32 %f1626, %f1160, %f1162, %f1562; fma.rn.f32 %f1627, %f1160, %f1163, %f1563; fma.rn.f32 %f1628, %f1160, %f1164, %f1564; fma.rn.f32 %f1629, %f1160, %f1165, %f1565; fma.rn.f32 %f1630, %f1160, %f1166, %f1566; fma.rn.f32 %f1631, %f1160, %f1167, %f1567; fma.rn.f32 %f1632, %f1160, %f1168, %f1568; fma.rn.f32 %f257, %f1169, %f1177, %f1569; fma.rn.f32 %f258, %f1169, %f1178, %f1570; fma.rn.f32 %f259, %f1169, %f1179, %f1571; fma.rn.f32 %f260, %f1169, %f1180, %f1572; fma.rn.f32 %f261, %f1169, %f1181, %f1573; fma.rn.f32 %f262, %f1169, %f1182, %f1574; fma.rn.f32 %f263, %f1169, %f1183, %f1575; fma.rn.f32 %f264, %f1169, %f1184, %f1576; fma.rn.f32 %f265, %f1170, %f1177, %f1577; fma.rn.f32 %f266, %f1170, %f1178, %f1578; fma.rn.f32 %f267, %f1170, %f1179, %f1579; fma.rn.f32 %f268, %f1170, %f1180, %f1580; fma.rn.f32 %f269, %f1170, %f1181, %f1581; fma.rn.f32 %f270, %f1170, %f1182, %f1582; fma.rn.f32 %f271, %f1170, %f1183, %f1583; fma.rn.f32 %f272, %f1170, %f1184, %f1584; fma.rn.f32 %f273, %f1171, %f1177, %f1585; fma.rn.f32 %f274, %f1171, %f1178, %f1586; fma.rn.f32 %f275, %f1171, %f1179, %f1587; fma.rn.f32 %f276, %f1171, %f1180, %f1588; fma.rn.f32 %f277, %f1171, %f1181, %f1589; fma.rn.f32 %f278, %f1171, %f1182, %f1590; fma.rn.f32 %f279, %f1171, %f1183, %f1591; fma.rn.f32 %f280, %f1171, %f1184, %f1592; fma.rn.f32 %f281, %f1172, %f1177, %f1593; fma.rn.f32 %f282, %f1172, %f1178, %f1594; fma.rn.f32 %f283, %f1172, %f1179, %f1595; fma.rn.f32 %f284, %f1172, %f1180, %f1596; fma.rn.f32 %f285, %f1172, %f1181, %f1597; fma.rn.f32 %f286, %f1172, %f1182, %f1598; fma.rn.f32 %f287, %f1172, %f1183, %f1599; fma.rn.f32 %f288, %f1172, %f1184, %f1600; fma.rn.f32 %f289, %f1173, %f1177, %f1601; fma.rn.f32 %f290, %f1173, %f1178, %f1602; fma.rn.f32 %f291, %f1173, %f1179, %f1603; fma.rn.f32 %f292, %f1173, %f1180, %f1604; fma.rn.f32 %f293, %f1173, %f1181, %f1605; fma.rn.f32 %f294, %f1173, %f1182, %f1606; fma.rn.f32 %f295, %f1173, %f1183, %f1607; fma.rn.f32 %f296, %f1173, %f1184, %f1608; fma.rn.f32 %f297, %f1174, %f1177, %f1609; fma.rn.f32 %f298, %f1174, %f1178, %f1610; fma.rn.f32 %f299, %f1174, %f1179, %f1611; fma.rn.f32 %f300, %f1174, %f1180, %f1612; fma.rn.f32 %f301, %f1174, %f1181, %f1613; fma.rn.f32 %f302, %f1174, %f1182, %f1614; fma.rn.f32 %f303, %f1174, %f1183, %f1615; fma.rn.f32 %f304, %f1174, %f1184, %f1616; fma.rn.f32 %f305, %f1175, %f1177, %f1617; fma.rn.f32 %f306, %f1175, %f1178, %f1618; fma.rn.f32 %f307, %f1175, %f1179, %f1619; fma.rn.f32 %f308, %f1175, %f1180, %f1620; fma.rn.f32 %f309, %f1175, %f1181, %f1621; fma.rn.f32 %f310, %f1175, %f1182, %f1622; fma.rn.f32 %f311, %f1175, %f1183, %f1623; fma.rn.f32 %f312, %f1175, %f1184, %f1624; fma.rn.f32 %f313, %f1176, %f1177, %f1625; fma.rn.f32 %f314, %f1176, %f1178, %f1626; fma.rn.f32 %f315, %f1176, %f1179, %f1627; fma.rn.f32 %f316, %f1176, %f1180, %f1628; fma.rn.f32 %f317, %f1176, %f1181, %f1629; fma.rn.f32 %f318, %f1176, %f1182, %f1630; fma.rn.f32 %f319, %f1176, %f1183, %f1631; fma.rn.f32 %f320, %f1176, %f1184, %f1632; shl.b32 %r171, %r1, 11; cvt.u64.u32 %rd53, %r171; add.s64 %rd55, %rd34, %rd53; shl.b32 %r173, %r3, 5; or.b32 %r174, %r173, %r2; mul.wide.u32 %rd56, %r174, 16; add.s64 %rd52, %rd55, %rd56; // begin inline asm {.reg .u64 u64addr; cvta.to.shared.u64 u64addr, %rd52; cvt.u32.u64 %r170, u64addr;} // end inline asm add.s32 %r175, %r8, %r4; setp.ge.u32 %p29, %r175, %r34; @%p29 bra $L__BB1_7; ld.param.u32 %r335, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_3]; add.s32 %r182, %r175, 32; setp.gt.u32 %p30, %r182, %r335; @%p30 bra $L__BB1_6; bra.uni $L__BB1_5; $L__BB1_6: mov.u32 %r352, %tid.x; and.b32 %r351, %r352, 31; mov.u32 %r350, %ctaid.x; bfi.b32 %r349, %r350, %r351, 7, 25; mov.u32 %r344, %tid.x; ld.param.u64 %rd173, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_2]; shl.b32 %r343, %r344, 1; mov.u32 %r342, _ZZ22sgemm_128x128x8_kernelPKfS0_PfjjjjjE4smem; { .reg .b64 %tmp; cvt.u64.u32 %tmp, %r342; cvta.shared.u64 %rd172, %tmp; } ld.param.u32 %r341, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_3]; ld.param.u32 %r340, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_4]; mul.lo.s32 %r321, %r175, %r340; cvt.u64.u32 %rd156, %r321; and.b32 %r326, %r343, 64; or.b32 %r327, %r349, %r326; cvt.u64.u32 %rd157, %r327; add.s64 %rd158, %rd156, %rd157; shl.b64 %rd159, %rd158, 2; add.s64 %rd160, %rd173, %rd159; shl.b32 %r328, %r344, 6; and.b32 %r329, %r328, -2048; cvt.u64.u32 %rd161, %r329; add.s64 %rd163, %rd172, %rd161; shl.b32 %r331, %r344, 2; cvt.u64.u32 %rd164, %r331; and.b64 %rd165, %rd164, 124; add.s64 %rd166, %rd163, %rd165; { // callseq 0, 0 .reg .b32 temp_param_reg; .param .align 4 .b8 param0[64]; st.param.f32 [param0+0], %f257; st.param.f32 [param0+4], %f258; st.param.f32 [param0+8], %f259; st.param.f32 [param0+12], %f260; st.param.f32 [param0+16], %f265; st.param.f32 [param0+20], %f266; st.param.f32 [param0+24], %f267; st.param.f32 [param0+28], %f268; st.param.f32 [param0+32], %f273; st.param.f32 [param0+36], %f274; st.param.f32 [param0+40], %f275; st.param.f32 [param0+44], %f276; st.param.f32 [param0+48], %f281; st.param.f32 [param0+52], %f282; st.param.f32 [param0+56], %f283; st.param.f32 [param0+60], %f284; .param .b64 param1; st.param.b64 [param1+0], %rd160; .param .b64 param2; st.param.b64 [param2+0], %rd166; .param .b32 param3; st.param.b32 [param3+0], %r170; .param .b32 param4; st.param.b32 [param4+0], %r341; .param .b32 param5; st.param.b32 [param5+0], %r340; .param .b32 param6; st.param.b32 [param6+0], %r175; .param .b32 param7; st.param.b32 [param7+0], %r327; call.uni _Z9C_tile_wb7StgFragPfPKfjjjjj, ( param0, param1, param2, param3, param4, param5, param6, param7 ); } // callseq 0 add.s64 %rd167, %rd160, 128; or.b32 %r332, %r327, 32; { // callseq 1, 0 .reg .b32 temp_param_reg; .param .align 4 .b8 param0[64]; st.param.f32 [param0+0], %f261; st.param.f32 [param0+4], %f262; st.param.f32 [param0+8], %f263; st.param.f32 [param0+12], %f264; st.param.f32 [param0+16], %f269; st.param.f32 [param0+20], %f270; st.param.f32 [param0+24], %f271; st.param.f32 [param0+28], %f272; st.param.f32 [param0+32], %f277; st.param.f32 [param0+36], %f278; st.param.f32 [param0+40], %f279; st.param.f32 [param0+44], %f280; st.param.f32 [param0+48], %f285; st.param.f32 [param0+52], %f286; st.param.f32 [param0+56], %f287; st.param.f32 [param0+60], %f288; .param .b64 param1; st.param.b64 [param1+0], %rd167; .param .b64 param2; st.param.b64 [param2+0], %rd166; .param .b32 param3; st.param.b32 [param3+0], %r170; .param .b32 param4; st.param.b32 [param4+0], %r341; .param .b32 param5; st.param.b32 [param5+0], %r340; .param .b32 param6; st.param.b32 [param6+0], %r175; .param .b32 param7; st.param.b32 [param7+0], %r332; call.uni _Z9C_tile_wb7StgFragPfPKfjjjjj, ( param0, param1, param2, param3, param4, param5, param6, param7 ); } // callseq 1 shl.b32 %r333, %r340, 4; or.b32 %r334, %r175, 16; mul.wide.u32 %rd168, %r333, 4; add.s64 %rd169, %rd160, %rd168; { // callseq 2, 0 .reg .b32 temp_param_reg; .param .align 4 .b8 param0[64]; st.param.f32 [param0+0], %f289; st.param.f32 [param0+4], %f290; st.param.f32 [param0+8], %f291; st.param.f32 [param0+12], %f292; st.param.f32 [param0+16], %f297; st.param.f32 [param0+20], %f298; st.param.f32 [param0+24], %f299; st.param.f32 [param0+28], %f300; st.param.f32 [param0+32], %f305; st.param.f32 [param0+36], %f306; st.param.f32 [param0+40], %f307; st.param.f32 [param0+44], %f308; st.param.f32 [param0+48], %f313; st.param.f32 [param0+52], %f314; st.param.f32 [param0+56], %f315; st.param.f32 [param0+60], %f316; .param .b64 param1; st.param.b64 [param1+0], %rd169; .param .b64 param2; st.param.b64 [param2+0], %rd166; .param .b32 param3; st.param.b32 [param3+0], %r170; .param .b32 param4; st.param.b32 [param4+0], %r341; .param .b32 param5; st.param.b32 [param5+0], %r340; .param .b32 param6; st.param.b32 [param6+0], %r334; .param .b32 param7; st.param.b32 [param7+0], %r327; call.uni _Z9C_tile_wb7StgFragPfPKfjjjjj, ( param0, param1, param2, param3, param4, param5, param6, param7 ); } // callseq 2 add.s64 %rd170, %rd169, 128; { // callseq 3, 0 .reg .b32 temp_param_reg; .param .align 4 .b8 param0[64]; st.param.f32 [param0+0], %f293; st.param.f32 [param0+4], %f294; st.param.f32 [param0+8], %f295; st.param.f32 [param0+12], %f296; st.param.f32 [param0+16], %f301; st.param.f32 [param0+20], %f302; st.param.f32 [param0+24], %f303; st.param.f32 [param0+28], %f304; st.param.f32 [param0+32], %f309; st.param.f32 [param0+36], %f310; st.param.f32 [param0+40], %f311; st.param.f32 [param0+44], %f312; st.param.f32 [param0+48], %f317; st.param.f32 [param0+52], %f318; st.param.f32 [param0+56], %f319; st.param.f32 [param0+60], %f320; .param .b64 param1; st.param.b64 [param1+0], %rd170; .param .b64 param2; st.param.b64 [param2+0], %rd166; .param .b32 param3; st.param.b32 [param3+0], %r170; .param .b32 param4; st.param.b32 [param4+0], %r341; .param .b32 param5; st.param.b32 [param5+0], %r340; .param .b32 param6; st.param.b32 [param6+0], %r334; .param .b32 param7; st.param.b32 [param7+0], %r332; call.uni _Z9C_tile_wb7StgFragPfPKfjjjjj, ( param0, param1, param2, param3, param4, param5, param6, param7 ); } // callseq 3 bra.uni $L__BB1_7; $L__BB1_5: mov.u32 %r348, %tid.x; and.b32 %r347, %r348, 31; mov.u32 %r346, %ctaid.x; bfi.b32 %r345, %r346, %r347, 7, 25; mov.u32 %r339, %tid.x; ld.param.u64 %rd171, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_2]; shl.b32 %r338, %r339, 1; mov.u32 %r337, _ZZ22sgemm_128x128x8_kernelPKfS0_PfjjjjjE4smem; ld.param.u32 %r336, [_Z22sgemm_128x128x8_kernelPKfS0_Pfjjjjj_param_4]; shl.b32 %r264, %r339, 6; and.b32 %r265, %r264, -2048; add.s32 %r267, %r337, %r265; and.b32 %r272, %r338, 64; or.b32 %r273, %r345, %r272; max.u32 %r274, %r273, %r336; sub.s32 %r275, %r274, %r273; shl.b32 %r276, %r339, 2; and.b32 %r277, %r276, 124; add.s32 %r278, %r267, %r277; bar.sync 0; // begin inline asm st.shared.v4.f32 [%r170], {%f257, %f258, %f259, %f260}; // end inline asm add.s32 %r244, %r170, 128; // begin inline asm st.shared.v4.f32 [%r244], {%f265, %f266, %f267, %f268}; // end inline asm add.s32 %r245, %r170, 256; // begin inline asm st.shared.v4.f32 [%r245], {%f273, %f274, %f275, %f276}; // end inline asm add.s32 %r246, %r170, 384; // begin inline asm st.shared.v4.f32 [%r246], {%f281, %f282, %f283, %f284}; // end inline asm bar.sync 0; setp.ne.s32 %p31, %r275, 0; selp.u32 %r242, 1, 0, %p31; mul.lo.s32 %r284, %r175, %r336; cvt.u64.u32 %rd121, %r284; cvt.u64.u32 %rd122, %r273; add.s64 %rd123, %rd121, %rd122; shl.b64 %rd124, %rd123, 2; add.s64 %rd57, %rd171, %rd124; ld.shared.f32 %f1649, [%r278]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd57], %f1649;} // end inline asm mul.wide.u32 %rd125, %r336, 4; add.s64 %rd58, %rd57, %rd125; ld.shared.f32 %f1650, [%r278+128]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd58], %f1650;} // end inline asm shl.b32 %r285, %r336, 1; mul.wide.u32 %rd126, %r285, 4; add.s64 %rd59, %rd57, %rd126; ld.shared.f32 %f1651, [%r278+256]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd59], %f1651;} // end inline asm mul.lo.s32 %r286, %r336, 3; mul.wide.u32 %rd127, %r286, 4; add.s64 %rd60, %rd57, %rd127; ld.shared.f32 %f1652, [%r278+384]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd60], %f1652;} // end inline asm shl.b32 %r287, %r336, 2; mul.wide.u32 %rd128, %r287, 4; add.s64 %rd61, %rd57, %rd128; ld.shared.f32 %f1653, [%r278+512]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd61], %f1653;} // end inline asm mul.lo.s32 %r288, %r336, 5; mul.wide.u32 %rd129, %r288, 4; add.s64 %rd62, %rd57, %rd129; ld.shared.f32 %f1654, [%r278+640]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd62], %f1654;} // end inline asm mul.lo.s32 %r289, %r336, 6; mul.wide.u32 %rd130, %r289, 4; add.s64 %rd63, %rd57, %rd130; ld.shared.f32 %f1655, [%r278+768]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd63], %f1655;} // end inline asm mul.lo.s32 %r290, %r336, 7; mul.wide.u32 %rd131, %r290, 4; add.s64 %rd64, %rd57, %rd131; ld.shared.f32 %f1656, [%r278+896]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd64], %f1656;} // end inline asm shl.b32 %r291, %r336, 3; mul.wide.u32 %rd132, %r291, 4; add.s64 %rd65, %rd57, %rd132; ld.shared.f32 %f1657, [%r278+1024]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd65], %f1657;} // end inline asm mul.lo.s32 %r292, %r336, 9; mul.wide.u32 %rd133, %r292, 4; add.s64 %rd66, %rd57, %rd133; ld.shared.f32 %f1658, [%r278+1152]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd66], %f1658;} // end inline asm mul.lo.s32 %r293, %r336, 10; mul.wide.u32 %rd134, %r293, 4; add.s64 %rd67, %rd57, %rd134; ld.shared.f32 %f1659, [%r278+1280]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd67], %f1659;} // end inline asm mul.lo.s32 %r294, %r336, 11; mul.wide.u32 %rd135, %r294, 4; add.s64 %rd68, %rd57, %rd135; ld.shared.f32 %f1660, [%r278+1408]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd68], %f1660;} // end inline asm mul.lo.s32 %r295, %r336, 12; mul.wide.u32 %rd136, %r295, 4; add.s64 %rd69, %rd57, %rd136; ld.shared.f32 %f1661, [%r278+1536]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd69], %f1661;} // end inline asm mul.lo.s32 %r296, %r336, 13; mul.wide.u32 %rd137, %r296, 4; add.s64 %rd70, %rd57, %rd137; ld.shared.f32 %f1662, [%r278+1664]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd70], %f1662;} // end inline asm mul.lo.s32 %r297, %r336, 14; mul.wide.u32 %rd138, %r297, 4; add.s64 %rd71, %rd57, %rd138; ld.shared.f32 %f1663, [%r278+1792]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd71], %f1663;} // end inline asm mul.lo.s32 %r298, %r336, 15; mul.wide.u32 %rd139, %r298, 4; add.s64 %rd72, %rd57, %rd139; ld.shared.f32 %f1664, [%r278+1920]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd72], %f1664;} // end inline asm bar.sync 0; // begin inline asm st.shared.v4.f32 [%r170], {%f261, %f262, %f263, %f264}; // end inline asm // begin inline asm st.shared.v4.f32 [%r244], {%f269, %f270, %f271, %f272}; // end inline asm // begin inline asm st.shared.v4.f32 [%r245], {%f277, %f278, %f279, %f280}; // end inline asm // begin inline asm st.shared.v4.f32 [%r246], {%f285, %f286, %f287, %f288}; // end inline asm bar.sync 0; setp.gt.u32 %p32, %r275, 32; selp.u32 %r262, 1, 0, %p32; add.s64 %rd73, %rd57, 128; ld.shared.f32 %f1681, [%r278]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd73], %f1681;} // end inline asm add.s64 %rd74, %rd73, %rd125; ld.shared.f32 %f1682, [%r278+128]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd74], %f1682;} // end inline asm add.s64 %rd75, %rd73, %rd126; ld.shared.f32 %f1683, [%r278+256]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd75], %f1683;} // end inline asm add.s64 %rd76, %rd73, %rd127; ld.shared.f32 %f1684, [%r278+384]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd76], %f1684;} // end inline asm add.s64 %rd77, %rd73, %rd128; ld.shared.f32 %f1685, [%r278+512]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd77], %f1685;} // end inline asm add.s64 %rd78, %rd73, %rd129; ld.shared.f32 %f1686, [%r278+640]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd78], %f1686;} // end inline asm add.s64 %rd79, %rd73, %rd130; ld.shared.f32 %f1687, [%r278+768]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd79], %f1687;} // end inline asm add.s64 %rd80, %rd73, %rd131; ld.shared.f32 %f1688, [%r278+896]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd80], %f1688;} // end inline asm add.s64 %rd81, %rd73, %rd132; ld.shared.f32 %f1689, [%r278+1024]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd81], %f1689;} // end inline asm add.s64 %rd82, %rd73, %rd133; ld.shared.f32 %f1690, [%r278+1152]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd82], %f1690;} // end inline asm add.s64 %rd83, %rd73, %rd134; ld.shared.f32 %f1691, [%r278+1280]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd83], %f1691;} // end inline asm add.s64 %rd84, %rd73, %rd135; ld.shared.f32 %f1692, [%r278+1408]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd84], %f1692;} // end inline asm add.s64 %rd85, %rd73, %rd136; ld.shared.f32 %f1693, [%r278+1536]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd85], %f1693;} // end inline asm add.s64 %rd86, %rd73, %rd137; ld.shared.f32 %f1694, [%r278+1664]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd86], %f1694;} // end inline asm add.s64 %rd87, %rd73, %rd138; ld.shared.f32 %f1695, [%r278+1792]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd87], %f1695;} // end inline asm add.s64 %rd88, %rd73, %rd139; ld.shared.f32 %f1696, [%r278+1920]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd88], %f1696;} // end inline asm shl.b32 %r299, %r336, 4; bar.sync 0; // begin inline asm st.shared.v4.f32 [%r170], {%f289, %f290, %f291, %f292}; // end inline asm // begin inline asm st.shared.v4.f32 [%r244], {%f297, %f298, %f299, %f300}; // end inline asm // begin inline asm st.shared.v4.f32 [%r245], {%f305, %f306, %f307, %f308}; // end inline asm // begin inline asm st.shared.v4.f32 [%r246], {%f313, %f314, %f315, %f316}; // end inline asm bar.sync 0; mul.wide.u32 %rd140, %r299, 4; add.s64 %rd89, %rd57, %rd140; ld.shared.f32 %f1713, [%r278]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd89], %f1713;} // end inline asm mul.lo.s32 %r300, %r336, 17; mul.wide.u32 %rd141, %r300, 4; add.s64 %rd90, %rd57, %rd141; ld.shared.f32 %f1714, [%r278+128]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd90], %f1714;} // end inline asm mul.lo.s32 %r301, %r336, 18; mul.wide.u32 %rd142, %r301, 4; add.s64 %rd91, %rd57, %rd142; ld.shared.f32 %f1715, [%r278+256]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd91], %f1715;} // end inline asm mul.lo.s32 %r302, %r336, 19; mul.wide.u32 %rd143, %r302, 4; add.s64 %rd92, %rd57, %rd143; ld.shared.f32 %f1716, [%r278+384]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd92], %f1716;} // end inline asm mul.lo.s32 %r303, %r336, 20; mul.wide.u32 %rd144, %r303, 4; add.s64 %rd93, %rd57, %rd144; ld.shared.f32 %f1717, [%r278+512]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd93], %f1717;} // end inline asm mul.lo.s32 %r304, %r336, 21; mul.wide.u32 %rd145, %r304, 4; add.s64 %rd94, %rd57, %rd145; ld.shared.f32 %f1718, [%r278+640]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd94], %f1718;} // end inline asm mul.lo.s32 %r305, %r336, 22; mul.wide.u32 %rd146, %r305, 4; add.s64 %rd95, %rd57, %rd146; ld.shared.f32 %f1719, [%r278+768]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd95], %f1719;} // end inline asm mul.lo.s32 %r306, %r336, 23; mul.wide.u32 %rd147, %r306, 4; add.s64 %rd96, %rd57, %rd147; ld.shared.f32 %f1720, [%r278+896]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd96], %f1720;} // end inline asm mul.lo.s32 %r307, %r336, 24; mul.wide.u32 %rd148, %r307, 4; add.s64 %rd97, %rd57, %rd148; ld.shared.f32 %f1721, [%r278+1024]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd97], %f1721;} // end inline asm mul.lo.s32 %r308, %r336, 25; mul.wide.u32 %rd149, %r308, 4; add.s64 %rd98, %rd57, %rd149; ld.shared.f32 %f1722, [%r278+1152]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd98], %f1722;} // end inline asm mul.lo.s32 %r309, %r336, 26; mul.wide.u32 %rd150, %r309, 4; add.s64 %rd99, %rd57, %rd150; ld.shared.f32 %f1723, [%r278+1280]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd99], %f1723;} // end inline asm mul.lo.s32 %r310, %r336, 27; mul.wide.u32 %rd151, %r310, 4; add.s64 %rd100, %rd57, %rd151; ld.shared.f32 %f1724, [%r278+1408]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd100], %f1724;} // end inline asm mul.lo.s32 %r311, %r336, 28; mul.wide.u32 %rd152, %r311, 4; add.s64 %rd101, %rd57, %rd152; ld.shared.f32 %f1725, [%r278+1536]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd101], %f1725;} // end inline asm mul.lo.s32 %r312, %r336, 29; mul.wide.u32 %rd153, %r312, 4; add.s64 %rd102, %rd57, %rd153; ld.shared.f32 %f1726, [%r278+1664]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd102], %f1726;} // end inline asm mul.lo.s32 %r313, %r336, 30; mul.wide.u32 %rd154, %r313, 4; add.s64 %rd103, %rd57, %rd154; ld.shared.f32 %f1727, [%r278+1792]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd103], %f1727;} // end inline asm mul.lo.s32 %r314, %r336, 31; mul.wide.u32 %rd155, %r314, 4; add.s64 %rd104, %rd57, %rd155; ld.shared.f32 %f1728, [%r278+1920]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r242, 0; @p st.global.f32 [%rd104], %f1728;} // end inline asm bar.sync 0; // begin inline asm st.shared.v4.f32 [%r170], {%f293, %f294, %f295, %f296}; // end inline asm // begin inline asm st.shared.v4.f32 [%r244], {%f301, %f302, %f303, %f304}; // end inline asm // begin inline asm st.shared.v4.f32 [%r245], {%f309, %f310, %f311, %f312}; // end inline asm // begin inline asm st.shared.v4.f32 [%r246], {%f317, %f318, %f319, %f320}; // end inline asm bar.sync 0; add.s64 %rd105, %rd73, %rd140; ld.shared.f32 %f1745, [%r278]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd105], %f1745;} // end inline asm add.s64 %rd106, %rd73, %rd141; ld.shared.f32 %f1746, [%r278+128]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd106], %f1746;} // end inline asm add.s64 %rd107, %rd73, %rd142; ld.shared.f32 %f1747, [%r278+256]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd107], %f1747;} // end inline asm add.s64 %rd108, %rd73, %rd143; ld.shared.f32 %f1748, [%r278+384]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd108], %f1748;} // end inline asm add.s64 %rd109, %rd73, %rd144; ld.shared.f32 %f1749, [%r278+512]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd109], %f1749;} // end inline asm add.s64 %rd110, %rd73, %rd145; ld.shared.f32 %f1750, [%r278+640]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd110], %f1750;} // end inline asm add.s64 %rd111, %rd73, %rd146; ld.shared.f32 %f1751, [%r278+768]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd111], %f1751;} // end inline asm add.s64 %rd112, %rd73, %rd147; ld.shared.f32 %f1752, [%r278+896]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd112], %f1752;} // end inline asm add.s64 %rd113, %rd73, %rd148; ld.shared.f32 %f1753, [%r278+1024]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd113], %f1753;} // end inline asm add.s64 %rd114, %rd73, %rd149; ld.shared.f32 %f1754, [%r278+1152]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd114], %f1754;} // end inline asm add.s64 %rd115, %rd73, %rd150; ld.shared.f32 %f1755, [%r278+1280]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd115], %f1755;} // end inline asm add.s64 %rd116, %rd73, %rd151; ld.shared.f32 %f1756, [%r278+1408]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd116], %f1756;} // end inline asm add.s64 %rd117, %rd73, %rd152; ld.shared.f32 %f1757, [%r278+1536]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd117], %f1757;} // end inline asm add.s64 %rd118, %rd73, %rd153; ld.shared.f32 %f1758, [%r278+1664]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd118], %f1758;} // end inline asm add.s64 %rd119, %rd73, %rd154; ld.shared.f32 %f1759, [%r278+1792]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd119], %f1759;} // end inline asm add.s64 %rd120, %rd73, %rd155; ld.shared.f32 %f1760, [%r278+1920]; // begin inline asm {.reg .pred p; setp.ne.b32 p, %r262, 0; @p st.global.f32 [%rd120], %f1760;} // end inline asm $L__BB1_7: ret; }