//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: UNKNOWN
// Unknown Toolkit Version
// Based on NVVM 7.0.1
//

.version 9.2
.target sm_120, texmode_independent
.address_size 64

	// .globl	DynamicKernel_nop_fsum_fsum_fsum

.entry DynamicKernel_nop_fsum_fsum_fsum(
	.param .u64 .ptr .global .align 8 DynamicKernel_nop_fsum_fsum_fsum_param_0,
	.param .u64 .ptr .global .align 8 DynamicKernel_nop_fsum_fsum_fsum_param_1,
	.param .u64 .ptr .global .align 8 DynamicKernel_nop_fsum_fsum_fsum_param_2
)
{
	.reg .pred 	%p<131>;
	.reg .b32 	%r<38>;
	.reg .f64 	%fd<97>;
	.reg .b64 	%rd<58>;


	ld.param.u64 	%rd16, [DynamicKernel_nop_fsum_fsum_fsum_param_0];
	ld.param.u64 	%rd57, [DynamicKernel_nop_fsum_fsum_fsum_param_1];
	ld.param.u64 	%rd18, [DynamicKernel_nop_fsum_fsum_fsum_param_2];
	mov.b32 	%r1, %envreg3;
	cvt.s64.s32 	%rd19, %r1;
	mov.u32 	%r11, %ntid.x;
	mov.u32 	%r12, %ctaid.x;
	mul.wide.s32 	%rd20, %r11, %r12;
	mov.u32 	%r13, %tid.x;
	cvt.s64.s32 	%rd21, %r13;
	add.s64 	%rd22, %rd21, %rd19;
	add.s64 	%rd1, %rd22, %rd20;
	cvt.u32.u64 	%r2, %rd1;
	setp.gt.s32 	%p1, %r2, 3;
	mov.f64 	%fd94, 0d0000000000000000;
	mov.f64 	%fd91, %fd94;
	@%p1 bra 	$L__BB0_38;

	and.b32  	%r14, %r2, 1;
	setp.eq.b32 	%p2, %r14, 1;
	mov.pred 	%p3, 0;
	xor.pred  	%p4, %p2, %p3;
	not.pred 	%p5, %p4;
	mov.f64 	%fd91, 0d0000000000000000;
	mov.u32 	%r35, %r2;
	@%p5 bra 	$L__BB0_3;

	shl.b64 	%rd23, %rd1, 32;
	shr.s64 	%rd24, %rd23, 29;
	add.s64 	%rd25, %rd18, %rd24;
	ld.global.f64 	%fd42, [%rd25];
	abs.f64 	%fd43, %fd42;
	setp.gtu.f64 	%p6, %fd43, 0d7FF0000000000000;
	add.f64 	%fd44, %fd42, 0d0000000000000000;
	selp.f64 	%fd91, 0d0000000000000000, %fd44, %p6;
	add.s32 	%r35, %r2, 1;

$L__BB0_3:
	setp.eq.s32 	%p7, %r2, 3;
	@%p7 bra 	$L__BB0_38;

	add.s32 	%r36, %r35, -2;
	mul.wide.s32 	%rd26, %r35, 8;
	add.s64 	%rd27, %rd18, %rd26;
	add.s64 	%rd56, %rd27, 8;

$L__BB0_5:
	add.s64 	%rd4, %rd56, -8;
	ld.global.f64 	%fd5, [%rd56+-8];
	abs.f64 	%fd6, %fd5;
	setp.gtu.f64 	%p8, %fd6, 0d7FF0000000000000;
	@%p8 bra 	$L__BB0_20;
	bra.uni 	$L__BB0_6;

$L__BB0_20:
	add.f64 	%fd90, %fd91, 0d0000000000000000;
	bra.uni 	$L__BB0_21;

$L__BB0_6:
	setp.lt.f64 	%p9, %fd5, 0d0000000000000000;
	setp.gt.f64 	%p10, %fd91, 0d0000000000000000;
	and.pred  	%p11, %p10, %p9;
	@%p11 bra 	$L__BB0_8;

	setp.geu.f64 	%p12, %fd91, 0d0000000000000000;
	setp.leu.f64 	%p13, %fd5, 0d0000000000000000;
	or.pred  	%p14, %p12, %p13;
	@%p14 bra 	$L__BB0_19;

$L__BB0_8:
	neg.f64 	%fd7, %fd91;
	setp.eq.f64 	%p15, %fd5, %fd7;
	mov.f64 	%fd90, 0d0000000000000000;
	@%p15 bra 	$L__BB0_21;

	setp.eq.f64 	%p16, %fd5, 0d0000000000000000;
	setp.eq.f64 	%p17, %fd91, 0d8000000000000000;
	or.pred  	%p18, %p17, %p16;
	@%p18 bra 	$L__BB0_19;

	add.f64 	%fd46, %fd91, %fd5;
	abs.f64 	%fd8, %fd46;
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r15}, %fd8;
	}
	and.b32  	%r16, %r15, 2146435072;
	setp.eq.s32 	%p19, %r16, 2146435072;
	mul.f64 	%fd47, %fd6, 0d3D30000000000000;
	setp.gt.f64 	%p20, %fd8, %fd47;
	or.pred  	%p21, %p19, %p20;
	@%p21 bra 	$L__BB0_19;

	abs.f64 	%fd9, %fd7;
	mul.f64 	%fd48, %fd9, 0d3D30000000000000;
	setp.gt.f64 	%p22, %fd8, %fd48;
	@%p22 bra 	$L__BB0_19;

	setp.gtu.f64 	%p23, %fd8, 0d433FFFFFFFFFFFFF;
	@%p23 bra 	$L__BB0_18;

	cvt.rzi.s64.f64 	%rd5, %fd8;
	setp.gt.s64 	%p24, %rd5, 9007199254740991;
	@%p24 bra 	$L__BB0_18;

	cvt.rn.f64.s64 	%fd49, %rd5;
	setp.ne.f64 	%p25, %fd8, %fd49;
	setp.gtu.f64 	%p26, %fd6, 0d433FFFFFFFFFFFFF;
	or.pred  	%p27, %p25, %p26;
	@%p27 bra 	$L__BB0_18;

	cvt.rzi.s64.f64 	%rd6, %fd6;
	setp.gt.s64 	%p28, %rd6, 9007199254740991;
	@%p28 bra 	$L__BB0_18;

	cvt.rn.f64.s64 	%fd50, %rd6;
	setp.ne.f64 	%p29, %fd6, %fd50;
	setp.gtu.f64 	%p30, %fd9, 0d433FFFFFFFFFFFFF;
	or.pred  	%p31, %p29, %p30;
	@%p31 bra 	$L__BB0_18;

	cvt.rzi.s64.f64 	%rd28, %fd9;
	setp.lt.s64 	%p32, %rd28, 9007199254740992;
	cvt.rn.f64.s64 	%fd51, %rd28;
	setp.equ.f64 	%p33, %fd9, %fd51;
	and.pred  	%p34, %p32, %p33;
	@%p34 bra 	$L__BB0_19;

