//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-35583870
// Cuda compilation tools, release 12.8, V12.8.93
// Based on NVVM 7.0.1
//

.version 8.7
.target sm_90
.address_size 64

	// .globl	_Z7add_onePKfPfi
// _ZZ12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4_E3job has been demoted
// _ZZ12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4_E3job has been demoted

.visible .entry _Z7add_onePKfPfi(
	.param .u64 _Z7add_onePKfPfi_param_0,
	.param .u64 _Z7add_onePKfPfi_param_1,
	.param .u32 _Z7add_onePKfPfi_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<4>;
	.reg .b32 	%r<6>;
	.reg .b64 	%rd<8>;
	.loc	1 23 0


	ld.param.u64 	%rd1, [_Z7add_onePKfPfi_param_0];
	ld.param.u64 	%rd2, [_Z7add_onePKfPfi_param_1];
	ld.param.u32 	%r2, [_Z7add_onePKfPfi_param_2];
	.loc	1 23 30
	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 	$L__BB0_2;

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

$L__BB0_2:
	ret;

}
	// .globl	_Z6squarePKfPfi
.visible .entry _Z6squarePKfPfi(
	.param .u64 _Z6squarePKfPfi_param_0,
	.param .u64 _Z6squarePKfPfi_param_1,
	.param .u32 _Z6squarePKfPfi_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<3>;
	.reg .b32 	%r<6>;
	.reg .b64 	%rd<8>;
	.loc	1 24 0


	ld.param.u64 	%rd1, [_Z6squarePKfPfi_param_0];
	ld.param.u64 	%rd2, [_Z6squarePKfPfi_param_1];
	ld.param.u32 	%r2, [_Z6squarePKfPfi_param_2];
	.loc	1 24 30
	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 	$L__BB1_2;

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

$L__BB1_2:
	ret;

}
	// .globl	_Z10half_valuePKfPfi
.visible .entry _Z10half_valuePKfPfi(
	.param .u64 _Z10half_valuePKfPfi_param_0,
	.param .u64 _Z10half_valuePKfPfi_param_1,
	.param .u32 _Z10half_valuePKfPfi_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<4>;
	.reg .b32 	%r<6>;
	.reg .b64 	%rd<8>;
	.loc	1 25 0


	ld.param.u64 	%rd1, [_Z10half_valuePKfPfi_param_0];
	ld.param.u64 	%rd2, [_Z10half_valuePKfPfi_param_1];
	ld.param.u32 	%r2, [_Z10half_valuePKfPfi_param_2];
	.loc	1 25 30
	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 	$L__BB2_2;

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

$L__BB2_2:
	ret;

}
	// .globl	_Z15fused_pointwisePKfPfi
.visible .entry _Z15fused_pointwisePKfPfi(
	.param .u64 _Z15fused_pointwisePKfPfi_param_0,
	.param .u64 _Z15fused_pointwisePKfPfi_param_1,
	.param .u32 _Z15fused_pointwisePKfPfi_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<7>;
	.reg .b32 	%r<6>;
	.reg .b64 	%rd<8>;
	.loc	1 26 0


	ld.param.u64 	%rd1, [_Z15fused_pointwisePKfPfi_param_0];
	ld.param.u64 	%rd2, [_Z15fused_pointwisePKfPfi_param_1];
	ld.param.u32 	%r2, [_Z15fused_pointwisePKfPfi_param_2];
	.loc	1 26 30
	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 	$L__BB3_2;

	cvta.to.global.u64 	%rd3, %rd1;
	mul.wide.s32 	%rd4, %r1, 4;
	add.s64 	%rd5, %rd3, %rd4;
	ld.global.f32 	%f1, [%rd5];
	mov.f32 	%f2, 0f3F800000;
	add.rn.f32 	%f3, %f1, %f2;
	mul.rn.f32 	%f4, %f3, %f3;
	mov.f32 	%f5, 0f3F000000;
	mul.rn.f32 	%f6, %f4, %f5;
	cvta.to.global.u64 	%rd6, %rd2;
	add.s64 	%rd7, %rd6, %rd4;
	st.global.f32 	[%rd7], %f6;

$L__BB3_2:
	ret;

}
	// .globl	_Z8ragged_aILb1EEvPKjPVjPKiP5Stamp
