la_helper_funs.ptx

//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Tue Apr  1 01:34:02 2014 (1396341242)
// Cuda compilation tools, release 6.0, V6.0.1
//

.version 4.0
.target sm_30
.address_size 64

.func  (.param .b64 func_retval0) __internal_accurate_pow
(
	.param .b64 __internal_accurate_pow_param_0,
	.param .b64 __internal_accurate_pow_param_1
)
;

.visible .entry vectorScalarSet(
	.param .u64 vectorScalarSet_param_0,
	.param .f32 vectorScalarSet_param_1,
	.param .u32 vectorScalarSet_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<2>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<5>;


	ld.param.u64 	%rd1, [vectorScalarSet_param_0];
	ld.param.f32 	%f1, [vectorScalarSet_param_1];
	ld.param.u32 	%r2, [vectorScalarSet_param_2];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB0_2;

	cvta.to.global.u64 	%rd2, %rd1;
	mul.wide.s32 	%rd3, %r1, 4;
	add.s64 	%rd4, %rd2, %rd3;
	st.global.f32 	[%rd4], %f1;

BB0_2:
	ret;
}

.visible .entry vectorScalarAdd(
	.param .u64 vectorScalarAdd_param_0,
	.param .u64 vectorScalarAdd_param_1,
	.param .f32 vectorScalarAdd_param_2,
	.param .u32 vectorScalarAdd_param_3
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<4>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [vectorScalarAdd_param_0];
	ld.param.u64 	%rd2, [vectorScalarAdd_param_1];
	ld.param.f32 	%f1, [vectorScalarAdd_param_2];
	ld.param.u32 	%r2, [vectorScalarAdd_param_3];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB1_2;

	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f2, [%rd6];
	add.ftz.f32 	%f3, %f2, %f1;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;

BB1_2:
	ret;
}

.visible .entry vectorLog(
	.param .u64 vectorLog_param_0,
	.param .u64 vectorLog_param_1,
	.param .u32 vectorLog_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<4>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [vectorLog_param_0];
	ld.param.u64 	%rd2, [vectorLog_param_1];
	ld.param.u32 	%r2, [vectorLog_param_2];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB2_2;

	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f1, [%rd6];
	lg2.approx.ftz.f32 	%f2, %f1;
	mul.ftz.f32 	%f3, %f2, 0f3F317218;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;

BB2_2:
	ret;
}

.visible .entry vectorExp(
	.param .u64 vectorExp_param_0,
	.param .u64 vectorExp_param_1,
	.param .u32 vectorExp_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<4>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [vectorExp_param_0];
	ld.param.u64 	%rd2, [vectorExp_param_1];
	ld.param.u32 	%r2, [vectorExp_param_2];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB3_2;

	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f1, [%rd6];
	mul.ftz.f32 	%f2, %f1, 0f3FB8AA3B;
	ex2.approx.ftz.f32 	%f3, %f2;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;

BB3_2:
	ret;
}

.visible .entry vectorSign(
	.param .u64 vectorSign_param_0,
	.param .u64 vectorSign_param_1,
	.param .u32 vectorSign_param_2
)
{
	.reg .pred 	%p<3>;
	.reg .f32 	%f<3>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [vectorSign_param_0];
	ld.param.u64 	%rd2, [vectorSign_param_1];
	ld.param.u32 	%r2, [vectorSign_param_2];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB4_2;

	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f1, [%rd6];
	setp.gt.ftz.f32	%p2, %f1, 0f00000000;
	selp.f32	%f2, 0f3F800000, 0fBF800000, %p2;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f2;

BB4_2:
	ret;
}

.visible .entry vectorAbs(
	.param .u64 vectorAbs_param_0,
	.param .u64 vectorAbs_param_1,
	.param .u32 vectorAbs_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<3>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [vectorAbs_param_0];
	ld.param.u64 	%rd2, [vectorAbs_param_1];
	ld.param.u32 	%r2, [vectorAbs_param_2];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB5_2;

	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f1, [%rd6];
	abs.ftz.f32 	%f2, %f1;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f2;

BB5_2:
	ret;
}