$L__BB0_18:
	mul.f64 	%fd53, %fd6, 0d3CF0000000000000;
	setp.lt.f64 	%p35, %fd8, %fd53;
	mul.f64 	%fd54, %fd9, 0d3CF0000000000000;
	setp.lt.f64 	%p36, %fd8, %fd54;
	and.pred  	%p37, %p35, %p36;
	@%p37 bra 	$L__BB0_21;

$L__BB0_19:
	add.f64 	%fd90, %fd91, %fd5;

$L__BB0_21:
	ld.global.f64 	%fd13, [%rd4+8];
	abs.f64 	%fd14, %fd13;
	setp.gtu.f64 	%p38, %fd14, 0d7FF0000000000000;
	@%p38 bra 	$L__BB0_36;
	bra.uni 	$L__BB0_22;

$L__BB0_36:
	add.f64 	%fd91, %fd90, 0d0000000000000000;
	bra.uni 	$L__BB0_37;

$L__BB0_22:
	setp.lt.f64 	%p39, %fd13, 0d0000000000000000;
	setp.gt.f64 	%p40, %fd90, 0d0000000000000000;
	and.pred  	%p41, %p40, %p39;
	@%p41 bra 	$L__BB0_24;

	setp.geu.f64 	%p42, %fd90, 0d0000000000000000;
	setp.leu.f64 	%p43, %fd13, 0d0000000000000000;
	or.pred  	%p44, %p42, %p43;
	@%p44 bra 	$L__BB0_35;