.visible .entry _Z8ragged_aILb1EEvPKjPVjPKiP5Stamp(
	.param .u64 _Z8ragged_aILb1EEvPKjPVjPKiP5Stamp_param_0,
	.param .u64 _Z8ragged_aILb1EEvPKjPVjPKiP5Stamp_param_1,
	.param .u64 _Z8ragged_aILb1EEvPKjPVjPKiP5Stamp_param_2,
	.param .u64 _Z8ragged_aILb1EEvPKjPVjPKiP5Stamp_param_3
)
{
	.reg .pred 	%p<5>;
	.reg .b32 	%r<17>;
	.reg .b64 	%rd<20>;
	.loc	1 47 0


	ld.param.u64 	%rd3, [_Z8ragged_aILb1EEvPKjPVjPKiP5Stamp_param_0];
	ld.param.u64 	%rd4, [_Z8ragged_aILb1EEvPKjPVjPKiP5Stamp_param_1];
	ld.param.u64 	%rd5, [_Z8ragged_aILb1EEvPKjPVjPKiP5Stamp_param_2];
	ld.param.u64 	%rd6, [_Z8ragged_aILb1EEvPKjPVjPKiP5Stamp_param_3];
	.loc	1 47 51
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 47 51
	mov.u32 	%r1, %tid.x;
	setp.eq.s32 	%p1, %r1, 0;
	.loc	1 47 51
	mov.u32 	%r2, %ctaid.x;
	cvta.to.global.u64 	%rd7, %rd6;
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 47 51
	mul.wide.s32 	%rd8, %r2, 48;
	add.s64 	%rd1, %rd7, %rd8;
	@%p1 bra 	$L__BB4_1;
	bra.uni 	$L__BB4_2;

$L__BB4_1:
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 47 51
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 36 5
	// begin inline asm
	mov.u64 %rd9, %globaltimer;
	// end inline asm
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 47 51
	st.global.u64 	[%rd1], %rd9;
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 47 51
	.loc	1 30 34, function_name $L__info_string2, inlined_at 1 36 5
	// begin inline asm
	mov.u32 %r10, %smid;
	// end inline asm
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 47 51
	st.global.v2.u32 	[%rd1+32], {%r10, %r2};

$L__BB4_2:
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 47 51
	shl.b32 	%r11, %r2, 8;
	add.s32 	%r12, %r11, %r1;
	cvt.s64.s32 	%rd2, %r12;
	.loc	1 47 51
	cvta.to.global.u64 	%rd10, %rd3;
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 47 51
	mul.wide.s32 	%rd11, %r12, 4;
	add.s64 	%rd12, %rd10, %rd11;
	ld.global.u32 	%r16, [%rd12];
	.loc	1 47 51
	cvta.to.global.u64 	%rd13, %rd5;
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 47 51
	mul.wide.s32 	%rd14, %r2, 4;
	add.s64 	%rd15, %rd13, %rd14;
	ld.global.u32 	%r4, [%rd15];
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 47 51
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 37 5
	setp.lt.s32 	%p2, %r4, 1;
	@%p2 bra 	$L__BB4_5;

	.loc	1 0 5
	mov.u32 	%r14, 0;

$L__BB4_4:
	.pragma "nounroll";
	.loc	1 33 5
	mad.lo.s32 	%r16, %r16, 1664525, 1013904223;
	add.s32 	%r14, %r14, 1;
	setp.lt.s32 	%p3, %r14, %r4;
	@%p3 bra 	$L__BB4_4;

$L__BB4_5:
	.loc	1 47 51
	cvta.to.global.u64 	%rd16, %rd4;
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 47 51
	shl.b64 	%rd17, %rd2, 2;
	add.s64 	%rd18, %rd16, %rd17;
	st.volatile.global.u32 	[%rd18], %r16;
	.loc	1 38 5, function_name $L__info_string0, inlined_at 1 47 51
	bar.sync 	0;
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 47 51
	setp.ne.s32 	%p4, %r1, 0;
	.loc	1 39 5, function_name $L__info_string0, inlined_at 1 47 51
	@%p4 bra 	$L__BB4_7;

	.loc	1 39 5, function_name $L__info_string0, inlined_at 1 47 51
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 39 5
	// begin inline asm
	mov.u64 %rd19, %globaltimer;
	// end inline asm
	.loc	1 39 5, function_name $L__info_string0, inlined_at 1 47 51
	st.global.u64 	[%rd1+8], %rd19;

$L__BB4_7:
	.loc	1 47 51
	ret;

}
	// .globl	_Z8ragged_bILb1EEvPVKjPjP5Stamp