.visible .entry vectorDiv(
	.param .u64 vectorDiv_param_0,
	.param .u64 vectorDiv_param_1,
	.param .u64 vectorDiv_param_2,
	.param .u32 vectorDiv_param_3
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<4>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<11>;


	ld.param.u64 	%rd1, [vectorDiv_param_0];
	ld.param.u64 	%rd2, [vectorDiv_param_1];
	ld.param.u64 	%rd3, [vectorDiv_param_2];
	ld.param.u32 	%r2, [vectorDiv_param_3];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB6_2;

	cvta.to.global.u64 	%rd4, %rd3;
	cvta.to.global.u64 	%rd5, %rd2;
	cvta.to.global.u64 	%rd6, %rd1;
	mul.wide.s32 	%rd7, %r1, 4;
	add.s64 	%rd8, %rd6, %rd7;
	add.s64 	%rd9, %rd5, %rd7;
	ld.global.f32 	%f1, [%rd9];
	ld.global.f32 	%f2, [%rd8];
	div.approx.ftz.f32 	%f3, %f2, %f1;
	add.s64 	%rd10, %rd4, %rd7;
	st.global.f32 	[%rd10], %f3;

BB6_2:
	ret;
}

.visible .entry vectorMul(
	.param .u64 vectorMul_param_0,
	.param .u64 vectorMul_param_1,
	.param .u64 vectorMul_param_2,
	.param .u32 vectorMul_param_3
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<4>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<11>;


	ld.param.u64 	%rd1, [vectorMul_param_0];
	ld.param.u64 	%rd2, [vectorMul_param_1];
	ld.param.u64 	%rd3, [vectorMul_param_2];
	ld.param.u32 	%r2, [vectorMul_param_3];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB7_2;

	cvta.to.global.u64 	%rd4, %rd3;
	cvta.to.global.u64 	%rd5, %rd2;
	cvta.to.global.u64 	%rd6, %rd1;
	mul.wide.s32 	%rd7, %r1, 4;
	add.s64 	%rd8, %rd6, %rd7;
	add.s64 	%rd9, %rd5, %rd7;
	ld.global.f32 	%f1, [%rd9];
	ld.global.f32 	%f2, [%rd8];
	mul.ftz.f32 	%f3, %f2, %f1;
	add.s64 	%rd10, %rd4, %rd7;
	st.global.f32 	[%rd10], %f3;

BB7_2:
	ret;
}

.visible .entry vectorMax(
	.param .u64 vectorMax_param_0,
	.param .u64 vectorMax_param_1,
	.param .f32 vectorMax_param_2,
	.param .u32 vectorMax_param_3
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<4>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [vectorMax_param_0];
	ld.param.u64 	%rd2, [vectorMax_param_1];
	ld.param.f32 	%f1, [vectorMax_param_2];
	ld.param.u32 	%r2, [vectorMax_param_3];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB8_2;

	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f2, [%rd6];
	max.ftz.f32 	%f3, %f2, %f1;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;

BB8_2:
	ret;
}

.visible .entry vectorMin(
	.param .u64 vectorMin_param_0,
	.param .u64 vectorMin_param_1,
	.param .f32 vectorMin_param_2,
	.param .u32 vectorMin_param_3
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<4>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [vectorMin_param_0];
	ld.param.u64 	%rd2, [vectorMin_param_1];
	ld.param.f32 	%f1, [vectorMin_param_2];
	ld.param.u32 	%r2, [vectorMin_param_3];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB9_2;

	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f2, [%rd6];
	min.ftz.f32 	%f3, %f2, %f1;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;

BB9_2:
	ret;
}

