kernel_1.ptx

//
// Generated by LLVM NVPTX Back-End
//

.version 6.0
.target sm_61
.address_size 64

.extern .func  (.param .b32 func_retval0) vprintf
(
	.param .b64 vprintf_param_0,
	.param .b64 vprintf_param_1
)
;
.func ptx_report_exception
(
	.param .b64 ptx_report_exception_param_0
)
;
.global .align 1 .b8 __unnamed_1[14] = {100, 105, 102, 102, 32, 61, 32, 37, 46, 49, 54, 101, 10, 0};
.global .align 1 .b8 exception[10] = {101, 120, 99, 101, 112, 116, 105, 111, 110, 0};
.global .align 1 .b8 __unnamed_2[108] = {69, 82, 82, 79, 82, 58, 32, 97, 32, 37, 115, 32, 119, 97, 115, 32, 116, 104, 114, 111, 119, 110, 32, 100, 117, 114, 105, 110, 103, 32, 107, 101, 114, 110, 101, 108, 32, 101, 120, 101, 99, 117, 116, 105, 111, 110, 46, 10, 32, 32, 32, 32, 32, 32, 32, 82, 117, 110, 32, 74, 117, 108, 105, 97, 32, 111, 110, 32, 100, 101, 98, 117, 103, 32, 108, 101, 118, 101, 108, 32, 50, 32, 102, 111, 114, 32, 100, 101, 118, 105, 99, 101, 32, 115, 116, 97, 99, 107, 32, 116, 114, 97, 99, 101, 115, 46, 10, 0};
                                        // -- Begin function julia_overdub_16970
                                        // @julia_overdub_16970
.func julia_overdub_16970()
{
	.reg .b64 	%rd<3>;

// %bb.0:                               // %top
	mov.u64 	%rd1, exception;
	cvta.global.u64 	%rd2, %rd1;
	{ // callseq 0, 0
	.reg .b32 temp_param_reg;
	.param .b64 param0;
	st.param.b64 	[param0+0], %rd2;
	call.uni 
	ptx_report_exception, 
	(
	param0
	);
	} // callseq 0
	// begin inline asm
	trap;
	// end inline asm
	ret;
}
                                        // -- End function
	// .globl	ptxcall_kernel_3 // -- Begin function ptxcall_kernel_3
.visible .entry ptxcall_kernel_3(
	.param .align 8 .b8 ptxcall_kernel_3_param_0[16],
	.param .align 8 .b8 ptxcall_kernel_3_param_1[16]
)                                       // @ptxcall_kernel_3
{
	.local .align 8 .b8 	__local_depot1[56];
	.reg .b64 	%SP;
	.reg .b64 	%SPL;
	.reg .pred 	%p<4>;
	.reg .b32 	%r<2>;
	.reg .f64 	%fd<7>;
	.reg .b64 	%rd<20>;

// %bb.0:                               // %entry
	mov.u64 	%SPL, __local_depot1;
	cvta.local.u64 	%SP, %SPL;
	ld.param.u64 	%rd4, [ptxcall_kernel_3_param_0+8];
	ld.param.u64 	%rd7, [ptxcall_kernel_3_param_0];
	ld.param.u64 	%rd6, [ptxcall_kernel_3_param_1+8];
	ld.param.u64 	%rd5, [ptxcall_kernel_3_param_1];
	add.u64 	%rd11, %SPL, 8;
	add.u64 	%rd3, %SPL, 16;
	add.u64 	%rd14, %SPL, 24;
	st.local.u64 	[%rd14], %rd7;
	st.local.u64 	[%rd14+8], %rd4;
	st.local.u64 	[%rd14+16], %rd5;
	st.local.u64 	[%rd14+24], %rd6;
	mov.u64 	%rd15, 1;
	st.local.u64 	[%rd11], %rd15;
	setp.lt.s64 	%p1, %rd7, 1;
	@%p1 bra 	LBB1_5;
	bra.uni 	LBB1_1;
LBB1_5:                                 // %L14.i
	{ // callseq 1, 0
	.reg .b32 temp_param_reg;
	call.uni 
	julia_overdub_16970, 
	(
	);
	} // callseq 1
	// begin inline asm
	trap;
	// end inline asm
LBB1_1:                                 // %L13.i
	cvta.to.global.u64 	%rd16, %rd4;
	ld.global.f64 	%fd1, [%rd16];
	st.local.u64 	[%rd3], %rd15;
	setp.lt.s64 	%p2, %rd5, 1;
	@%p2 bra 	LBB1_6;
	bra.uni 	LBB1_2;
LBB1_6:                                 // %L31.i
	{ // callseq 2, 0
	.reg .b32 temp_param_reg;
	call.uni 
	julia_overdub_16970, 
	(
	);
	} // callseq 2
	// begin inline asm
	trap;
	// end inline asm
LBB1_2:                                 // %L30.i
	cvta.to.global.u64 	%rd18, %rd6;
	ld.global.f64 	%fd3, [%rd18];
	mul.f64 	%fd4, %fd1, %fd3;
	neg.f64 	%fd5, %fd4;
	fma.rn.f64 	%fd2, %fd1, %fd3, %fd5;
	abs.f64 	%fd6, %fd2;
	setp.leu.f64 	%p3, %fd6, 0d3C9CD2B297D889BC;
	@%p3 bra 	LBB1_4;
// %bb.3:                               // %L45.i
	mov.u64 	%rd8, __unnamed_1;
	add.u64 	%rd9, %SP, 0;
	cvta.global.u64 	%rd1, %rd8;
	add.u64 	%rd2, %SPL, 0;
	st.local.f64 	[%rd2], %fd2;
	{ // callseq 3, 0
	.reg .b32 temp_param_reg;
	.param .b64 param0;
	st.param.b64 	[param0+0], %rd1;
	.param .b64 param1;
	st.param.b64 	[param1+0], %rd9;
	.param .b32 retval0;
	call.uni (retval0), 
	vprintf, 
	(
	param0, 
	param1
	);
	ld.param.b32 	%r1, [retval0+0];
	} // callseq 3
LBB1_4:                                 // %julia_kernel_3.exit
	ret;
}
                                        // -- End function
.func ptx_report_exception(
	.param .b64 ptx_report_exception_param_0
)                                       // -- Begin function ptx_report_exception
                                        // @ptx_report_exception
{
	.local .align 8 .b8 	__local_depot2[8];
	.reg .b64 	%SP;
	.reg .b64 	%SPL;
	.reg .b32 	%r<2>;
	.reg .b64 	%rd<6>;

// %bb.0:                               // %top
	mov.u64 	%SPL, __local_depot2;
	cvta.local.u64 	%SP, %SPL;
	ld.param.u64 	%rd1, [ptx_report_exception_param_0];
	mov.u64 	%rd2, __unnamed_2;
	cvta.global.u64 	%rd3, %rd2;
	add.u64 	%rd4, %SP, 0;
	add.u64 	%rd5, %SPL, 0;
	st.local.u64 	[%rd5], %rd1;
	{ // callseq 4, 0
	.reg .b32 temp_param_reg;
	.param .b64 param0;
	st.param.b64 	[param0+0], %rd3;
	.param .b64 param1;
	st.param.b64 	[param1+0], %rd4;
	.param .b32 retval0;
	call.uni (retval0), 
	vprintf, 
	(
	param0, 
	param1
	);
	ld.param.b32 	%r1, [retval0+0];
	} // callseq 4
	ret;
}
                                        // -- End function