.visible .entry _Z8ragged_bILb1EEvPVKjPjP5Stamp(
	.param .u64 _Z8ragged_bILb1EEvPVKjPjP5Stamp_param_0,
	.param .u64 _Z8ragged_bILb1EEvPVKjPjP5Stamp_param_1,
	.param .u64 _Z8ragged_bILb1EEvPVKjPjP5Stamp_param_2
)
{
	.reg .pred 	%p<4>;
	.reg .b32 	%r<14>;
	.reg .b64 	%rd<17>;
	.loc	1 48 0


	ld.param.u64 	%rd4, [_Z8ragged_bILb1EEvPVKjPjP5Stamp_param_0];
	ld.param.u64 	%rd5, [_Z8ragged_bILb1EEvPVKjPjP5Stamp_param_1];
	ld.param.u64 	%rd6, [_Z8ragged_bILb1EEvPVKjPjP5Stamp_param_2];
	.loc	1 48 51
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 48 51
	mov.u32 	%r1, %tid.x;
	setp.ne.s32 	%p1, %r1, 0;
	.loc	1 48 51
	mov.u32 	%r2, %ctaid.x;
	cvta.to.global.u64 	%rd7, %rd6;
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 48 51
	mul.wide.s32 	%rd8, %r2, 48;
	add.s64 	%rd9, %rd7, %rd8;
	add.s64 	%rd1, %rd9, 16;
	@%p1 bra 	$L__BB5_2;

	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 48 51
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 42 5
	// begin inline asm
	mov.u64 %rd10, %globaltimer;
	// end inline asm
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 48 51
	st.global.u64 	[%rd1], %rd10;
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 48 51
	.loc	1 30 34, function_name $L__info_string2, inlined_at 1 42 5
	// begin inline asm
	mov.u32 %r8, %smid;
	// end inline asm
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 48 51
	st.global.v2.u32 	[%rd1+24], {%r8, %r2};

$L__BB5_2:
	.loc	1 43 5, function_name $L__info_string4, inlined_at 1 48 51
	shl.b32 	%r10, %r2, 8;
	add.s32 	%r11, %r10, %r1;
	cvt.s64.s32 	%rd2, %r11;
	.loc	1 48 51
	cvta.to.global.u64 	%rd11, %rd4;
	.loc	1 43 5, function_name $L__info_string4, inlined_at 1 48 51
	mul.wide.s32 	%rd12, %r11, 4;
	add.s64 	%rd13, %rd11, %rd12;
	ld.volatile.global.u32 	%r13, [%rd13];
	.loc	1 48 51
	cvta.to.global.u64 	%rd3, %rd5;
	mov.u32 	%r12, 0;

$L__BB5_3:
	.pragma "nounroll";
	.loc	1 43 5, function_name $L__info_string4, inlined_at 1 48 51
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 43 5
	mad.lo.s32 	%r13, %r13, 22695477, 1;
	add.s32 	%r12, %r12, 1;
	setp.ne.s32 	%p2, %r12, 512;
	@%p2 bra 	$L__BB5_3;

	.loc	1 43 5, function_name $L__info_string4, inlined_at 1 48 51
	shl.b64 	%rd14, %rd2, 2;
	add.s64 	%rd15, %rd3, %rd14;
	st.global.u32 	[%rd15], %r13;
	.loc	1 44 5, function_name $L__info_string4, inlined_at 1 48 51
	bar.sync 	0;
	.loc	1 45 5, function_name $L__info_string4, inlined_at 1 48 51
	@%p1 bra 	$L__BB5_6;

	.loc	1 45 5, function_name $L__info_string4, inlined_at 1 48 51
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 45 5
	// begin inline asm
	mov.u64 %rd16, %globaltimer;
	// end inline asm
	.loc	1 45 5, function_name $L__info_string4, inlined_at 1 48 51
	st.global.u64 	[%rd1+8], %rd16;

$L__BB5_6:
	.loc	1 48 51
	ret;

}
	// .globl	_Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp
.visible .entry _Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp(
	.param .u64 _Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp_param_0,
	.param .u64 _Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp_param_1,
	.param .u64 _Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp_param_2,
	.param .u64 _Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp_param_3,
	.param .u64 _Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp_param_4
)
{
	.reg .pred 	%p<8>;
	.reg .b32 	%r<26>;
	.reg .b64 	%rd<26>;
	.loc	1 49 0


	ld.param.u64 	%rd5, [_Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp_param_0];
	ld.param.u64 	%rd6, [_Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp_param_1];
	ld.param.u64 	%rd7, [_Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp_param_2];
	ld.param.u64 	%rd8, [_Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp_param_3];
	ld.param.u64 	%rd9, [_Z12ragged_fusedILb1EEvPKjPVjPjPKiP5Stamp_param_4];
	.loc	1 49 51
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 49 51
	mov.u32 	%r1, %tid.x;
	setp.eq.s32 	%p1, %r1, 0;
	.loc	1 49 51
	mov.u32 	%r2, %ctaid.x;
	cvta.to.global.u64 	%rd10, %rd9;
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 49 51
	mul.wide.s32 	%rd11, %r2, 48;
	add.s64 	%rd1, %rd10, %rd11;
	@%p1 bra 	$L__BB6_1;
	bra.uni 	$L__BB6_2;

$L__BB6_1:
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 49 51
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 36 5
	// begin inline asm
	mov.u64 %rd12, %globaltimer;
	// end inline asm
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 49 51
	st.global.u64 	[%rd1], %rd12;
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 49 51
	.loc	1 30 34, function_name $L__info_string2, inlined_at 1 36 5
	// begin inline asm
	mov.u32 %r15, %smid;
	// end inline asm
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 49 51
	st.global.v2.u32 	[%rd1+32], {%r15, %r2};

$L__BB6_2:
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 49 51
	shl.b32 	%r16, %r2, 8;
	add.s32 	%r17, %r16, %r1;
	cvt.s64.s32 	%rd2, %r17;
	.loc	1 49 51
	cvta.to.global.u64 	%rd13, %rd5;
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 49 51
	mul.wide.s32 	%rd14, %r17, 4;
	add.s64 	%rd15, %rd13, %rd14;
	ld.global.u32 	%r23, [%rd15];
	.loc	1 49 51
	cvta.to.global.u64 	%rd16, %rd8;
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 49 51
	mul.wide.s32 	%rd17, %r2, 4;
	add.s64 	%rd18, %rd16, %rd17;
	ld.global.u32 	%r4, [%rd18];
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 49 51
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 37 5
	setp.lt.s32 	%p2, %r4, 1;
	@%p2 bra 	$L__BB6_5;

	.loc	1 0 5
	mov.u32 	%r21, 0;

$L__BB6_4:
	.pragma "nounroll";
	.loc	1 33 5
	mad.lo.s32 	%r23, %r23, 1664525, 1013904223;
	add.s32 	%r21, %r21, 1;
	setp.lt.s32 	%p3, %r21, %r4;
	@%p3 bra 	$L__BB6_4;

$L__BB6_5:
	.loc	1 49 51
	cvta.to.global.u64 	%rd19, %rd6;
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 49 51
	shl.b64 	%rd20, %rd2, 2;
	add.s64 	%rd3, %rd19, %rd20;
	st.volatile.global.u32 	[%rd3], %r23;
	.loc	1 38 5, function_name $L__info_string0, inlined_at 1 49 51
	bar.sync 	0;
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 49 51
	setp.ne.s32 	%p4, %r1, 0;
	.loc	1 39 5, function_name $L__info_string0, inlined_at 1 49 51
	@%p4 bra 	$L__BB6_7;

	.loc	1 39 5, function_name $L__info_string0, inlined_at 1 49 51
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 39 5
	// begin inline asm
	mov.u64 %rd21, %globaltimer;
	// end inline asm
	.loc	1 39 5, function_name $L__info_string0, inlined_at 1 49 51
	st.global.u64 	[%rd1+8], %rd21;

$L__BB6_7:
	.loc	1 49 51
	bar.sync 	0;
	.loc	1 49 51
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 49 51
	@%p4 bra 	$L__BB6_9;

	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 49 51
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 42 5
	// begin inline asm
	mov.u64 %rd22, %globaltimer;
	// end inline asm
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 49 51
	st.global.u64 	[%rd1+16], %rd22;
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 49 51
	.loc	1 30 34, function_name $L__info_string2, inlined_at 1 42 5
	// begin inline asm
	mov.u32 %r19, %smid;
	// end inline asm
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 49 51
	st.global.v2.u32 	[%rd1+40], {%r19, %r2};

$L__BB6_9:
	.loc	1 43 5, function_name $L__info_string4, inlined_at 1 49 51
	ld.volatile.global.u32 	%r25, [%rd3];
	.loc	1 49 51
	cvta.to.global.u64 	%rd4, %rd7;
	mov.u32 	%r24, 0;

$L__BB6_10:
	.pragma "nounroll";
	.loc	1 43 5, function_name $L__info_string4, inlined_at 1 49 51
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 43 5
	mad.lo.s32 	%r25, %r25, 22695477, 1;
	add.s32 	%r24, %r24, 1;
	setp.ne.s32 	%p6, %r24, 512;
	@%p6 bra 	$L__BB6_10;

	.loc	1 43 5, function_name $L__info_string4, inlined_at 1 49 51
	add.s64 	%rd24, %rd4, %rd20;
	st.global.u32 	[%rd24], %r25;
	.loc	1 44 5, function_name $L__info_string4, inlined_at 1 49 51
	bar.sync 	0;
	.loc	1 45 5, function_name $L__info_string4, inlined_at 1 49 51
	@%p4 bra 	$L__BB6_13;

	.loc	1 45 5, function_name $L__info_string4, inlined_at 1 49 51
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 45 5
	// begin inline asm
	mov.u64 %rd25, %globaltimer;
	// end inline asm
	.loc	1 45 5, function_name $L__info_string4, inlined_at 1 49 51
	st.global.u64 	[%rd1+24], %rd25;

$L__BB6_13:
	.loc	1 49 51
	ret;

}
	// .globl	_Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4_