.visible .entry vectorPow(
	.param .u64 vectorPow_param_0,
	.param .u64 vectorPow_param_1,
	.param .f32 vectorPow_param_2,
	.param .u32 vectorPow_param_3
)
{
	.reg .pred 	%p<23>;
	.reg .f32 	%f<4>;
	.reg .s32 	%r<26>;
	.reg .s64 	%rd<10>;
	.reg .f64 	%fd<20>;


	ld.param.u64 	%rd2, [vectorPow_param_0];
	ld.param.u64 	%rd3, [vectorPow_param_1];
	ld.param.f32 	%f2, [vectorPow_param_2];
	ld.param.u32 	%r4, [vectorPow_param_3];
	mov.u32 	%r5, %ctaid.x;
	mov.u32 	%r6, %ntid.x;
	mov.u32 	%r7, %tid.x;
	mad.lo.s32 	%r1, %r5, %r6, %r7;
	setp.ge.s32	%p1, %r1, %r4;
	@%p1 bra 	BB10_18;

	cvta.to.global.u64 	%rd4, %rd2;
	cvt.s64.s32	%rd1, %r1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f1, [%rd6];
	cvt.ftz.f64.f32	%fd1, %f1;
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r2}, %fd1;
	}
	cvt.ftz.f64.f32	%fd2, %f2;
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r3}, %fd2;
	}
	mul.f64 	%fd13, %fd2, 0d3FE0000000000000;
	cvt.rzi.f64.f64	%fd14, %fd13;
	fma.rn.f64 	%fd15, %fd14, 0dC000000000000000, %fd2;
	abs.f64 	%fd3, %fd15;
	setp.eq.ftz.f32	%p2, %f1, 0f3F800000;
	setp.eq.ftz.f32	%p3, %f2, 0f00000000;
	or.pred  	%p4, %p2, %p3;
	@!%p4 bra 	BB10_3;
	bra.uni 	BB10_2;

BB10_2:
	mov.f64 	%fd19, 0d3FF0000000000000;
	bra.uni 	BB10_17;

BB10_3:
	abs.f64 	%fd4, %fd1;
	setp.gtu.f64	%p5, %fd4, 0d7FF0000000000000;
	@%p5 bra 	BB10_16;

	abs.f64 	%fd5, %fd2;
	setp.gtu.f64	%p6, %fd5, 0d7FF0000000000000;
	@%p6 bra 	BB10_16;

	setp.eq.f64	%p7, %fd5, 0d7FF0000000000000;
	@%p7 bra 	BB10_15;

	setp.eq.f64	%p8, %fd4, 0d7FF0000000000000;
	@%p8 bra 	BB10_14;

	setp.eq.ftz.f32	%p9, %f1, 0f00000000;
	@%p9 bra 	BB10_13;

	setp.gt.s32	%p10, %r2, -1;
	@%p10 bra 	BB10_11;

	cvt.rzi.f64.f64	%fd16, %fd2;
	setp.eq.f64	%p11, %fd16, %fd2;
	@%p11 bra 	BB10_11;

	mov.f64 	%fd19, 0dFFF8000000000000;
	bra.uni 	BB10_17;

BB10_11:
	setp.lt.s32	%p12, %r2, 0;
	// Callseq Start 0
	{
	.reg .b32 temp_param_reg;
	.param .b64 param0;
	st.param.f64	[param0+0], %fd4;
	.param .b64 param1;
	st.param.f64	[param1+0], %fd2;
	.param .b64 retval0;
	call.uni (retval0), 
	__internal_accurate_pow, 
	(
	param0, 
	param1
	);
	ld.param.f64	%fd19, [retval0+0];
	}
	// Callseq End 0
	setp.eq.f64	%p13, %fd3, 0d3FF0000000000000;
	and.pred  	%p14, %p12, %p13;
	@!%p14 bra 	BB10_17;
	bra.uni 	BB10_12;

BB10_12:
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r8}, %fd19;
	}
	xor.b32  	%r9, %r8, -2147483648;
	{
	.reg .b32 %temp; 
	mov.b64 	{%r10, %temp}, %fd19;
	}
	mov.b64 	%fd19, {%r10, %r9};
	bra.uni 	BB10_17;

