cvta.to.global.u64 %rd1, %rd2; mov.u32 %r25, %ctaid.x; shl.b32 %r26, %r25, 9; mov.u32 %r1, %tid.x; add.s32 %r2, %r26, %r1; div.u32 %r4, %r2, %r16; setp.ge.u32 %p1, %r4, %r23; @%p1 bra $L__BB0_35; rem.u32 %r27, %r2, %r16; shl.b32 %r6, %r27, 4; mad.lo.s32 %r7, %r4, %r22, %r6; setp.ge.u32 %p2, %r6, %r22; mov.f32 %f106, 0f00000000; mov.f32 %f107, %f106; @%p2 bra $L__BB0_3; mul.wide.u32 %rd8, %r7, 4; add.s64 %rd9, %rd1, %rd8; ld.global.f32 %f70, [%rd9]; add.f32 %f106, %f70, 0f00000000; fma.rn.f32 %f107, %f70, %f70, 0f00000000; $L__BB0_3: add.s32 %r28, %r6, 1; setp.ge.u32 %p3, %r28, %r22; @%p3 bra $L__BB0_5; add.s32 %r29, %r7, 1; mul.wide.u32 %rd10, %r29, 4; add.s64 %rd11, %rd1, %rd10; ld.global.f32 %f71, [%rd11]; add.f32 %f106, %f106, %f71; fma.rn.f32 %f107, %f71, %f71, %f107; $L__BB0_5: add.s32 %r30, %r6, 2; setp.ge.u32 %p4, %r30, %r22; @%p4 bra $L__BB0_7; add.s32 %r31, %r7, 2; mul.wide.u32 %rd12, %r31, 4; add.s64 %rd13, %rd1, %rd12; ld.global.f32 %f72, [%rd13]; add.f32 %f106, %f106, %f72; fma.rn.f32 %f107, %f72, %f72, %f107; $L__BB0_7: add.s32 %r32, %r6, 3; setp.ge.u32 %p5, %r32, %r22; @%p5 bra $L__BB0_9; add.s32 %r33, %r7, 3; mul.wide.u32 %rd14, %r33, 4; add.s64 %rd15, %rd1, %rd14; ld.global.f32 %f73, [%rd15]; add.f32 %f106, %f106, %f73; fma.rn.f32 %f107, %f73, %f73, %f107; $L__BB0_9: add.s32 %r34, %r6, 4; setp.ge.u32 %p6, %r34, %r22; @%p6 bra $L__BB0_11; add.s32 %r35, %r7, 4; mul.wide.u32 %rd16, %r35, 4; add.s64 %rd17, %rd1, %rd16; ld.global.f32 %f74, [%rd17]; add.f32 %f106, %f106, %f74; fma.rn.f32 %f107, %f74, %f74, %f107; $L__BB0_11: add.s32 %r36, %r6, 5; setp.ge.u32 %p7, %r36, %r22; @%p7 bra $L__BB0_13; add.s32 %r37, %r7, 5; mul.wide.u32 %rd18, %r37, 4; add.s64 %rd19, %rd1, %rd18; ld.global.f32 %f75, [%rd19]; add.f32 %f106, %f106, %f75; fma.rn.f32 %f107, %f75, %f75, %f107; $L__BB0_13: add.s32 %r38, %r6, 6; setp.ge.u32 %p8, %r38, %r22; @%p8 bra $L__BB0_15; add.s32 %r39, %r7, 6; mul.wide.u32 %rd20, %r39, 4; add.s64 %rd21, %rd1, %rd20; ld.global.f32 %f76, [%rd21]; add.f32 %f106, %f106, %f76; fma.rn.f32 %f107, %f76, %f76, %f107; $L__BB0_15: add.s32 %r40, %r6, 7; setp.ge.u32 %p9, %r40, %r22; @%p9 bra $L__BB0_17; add.s32 %r41, %r7, 7; mul.wide.u32 %rd22, %r41, 4; add.s64 %rd23, %rd1, %rd22; ld.global.f32 %f77, [%rd23]; add.f32 %f106, %f106, %f77; fma.rn.f32 %f107, %f77, %f77, %f107; $L__BB0_17: add.s32 %r42, %r6, 8; setp.ge.u32 %p10, %r42, %r22; @%p10 bra $L__BB0_19; add.s32 %r43, %r7, 8; mul.wide.u32 %rd24, %r43, 4; add.s64 %rd25, %rd1, %rd24; ld.global.f32 %f78, [%rd25]; add.f32 %f106, %f106, %f78; fma.rn.f32 %f107, %f78, %f78, %f107; $L__BB0_19: add.s32 %r44, %r6, 9; setp.ge.u32 %p11, %r44, %r22; @%p11 bra $L__BB0_21; add.s32 %r45, %r7, 9; mul.wide.u32 %rd26, %r45, 4; add.s64 %rd27, %rd1, %rd26; ld.global.f32 %f79, [%rd27]; add.f32 %f106, %f106, %f79; fma.rn.f32 %f107, %f79, %f79, %f107; $L__BB0_21: add.s32 %r46, %r6, 10; setp.ge.u32 %p12, %r46, %r22; @%p12 bra $L__BB0_23; add.s32 %r47, %r7, 10; mul.wide.u32 %rd28, %r47, 4; add.s64 %rd29, %rd1, %rd28; ld.global.f32 %f80, [%rd29]; add.f32 %f106, %f106, %f80; fma.rn.f32 %f107, %f80, %f80, %f107; $L__BB0_23: add.s32 %r48, %r6, 11; setp.ge.u32 %p13, %r48, %r22; @%p13 bra $L__BB0_25; add.s32 %r49, %r7, 11; mul.wide.u32 %rd30, %r49, 4; add.s64 %rd31, %rd1, %rd30; ld.global.f32 %f81, [%rd31]; add.f32 %f106, %f106, %f81; fma.rn.f32 %f107, %f81, %f81, %f107; $L__BB0_25: add.s32 %r50, %r6, 12; setp.ge.u32 %p14, %r50, %r22; @%p14 bra $L__BB0_27; add.s32 %r51, %r7, 12; mul.wide.u32 %rd32, %r51, 4; add.s64 %rd33, %rd1, %rd32; ld.global.f32 %f82, [%rd33]; add.f32 %f106, %f106, %f82; fma.rn.f32 %f107, %f82, %f82, %f107; $L__BB0_27: add.s32 %r52, %r6, 13; setp.ge.u32 %p15, %r52, %r22; @%p15 bra $L__BB0_29; add.s32 %r53, %r7, 13; mul.wide.u32 %rd34, %r53, 4; add.s64 %rd35, %rd1, %rd34; ld.global.f32 %f83, [%rd35]; add.f32 %f106, %f106, %f83; fma.rn.f32 %f107, %f83, %f83, %f107; $L__BB0_29: add.s32 %r54, %r6, 14; setp.ge.u32 %p16, %r54, %r22; @%p16 bra $L__BB0_31; add.s32 %r55, %r7, 14; mul.wide.u32 %rd36, %r55, 4; add.s64 %rd37, %rd1, %rd36; ld.global.f32 %f84, [%rd37]; add.f32 %f106, %f106, %f84; fma.rn.f32 %f107, %f84, %f84, %f107; $L__BB0_31: add.s32 %r56, %r6, 15; setp.ge.u32 %p17, %r56, %r22; @%p17 bra $L__BB0_33; add.s32 %r57, %r7, 15; mul.wide.u32 %rd38, %r57, 4; add.s64 %rd39, %rd1, %rd38; ld.global.f32 %f85, [%rd39]; add.f32 %f106, %f106, %f85; fma.rn.f32 %f107, %f85, %f85, %f107; $L__BB0_33: mov.b32 %r58, %f106; mov.u32 %r59, 2; mov.u32 %r60, 31; mov.u32 %r61, 1; mov.u32 %r62, -1; shfl.sync.down.b32 %r63|%p18, %r58, %r61, %r60, %r62; mov.b32 %f86, %r63; add.f32 %f87, %f106, %f86; mov.b32 %r64, %f107; shfl.sync.down.b32 %r65|%p19, %r64, %r61, %r60, %r62; mov.b32 %f88, %r65; add.f32 %f89, %f107, %f88; mov.b32 %r66, %f87; shfl.sync.down.b32 %r67|%p20, %r66, %r59, %r60, %r62; mov.b32 %f90, %r67; add.f32 %f91, %f87, %f90; mov.b32 %r68, %f89; shfl.sync.down.b32 %r69|%p21, %r68, %r59, %r60, %r62; mov.b32 %f92, %r69; add.f32 %f93, %f89, %f92; mov.b32 %r70, %f91; mov.u32 %r71, 4; shfl.sync.down.b32 %r72|%p22, %r70, %r71, %r60, %r62; mov.b32 %f94, %r72; add.f32 %f95, %f91, %f94; mov.b32 %r73, %f93; shfl.sync.down.b32 %r74|%p23, %r73, %r71, %r60, %r62; mov.b32 %f96, %r74; add.f32 %f97, %f93, %f96; mov.b32 %r75, %f95; mov.u32 %r76, 8; shfl.sync.down.b32 %r77|%p24, %r75, %r76, %r60, %r62; mov.b32 %f98, %r77; add.f32 %f99, %f95, %f98; mov.b32 %r78, %f97; shfl.sync.down.b32 %r79|%p25, %r78, %r76, %r60, %r62; mov.b32 %f100, %r79; add.f32 %f101, %f97, %f100; mov.b32 %r80, %f99; mov.u32 %r81, 16; shfl.sync.down.b32 %r82|%p26, %r80, %r81, %r60, %r62; mov.b32 %f102, %r82; add.f32 %f65, %f99, %f102; mov.b32 %r83, %f101; shfl.sync.down.b32 %r84|%p27, %r83, %r81, %r60, %r62; mov.b32 %f103, %r84; add.f32 %f66, %f101, %f103; and.b32 %r85, %r1, 31; setp.ne.s32 %p28, %r85, 0; @%p28 bra $L__BB0_35; cvta.to.global.u64 %rd40, %rd3; mul.wide.u32 %rd41, %r4, 8; add.s64 %rd42, %rd40, %rd41; cvt.f64.f32 %fd1, %f65; atom.global.add.f64 %fd2, [%rd42], %fd1; cvta.to.global.u64 %rd43, %rd4; add.s64 %rd44, %rd43, %rd41; cvt.f64.f32 %fd3, %f66; atom.global.add.f64 %fd4, [%rd44], %fd3; $L__BB0_35: ret; }