.visible .entry _Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4_(
	.param .u64 _Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_0,
	.param .u64 _Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_1,
	.param .u64 _Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_2,
	.param .u64 _Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_3,
	.param .u64 _Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_4,
	.param .u32 _Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_5,
	.param .u64 _Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_6
)
{
	.reg .pred 	%p<10>;
	.reg .b32 	%r<34>;
	.reg .b64 	%rd<36>;
	.loc	1 50 0
	// demoted variable
	.shared .align 4 .u32 _ZZ12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4_E3job;

	ld.param.u64 	%rd13, [_Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_0];
	ld.param.u64 	%rd14, [_Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_1];
	ld.param.u64 	%rd15, [_Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_2];
	ld.param.u64 	%rd16, [_Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_3];
	ld.param.u64 	%rd17, [_Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_4];
	ld.param.u32 	%r15, [_Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_5];
	ld.param.u64 	%rd18, [_Z12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4__param_6];
	.loc	1 52 5
	cvta.to.global.u64 	%rd1, %rd15;
	cvta.to.global.u64 	%rd2, %rd14;
	cvta.to.global.u64 	%rd3, %rd16;
	cvta.to.global.u64 	%rd4, %rd13;
	cvta.to.global.u64 	%rd5, %rd17;
	cvta.to.global.u64 	%rd6, %rd18;
	mov.u32 	%r1, %tid.x;
	.loc	1 53 9
	mov.u32 	%r2, %ctaid.x;
	bra.uni 	$L__BB7_1;

$L__BB7_18:
	bar.sync 	0;

$L__BB7_1:
	.loc	1 0 9
	setp.ne.s32 	%p1, %r1, 0;
	.loc	1 52 5
	@%p1 bra 	$L__BB7_3;

	.loc	1 52 5
	.loc	2 112 3, function_name $L__info_string5, inlined_at 1 52 5
	atom.global.add.u32 	%r16, [%rd6], 1;
	.loc	1 52 5
	st.shared.u32 	[_ZZ12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4_E3job], %r16;

$L__BB7_3:
	bar.sync 	0;
	ld.shared.u32 	%r17, [_ZZ12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4_E3job];
	cvt.s64.s32 	%rd7, %r17;
	setp.lt.u32 	%p2, %r17, %r15;
	@%p2 bra 	$L__BB7_5;
	bra.uni 	$L__BB7_4;

$L__BB7_5:
	.loc	1 0 5
	setp.eq.s32 	%p3, %r1, 0;
	.loc	1 53 9
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 53 9
	mul.lo.s64 	%rd19, %rd7, 48;
	add.s64 	%rd8, %rd5, %rd19;
	@%p3 bra 	$L__BB7_6;
	bra.uni 	$L__BB7_7;

$L__BB7_6:
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 53 9
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 36 5
	// begin inline asm
	mov.u64 %rd20, %globaltimer;
	// end inline asm
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 53 9
	st.global.u64 	[%rd8], %rd20;
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 53 9
	.loc	1 30 34, function_name $L__info_string2, inlined_at 1 36 5
	// begin inline asm
	mov.u32 %r18, %smid;
	// end inline asm
	.loc	1 36 5, function_name $L__info_string0, inlined_at 1 53 9
	st.global.v2.u32 	[%rd8+32], {%r18, %r2};

$L__BB7_7:
	.loc	1 0 5
	cvt.u32.u64 	%r19, %rd7;
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 53 9
	shl.b32 	%r20, %r19, 8;
	add.s32 	%r21, %r20, %r1;
	cvt.s64.s32 	%rd9, %r21;
	mul.wide.s32 	%rd21, %r21, 4;
	add.s64 	%rd22, %rd4, %rd21;
	ld.global.u32 	%r31, [%rd22];
	mul.wide.s32 	%rd23, %r19, 4;
	add.s64 	%rd24, %rd3, %rd23;
	ld.global.u32 	%r4, [%rd24];
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 53 9
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 37 5
	setp.lt.s32 	%p4, %r4, 1;
	@%p4 bra 	$L__BB7_10;

	.loc	1 0 5
	mov.u32 	%r29, 0;

$L__BB7_9:
	.pragma "nounroll";
	.loc	1 33 5
	mad.lo.s32 	%r31, %r31, 1664525, 1013904223;
	add.s32 	%r29, %r29, 1;
	setp.lt.s32 	%p5, %r29, %r4;
	@%p5 bra 	$L__BB7_9;

$L__BB7_10:
	.loc	1 37 5, function_name $L__info_string0, inlined_at 1 53 9
	shl.b64 	%rd25, %rd9, 2;
	add.s64 	%rd26, %rd2, %rd25;
	st.volatile.global.u32 	[%rd26], %r31;
	.loc	1 38 5, function_name $L__info_string0, inlined_at 1 53 9
	bar.sync 	0;
	.loc	1 39 5, function_name $L__info_string0, inlined_at 1 53 9
	@%p1 bra 	$L__BB7_12;

	.loc	1 39 5, function_name $L__info_string0, inlined_at 1 53 9
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 39 5
	// begin inline asm
	mov.u64 %rd27, %globaltimer;
	// end inline asm
	.loc	1 39 5, function_name $L__info_string0, inlined_at 1 53 9
	st.global.u64 	[%rd8+8], %rd27;

$L__BB7_12:
	.loc	1 53 9
	bar.sync 	0;
	ld.shared.u32 	%r23, [_ZZ12ragged_queueILb1EEvPKjPVjPjPKiP5StampiS4_E3job];
	cvt.s64.s32 	%rd10, %r23;
	.loc	1 53 9
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 53 9
	mul.wide.s32 	%rd28, %r23, 48;
	add.s64 	%rd29, %rd5, %rd28;
	add.s64 	%rd11, %rd29, 16;
	@%p1 bra 	$L__BB7_14;

	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 53 9
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 42 5
	// begin inline asm
	mov.u64 %rd30, %globaltimer;
	// end inline asm
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 53 9
	st.global.u64 	[%rd11], %rd30;
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 53 9
	.loc	1 30 34, function_name $L__info_string2, inlined_at 1 42 5
	// begin inline asm
	mov.u32 %r24, %smid;
	// end inline asm
	.loc	1 42 5, function_name $L__info_string4, inlined_at 1 53 9
	st.global.v2.u32 	[%rd11+24], {%r24, %r2};

$L__BB7_14:
	.loc	1 0 5
	cvt.u32.u64 	%r26, %rd10;
	.loc	1 43 5, function_name $L__info_string4, inlined_at 1 53 9
	shl.b32 	%r27, %r26, 8;
	add.s32 	%r28, %r27, %r1;
	cvt.s64.s32 	%rd12, %r28;
	mul.wide.s32 	%rd31, %r28, 4;
	add.s64 	%rd32, %rd2, %rd31;
	ld.volatile.global.u32 	%r33, [%rd32];
	mov.u32 	%r32, 0;

$L__BB7_15:
	.pragma "nounroll";
	.loc	1 43 5, function_name $L__info_string4, inlined_at 1 53 9
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 43 5
	mad.lo.s32 	%r33, %r33, 22695477, 1;
	add.s32 	%r32, %r32, 1;
	setp.ne.s32 	%p8, %r32, 512;
	@%p8 bra 	$L__BB7_15;

	.loc	1 43 5, function_name $L__info_string4, inlined_at 1 53 9
	shl.b64 	%rd33, %rd12, 2;
	add.s64 	%rd34, %rd1, %rd33;
	st.global.u32 	[%rd34], %r33;
	.loc	1 44 5, function_name $L__info_string4, inlined_at 1 53 9
	bar.sync 	0;
	.loc	1 45 5, function_name $L__info_string4, inlined_at 1 53 9
	@%p1 bra 	$L__BB7_18;

	.loc	1 45 5, function_name $L__info_string4, inlined_at 1 53 9
	.loc	1 29 44, function_name $L__info_string1, inlined_at 1 45 5
	// begin inline asm
	mov.u64 %rd35, %globaltimer;
	// end inline asm
	.loc	1 45 5, function_name $L__info_string4, inlined_at 1 53 9
	st.global.u64 	[%rd11+8], %rd35;
	bra.uni 	$L__BB7_18;

$L__BB7_4:
	.loc	1 54 1
	ret;

}
	// .globl	_Z8ragged_aILb0EEvPKjPVjPKiP5Stamp