BB10_13:
	setp.eq.f64	%p15, %fd3, 0d3FF0000000000000;
	selp.b32	%r11, %r2, 0, %p15;
	mov.u32 	%r12, 0;
	or.b32  	%r13, %r11, 2146435072;
	setp.lt.s32	%p16, %r3, 0;
	selp.b32	%r14, %r13, %r11, %p16;
	mov.b64 	%fd19, {%r12, %r14};
	bra.uni 	BB10_17;

BB10_14:
	setp.eq.f64	%p17, %fd3, 0d3FF0000000000000;
	shr.s32 	%r15, %r3, 31;
	and.b32  	%r16, %r15, -2146435072;
	add.s32 	%r17, %r16, 2146435072;
	setp.lt.s32	%p18, %r2, 0;
	mov.u32 	%r18, 0;
	and.pred  	%p19, %p18, %p17;
	or.b32  	%r19, %r17, -2147483648;
	selp.b32	%r20, %r19, %r17, %p19;
	mov.b64 	%fd19, {%r18, %r20};
	bra.uni 	BB10_17;

BB10_15:
	setp.gt.f64	%p20, %fd4, 0d3FF0000000000000;
	selp.b32	%r21, 2146435072, 0, %p20;
	mov.u32 	%r22, 0;
	xor.b32  	%r23, %r21, 2146435072;
	setp.lt.s32	%p21, %r3, 0;
	selp.b32	%r24, %r23, %r21, %p21;
	setp.eq.ftz.f32	%p22, %f1, 0fBF800000;
	selp.b32	%r25, 1072693248, %r24, %p22;
	mov.b64 	%fd19, {%r22, %r25};
	bra.uni 	BB10_17;

BB10_16:
	add.f64 	%fd19, %fd1, %fd2;

BB10_17:
	cvta.to.global.u64 	%rd7, %rd3;
	shl.b64 	%rd8, %rd1, 2;
	add.s64 	%rd9, %rd7, %rd8;
	cvt.rn.ftz.f32.f64	%f3, %fd19;
	st.global.f32 	[%rd9], %f3;

BB10_18:
	ret;
}

.visible .entry vectorSqr(
	.param .u64 vectorSqr_param_0,
	.param .u64 vectorSqr_param_1,
	.param .u32 vectorSqr_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<3>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [vectorSqr_param_0];
	ld.param.u64 	%rd2, [vectorSqr_param_1];
	ld.param.u32 	%r2, [vectorSqr_param_2];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB11_2;

	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f1, [%rd6];
	mul.ftz.f32 	%f2, %f1, %f1;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f2;

BB11_2:
	ret;
}

.visible .entry vectorSqrt(
	.param .u64 vectorSqrt_param_0,
	.param .u64 vectorSqrt_param_1,
	.param .u32 vectorSqrt_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<3>;
	.reg .s32 	%r<6>;
	.reg .s64 	%rd<8>;


	ld.param.u64 	%rd1, [vectorSqrt_param_0];
	ld.param.u64 	%rd2, [vectorSqrt_param_1];
	ld.param.u32 	%r2, [vectorSqrt_param_2];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB12_2;

	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f1, [%rd6];
	sqrt.approx.ftz.f32 	%f2, %f1;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f2;

BB12_2:
	ret;
}

.func  (.param .b64 func_retval0) __internal_accurate_pow(
	.param .b64 __internal_accurate_pow_param_0,
	.param .b64 __internal_accurate_pow_param_1
)
{
	.reg .pred 	%p<11>;
	.reg .f32 	%f<5>;
	.reg .s32 	%r<39>;
	.reg .f64 	%fd<144>;


	ld.param.f64 	%fd14, [__internal_accurate_pow_param_0];
	ld.param.f64 	%fd15, [__internal_accurate_pow_param_1];
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r35}, %fd14;
	}
	{
	.reg .b32 %temp; 
	mov.b64 	{%r34, %temp}, %fd14;
	}
	shr.u32 	%r18, %r35, 20;
	and.b32  	%r36, %r18, 2047;
	setp.ne.s32	%p1, %r36, 0;
	@%p1 bra 	BB13_2;

	mul.f64 	%fd16, %fd14, 0d4350000000000000;
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r35}, %fd16;
	}
	{
	.reg .b32 %temp; 
	mov.b64 	{%r34, %temp}, %fd16;
	}
	shr.u32 	%r19, %r35, 20;
	and.b32  	%r20, %r19, 2047;
	add.s32 	%r36, %r20, -54;