$L__BB0_24:
	neg.f64 	%fd15, %fd90;
	setp.eq.f64 	%p45, %fd13, %fd15;
	mov.f64 	%fd91, 0d0000000000000000;
	@%p45 bra 	$L__BB0_37;

	setp.eq.f64 	%p46, %fd13, 0d0000000000000000;
	setp.eq.f64 	%p47, %fd90, 0d8000000000000000;
	or.pred  	%p48, %p47, %p46;
	@%p48 bra 	$L__BB0_35;

	add.f64 	%fd56, %fd90, %fd13;
	abs.f64 	%fd16, %fd56;
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r17}, %fd16;
	}
	and.b32  	%r18, %r17, 2146435072;
	setp.eq.s32 	%p49, %r18, 2146435072;
	mul.f64 	%fd57, %fd14, 0d3D30000000000000;
	setp.gt.f64 	%p50, %fd16, %fd57;
	or.pred  	%p51, %p49, %p50;
	@%p51 bra 	$L__BB0_35;

	abs.f64 	%fd17, %fd15;
	mul.f64 	%fd58, %fd17, 0d3D30000000000000;
	setp.gt.f64 	%p52, %fd16, %fd58;
	@%p52 bra 	$L__BB0_35;

	setp.gtu.f64 	%p53, %fd16, 0d433FFFFFFFFFFFFF;
	@%p53 bra 	$L__BB0_34;

	cvt.rzi.s64.f64 	%rd7, %fd16;
	setp.gt.s64 	%p54, %rd7, 9007199254740991;
	@%p54 bra 	$L__BB0_34;

	cvt.rn.f64.s64 	%fd59, %rd7;
	setp.ne.f64 	%p55, %fd16, %fd59;
	setp.gtu.f64 	%p56, %fd14, 0d433FFFFFFFFFFFFF;
	or.pred  	%p57, %p55, %p56;
	@%p57 bra 	$L__BB0_34;

	cvt.rzi.s64.f64 	%rd8, %fd14;
	setp.gt.s64 	%p58, %rd8, 9007199254740991;
	@%p58 bra 	$L__BB0_34;

	cvt.rn.f64.s64 	%fd60, %rd8;
	setp.ne.f64 	%p59, %fd14, %fd60;
	setp.gtu.f64 	%p60, %fd17, 0d433FFFFFFFFFFFFF;
	or.pred  	%p61, %p59, %p60;
	@%p61 bra 	$L__BB0_34;

	cvt.rzi.s64.f64 	%rd29, %fd17;
	setp.lt.s64 	%p62, %rd29, 9007199254740992;
	cvt.rn.f64.s64 	%fd61, %rd29;
	setp.equ.f64 	%p63, %fd17, %fd61;
	and.pred  	%p64, %p62, %p63;
	@%p64 bra 	$L__BB0_35;

$L__BB0_34:
	mul.f64 	%fd63, %fd14, 0d3CF0000000000000;
	setp.lt.f64 	%p65, %fd16, %fd63;
	mul.f64 	%fd64, %fd17, 0d3CF0000000000000;
	setp.lt.f64 	%p66, %fd16, %fd64;
	and.pred  	%p67, %p65, %p66;
	@%p67 bra 	$L__BB0_37;

$L__BB0_35:
	add.f64 	%fd91, %fd90, %fd13;

$L__BB0_37:
	add.s64 	%rd56, %rd56, 16;
	add.s32 	%r36, %r36, 2;
	setp.lt.s32 	%p68, %r36, 2;
	@%p68 bra 	$L__BB0_5;

$L__BB0_38:
	add.f64 	%fd22, %fd91, 0d0000000000000000;
	cvt.u64.u32 	%rd30, %r13;
	cvt.u64.u32 	%rd31, %r1;
	add.s64 	%rd32, %rd30, %rd31;
	cvt.u64.u32 	%rd33, %r11;
	cvt.u64.u32 	%rd34, %r12;
	mul.lo.s64 	%rd35, %rd33, %rd34;
	add.s64 	%rd36, %rd32, %rd35;
	cvt.u32.u64 	%r22, %rd36;
	setp.lt.s32 	%p69, %r22, -4;
	@%p69 bra 	$L__BB0_57;

	add.s32 	%r8, %r22, 4;
	mov.f64 	%fd94, 0d0000000000000000;
	mov.u32 	%r37, 0;