.visible .entry _Z8ragged_aILb0EEvPKjPVjPKiP5Stamp(
	.param .u64 _Z8ragged_aILb0EEvPKjPVjPKiP5Stamp_param_0,
	.param .u64 _Z8ragged_aILb0EEvPKjPVjPKiP5Stamp_param_1,
	.param .u64 _Z8ragged_aILb0EEvPKjPVjPKiP5Stamp_param_2,
	.param .u64 _Z8ragged_aILb0EEvPKjPVjPKiP5Stamp_param_3
)
{
	.reg .pred 	%p<3>;
	.reg .b32 	%r<16>;
	.reg .b64 	%rd<14>;
	.loc	1 47 0


	ld.param.u64 	%rd3, [_Z8ragged_aILb0EEvPKjPVjPKiP5Stamp_param_0];
	ld.param.u64 	%rd2, [_Z8ragged_aILb0EEvPKjPVjPKiP5Stamp_param_1];
	ld.param.u64 	%rd4, [_Z8ragged_aILb0EEvPKjPVjPKiP5Stamp_param_2];
	.loc	1 47 51
	mov.u32 	%r8, %ctaid.x;
	.loc	1 47 51
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 47 51
	shl.b32 	%r9, %r8, 8;
	mov.u32 	%r10, %tid.x;
	add.s32 	%r11, %r9, %r10;
	cvt.s64.s32 	%rd1, %r11;
	.loc	1 47 51
	cvta.to.global.u64 	%rd5, %rd3;
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 47 51
	mul.wide.s32 	%rd6, %r11, 4;
	add.s64 	%rd7, %rd5, %rd6;
	ld.global.u32 	%r15, [%rd7];
	.loc	1 47 51
	cvta.to.global.u64 	%rd8, %rd4;
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 47 51
	mul.wide.s32 	%rd9, %r8, 4;
	add.s64 	%rd10, %rd8, %rd9;
	ld.global.u32 	%r2, [%rd10];
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 47 51
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 37 5
	setp.lt.s32 	%p1, %r2, 1;
	@%p1 bra 	$L__BB8_3;

	.loc	1 0 5
	mov.u32 	%r13, 0;

$L__BB8_2:
	.pragma "nounroll";
	.loc	1 33 5
	mad.lo.s32 	%r15, %r15, 1664525, 1013904223;
	add.s32 	%r13, %r13, 1;
	setp.lt.s32 	%p2, %r13, %r2;
	@%p2 bra 	$L__BB8_2;

$L__BB8_3:
	.loc	1 47 51
	cvta.to.global.u64 	%rd11, %rd2;
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 47 51
	shl.b64 	%rd12, %rd1, 2;
	add.s64 	%rd13, %rd11, %rd12;
	st.volatile.global.u32 	[%rd13], %r15;
	.loc	1 38 5, function_name $L__info_string6, inlined_at 1 47 51
	bar.sync 	0;
	.loc	1 47 51
	ret;

}
	// .globl	_Z8ragged_bILb0EEvPVKjPjP5Stamp