BB13_2:
	add.s32 	%r37, %r36, -1023;
	and.b32  	%r21, %r35, -2146435073;
	or.b32  	%r22, %r21, 1072693248;
	mov.b64 	%fd141, {%r34, %r22};
	setp.lt.u32	%p2, %r22, 1073127583;
	@%p2 bra 	BB13_4;

	{
	.reg .b32 %temp; 
	mov.b64 	{%r23, %temp}, %fd141;
	}
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r24}, %fd141;
	}
	add.s32 	%r25, %r24, -1048576;
	mov.b64 	%fd141, {%r23, %r25};
	add.s32 	%r37, %r36, -1022;

BB13_4:
	add.f64 	%fd17, %fd141, 0d3FF0000000000000;
	mov.f64 	%fd19, 0d3FF0000000000000;
	// inline asm
	cvt.rn.f32.f64     %f1,%fd17;
	// inline asm
	// inline asm
	rcp.approx.ftz.f32 %f2,%f1;
	// inline asm
	// inline asm
	cvt.f64.f32        %fd18,%f2;
	// inline asm
	neg.f64 	%fd20, %fd17;
	fma.rn.f64 	%fd21, %fd20, %fd18, %fd19;
	fma.rn.f64 	%fd22, %fd21, %fd21, %fd21;
	fma.rn.f64 	%fd23, %fd22, %fd18, %fd18;
	add.f64 	%fd24, %fd141, 0dBFF0000000000000;
	mul.f64 	%fd25, %fd24, %fd23;
	fma.rn.f64 	%fd26, %fd24, %fd23, %fd25;
	mul.f64 	%fd27, %fd26, %fd26;
	mov.f64 	%fd28, 0d3ED0F5D241AD3B5A;
	mov.f64 	%fd29, 0d3EB0F5FF7D2CAFE2;
	fma.rn.f64 	%fd30, %fd29, %fd27, %fd28;
	mov.f64 	%fd31, 0d3EF3B20A75488A3F;
	fma.rn.f64 	%fd32, %fd30, %fd27, %fd31;
	mov.f64 	%fd33, 0d3F1745CDE4FAECD5;
	fma.rn.f64 	%fd34, %fd32, %fd27, %fd33;
	mov.f64 	%fd35, 0d3F3C71C7258A578B;
	fma.rn.f64 	%fd36, %fd34, %fd27, %fd35;
	mov.f64 	%fd37, 0d3F6249249242B910;
	fma.rn.f64 	%fd38, %fd36, %fd27, %fd37;
	mov.f64 	%fd39, 0d3F89999999999DFB;
	fma.rn.f64 	%fd40, %fd38, %fd27, %fd39;
	sub.f64 	%fd41, %fd24, %fd26;
	add.f64 	%fd42, %fd41, %fd41;
	neg.f64 	%fd43, %fd26;
	fma.rn.f64 	%fd44, %fd43, %fd24, %fd42;
	mul.f64 	%fd45, %fd23, %fd44;
	fma.rn.f64 	%fd46, %fd40, %fd27, 0d3FB5555555555555;
	mov.f64 	%fd47, 0d3FB5555555555555;
	sub.f64 	%fd48, %fd47, %fd46;
	fma.rn.f64 	%fd49, %fd40, %fd27, %fd48;
	add.f64 	%fd50, %fd49, 0d0000000000000000;
	add.f64 	%fd51, %fd50, 0dBC46A4CB00B9E7B0;
	add.f64 	%fd52, %fd46, %fd51;
	sub.f64 	%fd53, %fd46, %fd52;
	add.f64 	%fd54, %fd53, %fd51;
	mul.rn.f64 	%fd55, %fd52, %fd26;
	neg.f64 	%fd56, %fd55;
	fma.rn.f64 	%fd57, %fd52, %fd26, %fd56;
	fma.rn.f64 	%fd58, %fd52, %fd45, %fd57;
	fma.rn.f64 	%fd59, %fd54, %fd26, %fd58;
	add.f64 	%fd60, %fd55, %fd59;
	sub.f64 	%fd61, %fd55, %fd60;
	add.f64 	%fd62, %fd61, %fd59;
	mul.rn.f64 	%fd63, %fd60, %fd26;
	neg.f64 	%fd64, %fd63;
	fma.rn.f64 	%fd65, %fd60, %fd26, %fd64;
	fma.rn.f64 	%fd66, %fd60, %fd45, %fd65;
	fma.rn.f64 	%fd67, %fd62, %fd26, %fd66;
	add.f64 	%fd68, %fd63, %fd67;
	sub.f64 	%fd69, %fd63, %fd68;
	add.f64 	%fd70, %fd69, %fd67;
	mul.rn.f64 	%fd71, %fd68, %fd26;
	neg.f64 	%fd72, %fd71;
	fma.rn.f64 	%fd73, %fd68, %fd26, %fd72;
	fma.rn.f64 	%fd74, %fd68, %fd45, %fd73;
	fma.rn.f64 	%fd75, %fd70, %fd26, %fd74;
	add.f64 	%fd76, %fd71, %fd75;
	sub.f64 	%fd77, %fd71, %fd76;
	add.f64 	%fd78, %fd77, %fd75;
	add.f64 	%fd79, %fd26, %fd76;
	sub.f64 	%fd80, %fd26, %fd79;
	add.f64 	%fd81, %fd80, %fd76;
	add.f64 	%fd82, %fd81, %fd78;
	add.f64 	%fd83, %fd82, %fd45;
	add.f64 	%fd84, %fd79, %fd83;
	sub.f64 	%fd85, %fd79, %fd84;
	add.f64 	%fd86, %fd85, %fd83;
	cvt.rn.f64.s32	%fd87, %r37;
	mov.f64 	%fd88, 0d3FE62E42FEFA3000;
	mul.rn.f64 	%fd89, %fd87, %fd88;
	mov.f64 	%fd90, 0d3D53DE6AF278ECE6;
	mul.rn.f64 	%fd91, %fd87, %fd90;
	add.f64 	%fd92, %fd89, %fd84;
	sub.f64 	%fd93, %fd89, %fd92;
	add.f64 	%fd94, %fd93, %fd84;
	add.f64 	%fd95, %fd94, %fd86;
	add.f64 	%fd96, %fd95, %fd91;
	add.f64 	%fd97, %fd92, %fd96;
	sub.f64 	%fd98, %fd92, %fd97;
	add.f64 	%fd99, %fd98, %fd96;
	abs.f64 	%fd100, %fd15;
	setp.gt.f64	%p3, %fd100, 0d7F0D2A1BE4048F90;
	mul.f64 	%fd101, %fd15, 0d3F20000000000000;
	selp.f64	%fd102, %fd101, %fd15, %p3;
	mul.rn.f64 	%fd103, %fd97, %fd102;
	neg.f64 	%fd104, %fd103;
	fma.rn.f64 	%fd105, %fd97, %fd102, %fd104;
	fma.rn.f64 	%fd106, %fd99, %fd102, %fd105;
	add.f64 	%fd4, %fd103, %fd106;
	sub.f64 	%fd107, %fd103, %fd4;
	add.f64 	%fd5, %fd107, %fd106;
	{
	.reg .b32 %temp; 
	mov.b64 	{%temp, %r13}, %fd4;
	}
	setp.lt.u32	%p4, %r13, 1082535491;
	setp.lt.s32	%p5, %r13, -1064875759;
	or.pred  	%p6, %p4, %p5;
	@%p6 bra 	BB13_6;

	setp.lt.s32	%p7, %r13, 0;
	selp.f64	%fd108, 0d0000000000000000, 0d7FF0000000000000, %p7;
	abs.f64 	%fd109, %fd4;
	setp.gtu.f64	%p8, %fd109, 0d7FF0000000000000;
	add.f64 	%fd110, %fd4, %fd4;
	selp.f64	%fd143, %fd110, %fd108, %p8;
	bra.uni 	BB13_10;