$L__BB0_40:
	mov.f64 	%fd23, %fd94;
	mov.u32 	%r9, %r37;
	ld.global.f64 	%fd24, [%rd57];
	abs.f64 	%fd25, %fd24;
	setp.gtu.f64 	%p70, %fd25, 0d7FF0000000000000;
	@%p70 bra 	$L__BB0_55;
	bra.uni 	$L__BB0_41;

$L__BB0_55:
	add.f64 	%fd94, %fd23, 0d0000000000000000;
	bra.uni 	$L__BB0_56;

$L__BB0_41:
	setp.lt.f64 	%p71, %fd24, 0d0000000000000000;
	setp.gt.f64 	%p72, %fd23, 0d0000000000000000;
	and.pred  	%p73, %p72, %p71;
	@%p73 bra 	$L__BB0_43;

	setp.geu.f64 	%p74, %fd23, 0d0000000000000000;
	setp.leu.f64 	%p75, %fd24, 0d0000000000000000;
	or.pred  	%p76, %p74, %p75;
	@%p76 bra 	$L__BB0_54;

$L__BB0_43:
	neg.f64 	%fd26, %fd23;
	setp.eq.f64 	%p77, %fd24, %fd26;
	mov.f64 	%fd94, 0d0000000000000000;
	@%p77 bra 	$L__BB0_56;

	setp.eq.f64 	%p78, %fd24, 0d0000000000000000;
	setp.eq.f64 	%p79, %fd23, 0d8000000000000000;
	or.pred  	%p80, %p79, %p78;
	@%p80 bra 	$L__BB0_54;

	add.f64 	%fd68, %fd23, %fd24;
	abs.f64 	%fd27, %fd68;
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r28}, %fd27;
	}
	and.b32  	%r29, %r28, 2146435072;
	setp.eq.s32 	%p81, %r29, 2146435072;
	mul.f64 	%fd69, %fd25, 0d3D30000000000000;
	setp.gt.f64 	%p82, %fd27, %fd69;
	or.pred  	%p83, %p81, %p82;
	@%p83 bra 	$L__BB0_54;

	abs.f64 	%fd28, %fd26;
	mul.f64 	%fd70, %fd28, 0d3D30000000000000;
	setp.gt.f64 	%p84, %fd27, %fd70;
	@%p84 bra 	$L__BB0_54;

	setp.gtu.f64 	%p85, %fd27, 0d433FFFFFFFFFFFFF;
	@%p85 bra 	$L__BB0_53;

	cvt.rzi.s64.f64 	%rd11, %fd27;
	setp.gt.s64 	%p86, %rd11, 9007199254740991;
	@%p86 bra 	$L__BB0_53;

	cvt.rn.f64.s64 	%fd71, %rd11;
	setp.ne.f64 	%p87, %fd27, %fd71;
	setp.gtu.f64 	%p88, %fd25, 0d433FFFFFFFFFFFFF;
	or.pred  	%p89, %p87, %p88;
	@%p89 bra 	$L__BB0_53;

	cvt.rzi.s64.f64 	%rd12, %fd25;
	setp.gt.s64 	%p90, %rd12, 9007199254740991;
	@%p90 bra 	$L__BB0_53;

	cvt.rn.f64.s64 	%fd72, %rd12;
	setp.ne.f64 	%p91, %fd25, %fd72;
	setp.gtu.f64 	%p92, %fd28, 0d433FFFFFFFFFFFFF;
	or.pred  	%p93, %p91, %p92;
	@%p93 bra 	$L__BB0_53;

	cvt.rzi.s64.f64 	%rd44, %fd28;
	setp.lt.s64 	%p94, %rd44, 9007199254740992;
	cvt.rn.f64.s64 	%fd73, %rd44;
	setp.equ.f64 	%p95, %fd28, %fd73;
	and.pred  	%p96, %p94, %p95;
	@%p96 bra 	$L__BB0_54;

$L__BB0_53:
	mul.f64 	%fd75, %fd25, 0d3CF0000000000000;
	setp.lt.f64 	%p97, %fd27, %fd75;
	mul.f64 	%fd76, %fd28, 0d3CF0000000000000;
	setp.lt.f64 	%p98, %fd27, %fd76;
	and.pred  	%p99, %p97, %p98;
	@%p99 bra 	$L__BB0_56;

$L__BB0_54:
	add.f64 	%fd94, %fd23, %fd24;

