.visible .entry main_graph_async_dispatch_2_elementwise_broadcast_D_i64(
	.param .u64 .ptr .global .align 16 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_0,
	.param .u32 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_1,
	.param .u32 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_2,
	.param .u32 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_3,
	.param .u32 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_4,
	.param .u32 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_5,
	.param .u32 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_6,
	.param .u32 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_7,
	.param .u32 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_8,
	.param .u32 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_9,
	.param .u32 main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_10
)
.maxntid 32, 1, 1                       // @main_graph_async_dispatch_2_elementwise_broadcast_D_i64
{
	.reg .pred 	%p<7>;
	.reg .b32 	%r<9>;
	.reg .b64 	%rd<47>;

// %bb.0:
	ld.param.b32 	%rd12, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_3];
	ld.param.b32 	%rd18, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_4];
	shl.b64 	%rd19, %rd18, 32;
	or.b64 	%rd1, %rd19, %rd12;
	setp.eq.b64 	%p1, %rd1, 0;
	max.u64 	%rd28, %rd1, 1;
	add.s64 	%rd29, %rd28, -1;
	shr.u64 	%rd30, %rd29, 5;
	add.s64 	%rd31, %rd30, 1;
	selp.b64 	%rd6, 0, %rd31, %p1;
	mov.u32 	%r3, %ctaid.x;
	cvt.u64.u32 	%rd45, %r3;
	setp.le.u64 	%p2, %rd6, %rd45;
	@%p2 bra 	$L__BB2_9;
// %bb.1:                               // %.lr.ph3.preheader
	ld.param.b64 	%rd10, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_0];
	mov.u32 	%r1, %tid.x;
	ld.param.b32 	%rd11, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_1];
	ld.param.b32 	%rd13, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_2];
	shl.b64 	%rd14, %rd13, 32;
	or.b64 	%rd15, %rd14, %rd11;
	ld.param.b32 	%rd16, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_5];
	ld.param.b32 	%rd17, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_7];
	ld.param.b32 	%rd20, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_9];
	ld.param.b32 	%rd21, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_6];
	shl.b64 	%rd22, %rd21, 32;
	or.b64 	%rd2, %rd22, %rd16;
	ld.param.b32 	%rd23, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_8];
	shl.b64 	%rd24, %rd23, 32;
	or.b64 	%rd3, %rd24, %rd17;
	ld.param.b32 	%rd25, [main_graph_async_dispatch_2_elementwise_broadcast_D_i64_param_10];
	shl.b64 	%rd26, %rd25, 32;
	or.b64 	%rd4, %rd26, %rd20;
	and.b64 	%rd27, %rd15, 2305843009213693888;
	add.s64 	%rd5, %rd10, %rd27;
	mov.u32 	%r4, %nctaid.x;
	cvt.u64.u32 	%rd7, %r4;
	cvt.u32.u64 	%r5, %rd2;
	bra.uni 	$L__BB2_2;
$L__BB2_8:                              // %._crit_edge
                                        //   in Loop: Header=BB2_2 Depth=1
	add.s64 	%rd45, %rd45, %rd7;
	setp.lt.u64 	%p6, %rd45, %rd6;
	@!%p6 bra 	$L__BB2_9;
$L__BB2_2:                              // %.lr.ph3
                                        // =>This Loop Header: Depth=1
                                        //     Child Loop BB2_4 Depth 2
	shl.b64 	%rd8, %rd45, 5;
	sub.s64 	%rd32, %rd1, %rd8;
	min.s64 	%rd33, %rd32, 32;
	cvt.u32.u64 	%r2, %rd33;
	setp.ge.s32 	%p3, %r1, %r2;
	@%p3 bra 	$L__BB2_8;
// %bb.3:                               // %.lr.ph.preheader
                                        //   in Loop: Header=BB2_2 Depth=1
	mov.b32 	%r8, %r1;
	bra.uni 	$L__BB2_4;
$L__BB2_6:                              //   in Loop: Header=BB2_4 Depth=2
	div.s64 	%rd46, %rd9, %rd2;
$L__BB2_7:                              // %.lr.ph.split
                                        //   in Loop: Header=BB2_4 Depth=2
	mul.lo.s64 	%rd37, %rd46, %rd2;
	sub.s64 	%rd38, %rd9, %rd37;
	mad.lo.s64 	%rd39, %rd46, %rd3, %rd38;
	shr.s64 	%rd40, %rd39, 63;
	and.b64 	%rd41, %rd40, %rd4;
	add.s64 	%rd42, %rd41, %rd39;
	shl.b64 	%rd43, %rd9, 3;
	add.s64 	%rd44, %rd5, %rd43;
	st.global.b64 	[%rd44], %rd42;
	add.s32 	%r8, %r8, 32;
	setp.lt.s32 	%p5, %r8, %r2;
	@!%p5 bra 	$L__BB2_8;
$L__BB2_4:                              // %.lr.ph
                                        //   Parent Loop BB2_2 Depth=1
                                        // =>  This Inner Loop Header: Depth=2
	cvt.u64.u32 	%rd34, %r8;
	add.s64 	%rd9, %rd8, %rd34;
	or.b64 	%rd35, %rd9, %rd2;
	and.b64 	%rd36, %rd35, -4294967296;
	setp.ne.b64 	%p4, %rd36, 0;
	@%p4 bra 	$L__BB2_6;
// %bb.5:                               //   in Loop: Header=BB2_4 Depth=2
	cvt.u32.u64 	%r6, %rd9;
	div.u32 	%r7, %r6, %r5;
	cvt.u64.u32 	%rd46, %r7;
	bra.uni 	$L__BB2_7;
$L__BB2_9:                              // %._crit_edge4
	ret;
                                        // -- End function
}