BB13_6:
	mul.f64 	%fd111, %fd4, 0d3FF71547652B82FE;
	cvt.rni.f64.f64	%fd112, %fd111;
	cvt.rzi.s32.f64	%r14, %fd112;
	mov.f64 	%fd113, 0dBFE62E42FEFA39EF;
	fma.rn.f64 	%fd114, %fd112, %fd113, %fd4;
	mov.f64 	%fd115, 0dBC7ABC9E3B39803F;
	fma.rn.f64 	%fd116, %fd112, %fd115, %fd114;
	mov.f64 	%fd117, 0d3E928A27E30F5561;
	mov.f64 	%fd118, 0d3E5AE6449C0686C0;
	fma.rn.f64 	%fd119, %fd118, %fd116, %fd117;
	mov.f64 	%fd120, 0d3EC71DE8E6486D6B;
	fma.rn.f64 	%fd121, %fd119, %fd116, %fd120;
	mov.f64 	%fd122, 0d3EFA019A6B2464C5;
	fma.rn.f64 	%fd123, %fd121, %fd116, %fd122;
	mov.f64 	%fd124, 0d3F2A01A0171064A5;
	fma.rn.f64 	%fd125, %fd123, %fd116, %fd124;
	mov.f64 	%fd126, 0d3F56C16C17F29C8D;
	fma.rn.f64 	%fd127, %fd125, %fd116, %fd126;
	mov.f64 	%fd128, 0d3F8111111111A24E;
	fma.rn.f64 	%fd129, %fd127, %fd116, %fd128;
	mov.f64 	%fd130, 0d3FA555555555211D;
	fma.rn.f64 	%fd131, %fd129, %fd116, %fd130;
	mov.f64 	%fd132, 0d3FC5555555555530;
	fma.rn.f64 	%fd133, %fd131, %fd116, %fd132;
	mov.f64 	%fd134, 0d3FE0000000000005;
	fma.rn.f64 	%fd135, %fd133, %fd116, %fd134;
	fma.rn.f64 	%fd137, %fd135, %fd116, %fd19;
	fma.rn.f64 	%fd142, %fd137, %fd116, %fd19;
	abs.s32 	%r26, %r14;
	setp.lt.s32	%p9, %r26, 1023;
	@%p9 bra 	BB13_8;

	add.s32 	%r27, %r14, 2046;
	shl.b32 	%r28, %r27, 19;
	and.b32  	%r29, %r28, -1048576;
	shl.b32 	%r30, %r27, 20;
	sub.s32 	%r38, %r30, %r29;
	mov.u32 	%r31, 0;
	mov.b64 	%fd138, {%r31, %r29};
	mul.f64 	%fd142, %fd142, %fd138;
	bra.uni 	BB13_9;

BB13_8:
	shl.b32 	%r32, %r14, 20;
	add.s32 	%r38, %r32, 1072693248;

BB13_9:
	mov.u32 	%r33, 0;
	mov.b64 	%fd139, {%r33, %r38};
	mul.f64 	%fd143, %fd142, %fd139;

BB13_10:
	abs.f64 	%fd140, %fd143;
	setp.eq.f64	%p10, %fd140, 0d7FF0000000000000;
	@%p10 bra 	BB13_12;

	fma.rn.f64 	%fd143, %fd143, %fd5, %fd143;

BB13_12:
	st.param.f64	[func_retval0+0], %fd143;
	ret;
}