$L__BB0_56:
	add.s32 	%r37, %r9, 1;
	setp.lt.u32 	%p100, %r37, 6;
	setp.lt.s32 	%p101, %r9, %r8;
	and.pred  	%p102, %p100, %p101;
	add.s64 	%rd57, %rd57, 8;
	@%p102 bra 	$L__BB0_40;

$L__BB0_57:
	setp.gt.f64 	%p103, %fd22, 0d0000000000000000;
	setp.lt.f64 	%p104, %fd94, 0d0000000000000000;
	and.pred  	%p105, %p103, %p104;
	@%p105 bra 	$L__BB0_59;

	setp.geu.f64 	%p106, %fd22, 0d0000000000000000;
	setp.leu.f64 	%p107, %fd94, 0d0000000000000000;
	or.pred  	%p108, %p106, %p107;
	@%p108 bra 	$L__BB0_71;

$L__BB0_59:
	neg.f64 	%fd33, %fd22;
	setp.eq.f64 	%p109, %fd94, %fd33;
	mov.f64 	%fd96, 0d0000000000000000;
	@%p109 bra 	$L__BB0_72;

	setp.eq.f64 	%p110, %fd94, 0d0000000000000000;
	setp.eq.f64 	%p111, %fd22, 0d8000000000000000;
	or.pred  	%p112, %p111, %p110;
	@%p112 bra 	$L__BB0_71;

	add.f64 	%fd78, %fd22, %fd94;
	abs.f64 	%fd34, %fd78;
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r30}, %fd34;
	}
	and.b32  	%r31, %r30, 2146435072;
	setp.eq.s32 	%p113, %r31, 2146435072;
	@%p113 bra 	$L__BB0_71;

	abs.f64 	%fd35, %fd94;
	mul.f64 	%fd79, %fd35, 0d3D30000000000000;
	setp.gt.f64 	%p114, %fd34, %fd79;
	@%p114 bra 	$L__BB0_71;

	abs.f64 	%fd36, %fd33;
	mul.f64 	%fd80, %fd36, 0d3D30000000000000;
	setp.gt.f64 	%p115, %fd34, %fd80;
	@%p115 bra 	$L__BB0_71;

	setp.gtu.f64 	%p116, %fd34, 0d433FFFFFFFFFFFFF;
	@%p116 bra 	$L__BB0_70;

	cvt.rzi.s64.f64 	%rd14, %fd34;
	setp.gt.s64 	%p117, %rd14, 9007199254740991;
	@%p117 bra 	$L__BB0_70;

	cvt.rn.f64.s64 	%fd81, %rd14;
	setp.ne.f64 	%p118, %fd34, %fd81;
	setp.gtu.f64 	%p119, %fd35, 0d433FFFFFFFFFFFFF;
	or.pred  	%p120, %p118, %p119;
	@%p120 bra 	$L__BB0_70;

	cvt.rzi.s64.f64 	%rd15, %fd35;
	setp.gt.s64 	%p121, %rd15, 9007199254740991;
	@%p121 bra 	$L__BB0_70;

	cvt.rn.f64.s64 	%fd82, %rd15;
	setp.ne.f64 	%p122, %fd35, %fd82;
	setp.gtu.f64 	%p123, %fd36, 0d433FFFFFFFFFFFFF;
	or.pred  	%p124, %p122, %p123;
	@%p124 bra 	$L__BB0_70;

	cvt.rzi.s64.f64 	%rd45, %fd36;
	setp.lt.s64 	%p125, %rd45, 9007199254740992;
	cvt.rn.f64.s64 	%fd83, %rd45;
	setp.equ.f64 	%p126, %fd36, %fd83;
	and.pred  	%p127, %p125, %p126;
	@%p127 bra 	$L__BB0_71;

$L__BB0_70:
	mul.f64 	%fd85, %fd35, 0d3CF0000000000000;
	setp.lt.f64 	%p128, %fd34, %fd85;
	mul.f64 	%fd86, %fd36, 0d3CF0000000000000;
	setp.lt.f64 	%p129, %fd34, %fd86;
	and.pred  	%p130, %p128, %p129;
	@%p130 bra 	$L__BB0_72;

$L__BB0_71:
	add.f64 	%fd96, %fd22, %fd94;

$L__BB0_72:
	shl.b64 	%rd53, %rd36, 32;
	shr.s64 	%rd54, %rd53, 29;
	add.s64 	%rd55, %rd16, %rd54;
	st.global.f64 	[%rd55], %fd96;
	ret;

}

  