.visible .entry _Z8ragged_bILb0EEvPVKjPjP5Stamp(
	.param .u64 _Z8ragged_bILb0EEvPVKjPjP5Stamp_param_0,
	.param .u64 _Z8ragged_bILb0EEvPVKjPjP5Stamp_param_1,
	.param .u64 _Z8ragged_bILb0EEvPVKjPjP5Stamp_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .b32 	%r<13>;
	.reg .b64 	%rd<10>;
	.loc	1 48 0


	ld.param.u64 	%rd3, [_Z8ragged_bILb0EEvPVKjPjP5Stamp_param_0];
	ld.param.u64 	%rd2, [_Z8ragged_bILb0EEvPVKjPjP5Stamp_param_1];
	.loc	1 48 51
	cvta.to.global.u64 	%rd4, %rd3;
	mov.u32 	%r7, %ctaid.x;
	.loc	1 48 51
	.loc	1 43 5, function_name $L__info_string7, inlined_at 1 48 51
	shl.b32 	%r8, %r7, 8;
	mov.u32 	%r9, %tid.x;
	add.s32 	%r10, %r8, %r9;
	cvt.s64.s32 	%rd1, %r10;
	mul.wide.s32 	%rd5, %r10, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.volatile.global.u32 	%r12, [%rd6];
	mov.u32 	%r11, 0;

$L__BB9_1:
	.pragma "nounroll";
	.loc	1 43 5, function_name $L__info_string7, inlined_at 1 48 51
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 43 5
	mad.lo.s32 	%r12, %r12, 22695477, 1;
	add.s32 	%r11, %r11, 1;
	setp.ne.s32 	%p1, %r11, 512;
	@%p1 bra 	$L__BB9_1;

	.loc	1 48 51
	cvta.to.global.u64 	%rd7, %rd2;
	.loc	1 43 5, function_name $L__info_string7, inlined_at 1 48 51
	shl.b64 	%rd8, %rd1, 2;
	add.s64 	%rd9, %rd7, %rd8;
	st.global.u32 	[%rd9], %r12;
	.loc	1 44 5, function_name $L__info_string7, inlined_at 1 48 51
	bar.sync 	0;
	.loc	1 48 51
	ret;

}
	// .globl	_Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp
.visible .entry _Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp(
	.param .u64 _Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp_param_0,
	.param .u64 _Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp_param_1,
	.param .u64 _Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp_param_2,
	.param .u64 _Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp_param_3,
	.param .u64 _Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp_param_4
)
{
	.reg .pred 	%p<4>;
	.reg .b32 	%r<24>;
	.reg .b64 	%rd<18>;
	.loc	1 49 0


	ld.param.u64 	%rd5, [_Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp_param_0];
	ld.param.u64 	%rd3, [_Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp_param_1];
	ld.param.u64 	%rd4, [_Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp_param_2];
	ld.param.u64 	%rd6, [_Z12ragged_fusedILb0EEvPKjPVjPjPKiP5Stamp_param_3];
	.loc	1 49 51
	mov.u32 	%r13, %ctaid.x;
	.loc	1 49 51
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 49 51
	shl.b32 	%r14, %r13, 8;
	mov.u32 	%r15, %tid.x;
	add.s32 	%r16, %r14, %r15;
	cvt.s64.s32 	%rd1, %r16;
	.loc	1 49 51
	cvta.to.global.u64 	%rd7, %rd5;
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 49 51
	mul.wide.s32 	%rd8, %r16, 4;
	add.s64 	%rd9, %rd7, %rd8;
	ld.global.u32 	%r21, [%rd9];
	.loc	1 49 51
	cvta.to.global.u64 	%rd10, %rd6;
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 49 51
	mul.wide.s32 	%rd11, %r13, 4;
	add.s64 	%rd12, %rd10, %rd11;
	ld.global.u32 	%r2, [%rd12];
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 49 51
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 37 5
	setp.lt.s32 	%p1, %r2, 1;
	@%p1 bra 	$L__BB10_3;

	.loc	1 0 5
	mov.u32 	%r19, 0;

$L__BB10_2:
	.pragma "nounroll";
	.loc	1 33 5
	mad.lo.s32 	%r21, %r21, 1664525, 1013904223;
	add.s32 	%r19, %r19, 1;
	setp.lt.s32 	%p2, %r19, %r2;
	@%p2 bra 	$L__BB10_2;

$L__BB10_3:
	.loc	1 49 51
	cvta.to.global.u64 	%rd13, %rd3;
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 49 51
	shl.b64 	%rd14, %rd1, 2;
	add.s64 	%rd15, %rd13, %rd14;
	st.volatile.global.u32 	[%rd15], %r21;
	.loc	1 38 5, function_name $L__info_string6, inlined_at 1 49 51
	bar.sync 	0;
	.loc	1 49 51
	bar.sync 	0;
	.loc	1 49 51
	.loc	1 43 5, function_name $L__info_string7, inlined_at 1 49 51
	ld.volatile.global.u32 	%r23, [%rd15];
	.loc	1 49 51
	cvta.to.global.u64 	%rd2, %rd4;
	mov.u32 	%r22, 0;

$L__BB10_4:
	.pragma "nounroll";
	.loc	1 43 5, function_name $L__info_string7, inlined_at 1 49 51
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 43 5
	mad.lo.s32 	%r23, %r23, 22695477, 1;
	add.s32 	%r22, %r22, 1;
	setp.ne.s32 	%p3, %r22, 512;
	@%p3 bra 	$L__BB10_4;

	.loc	1 43 5, function_name $L__info_string7, inlined_at 1 49 51
	add.s64 	%rd17, %rd2, %rd14;
	st.global.u32 	[%rd17], %r23;
	.loc	1 44 5, function_name $L__info_string7, inlined_at 1 49 51
	bar.sync 	0;
	.loc	1 49 51
	ret;

}
	// .globl	_Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4_
.visible .entry _Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4_(
	.param .u64 _Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_0,
	.param .u64 _Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_1,
	.param .u64 _Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_2,
	.param .u64 _Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_3,
	.param .u64 _Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_4,
	.param .u32 _Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_5,
	.param .u64 _Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_6
)
{
	.reg .pred 	%p<6>;
	.reg .b32 	%r<29>;
	.reg .b64 	%rd<23>;
	.loc	1 50 0
	// demoted variable
	.shared .align 4 .u32 _ZZ12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4_E3job;

	ld.param.u64 	%rd8, [_Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_0];
	ld.param.u64 	%rd9, [_Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_1];
	ld.param.u64 	%rd10, [_Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_2];
	ld.param.u64 	%rd11, [_Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_3];
	ld.param.u32 	%r15, [_Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_5];
	ld.param.u64 	%rd12, [_Z12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4__param_6];
	.loc	1 52 5
	cvta.to.global.u64 	%rd1, %rd10;
	cvta.to.global.u64 	%rd2, %rd9;
	cvta.to.global.u64 	%rd3, %rd11;
	cvta.to.global.u64 	%rd4, %rd8;
	cvta.to.global.u64 	%rd5, %rd12;
	mov.u32 	%r1, %tid.x;
	bra.uni 	$L__BB11_1;

$L__BB11_10:
	.loc	1 53 9
	.loc	1 43 5, function_name $L__info_string7, inlined_at 1 53 9
	shl.b64 	%rd21, %rd7, 2;
	add.s64 	%rd22, %rd1, %rd21;
	st.global.u32 	[%rd22], %r28;
	.loc	1 44 5, function_name $L__info_string7, inlined_at 1 53 9
	bar.sync 	0;
	.loc	1 53 9
	bar.sync 	0;

$L__BB11_1:
	.loc	1 0 9
	setp.ne.s32 	%p1, %r1, 0;
	.loc	1 52 5
	@%p1 bra 	$L__BB11_3;

	.loc	1 52 5
	.loc	2 112 3, function_name $L__info_string5, inlined_at 1 52 5
	atom.global.add.u32 	%r16, [%rd5], 1;
	.loc	1 52 5
	st.shared.u32 	[_ZZ12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4_E3job], %r16;

$L__BB11_3:
	bar.sync 	0;
	ld.shared.u32 	%r2, [_ZZ12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4_E3job];
	setp.lt.u32 	%p2, %r2, %r15;
	@%p2 bra 	$L__BB11_5;
	bra.uni 	$L__BB11_4;

$L__BB11_5:
	.loc	1 53 9
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 53 9
	shl.b32 	%r17, %r2, 8;
	add.s32 	%r18, %r17, %r1;
	cvt.s64.s32 	%rd6, %r18;
	mul.wide.s32 	%rd13, %r18, 4;
	add.s64 	%rd14, %rd4, %rd13;
	ld.global.u32 	%r26, [%rd14];
	mul.wide.s32 	%rd15, %r2, 4;
	add.s64 	%rd16, %rd3, %rd15;
	ld.global.u32 	%r4, [%rd16];
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 53 9
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 37 5
	setp.lt.s32 	%p3, %r4, 1;
	@%p3 bra 	$L__BB11_8;

	.loc	1 0 5
	mov.u32 	%r24, 0;

$L__BB11_7:
	.pragma "nounroll";
	.loc	1 33 5
	mad.lo.s32 	%r26, %r26, 1664525, 1013904223;
	add.s32 	%r24, %r24, 1;
	setp.lt.s32 	%p4, %r24, %r4;
	@%p4 bra 	$L__BB11_7;

$L__BB11_8:
	.loc	1 37 5, function_name $L__info_string6, inlined_at 1 53 9
	shl.b64 	%rd17, %rd6, 2;
	add.s64 	%rd18, %rd2, %rd17;
	st.volatile.global.u32 	[%rd18], %r26;
	.loc	1 38 5, function_name $L__info_string6, inlined_at 1 53 9
	bar.sync 	0;
	.loc	1 53 9
	bar.sync 	0;
	mov.u32 	%r27, 0;
	ld.shared.u32 	%r21, [_ZZ12ragged_queueILb0EEvPKjPVjPjPKiP5StampiS4_E3job];
	.loc	1 43 5, function_name $L__info_string7, inlined_at 1 53 9
	shl.b32 	%r22, %r21, 8;
	add.s32 	%r23, %r22, %r1;
	cvt.s64.s32 	%rd7, %r23;
	mul.wide.s32 	%rd19, %r23, 4;
	add.s64 	%rd20, %rd2, %rd19;
	ld.volatile.global.u32 	%r28, [%rd20];

$L__BB11_9:
	.pragma "nounroll";
	.loc	1 43 5, function_name $L__info_string7, inlined_at 1 53 9
	.loc	1 33 5, function_name $L__info_string3, inlined_at 1 43 5
	mad.lo.s32 	%r28, %r28, 22695477, 1;
	add.s32 	%r27, %r27, 1;
	setp.ne.s32 	%p5, %r27, 512;
	@%p5 bra 	$L__BB11_9;
	bra.uni 	$L__BB11_10;

$L__BB11_4:
	.loc	1 54 1
	ret;

}
	.file	1 "/workspace/bench.cu"
	.file	2 "/usr/local/cuda/bin/../targets/x86_64-linux/include/device_atomic_functions.hpp"
	.section	.debug_str
	{
$L__info_string0:
.b8 95,90,55,115,116,97,103,101,95,97,73,76,98,49,69,69,118,105,105,80,75,106,80,86,106,80,75,105,80,53,83,116,97,109,112,0
$L__info_string1:
.b8 95,90,53,116,105,109,101,114,118,0
$L__info_string2:
.b8 95,90,52,115,109,105,100,118,0
$L__info_string3:
.b8 95,90,55,105,116,101,114,97,116,101,106,105,106,106,0
$L__info_string4:
.b8 95,90,55,115,116,97,103,101,95,98,73,76,98,49,69,69,118,105,80,86,75,106,80,106,80,53,83,116,97,109,112,0
$L__info_string5:
.b8 95,90,78,51,56,95,73,78,84,69,82,78,65,76,95,100,102,53,49,98,48,99,53,95,56,95,98,101,110,99,104,95,99,117,95,100,50,57,57,99
.b8 99,52,54,57,97,116,111,109,105,99,65,100,100,69,80,106,106,0
$L__info_string6:
.b8 95,90,55,115,116,97,103,101,95,97,73,76,98,48,69,69,118,105,105,80,75,106,80,86,106,80,75,105,80,53,83,116,97,109,112,0
$L__info_string7:
.b8 95,90,55,115,116,97,103,101,95,98,73,76,98,48,69,69,118,105,80,86,75,106,80,106,80,53,83,116,97,109,112,0

	}