.version 1.4
	.target sm_10, map_f64_to_f32
	// compiled with /usr/local/cuda-6.0/open64/lib//be
	// nvopencc 4.1 built on 2014-03-13

	//-----------------------------------------------------------
	// Compiling /tmp/tmpxft_00006744_00000000-9_util-cuda-bin.cpp3.i (/tmp/ccBI#.6MXtxK)
	//-----------------------------------------------------------

	//-----------------------------------------------------------
	// Options:
	//-----------------------------------------------------------
	//  Target:ptx, ISA:sm_10, Endian:little, Pointer Size:64
	//  -O3	(Optimization level)
	//  -g0	(Debug level)
	//  -m2	(Report advisories)
	//-----------------------------------------------------------

	.file	1	"<command-line>"
	.file	2	"/tmp/tmpxft_00006744_00000000-8_util-cuda-bin.cudafe2.gpu"
	.file	3	"/tmp/tmpxft_00006744_00000000-3_util-cuda-bin.cudafe1.gpu"
	.file	4	"/usr/lib/gcc/x86_64-linux-gnu/4.6/include/stddef.h"
	.file	5	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/crt/device_runtime.h"
	.file	6	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/host_defines.h"
	.file	7	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/builtin_types.h"
	.file	8	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/device_types.h"
	.file	9	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/driver_types.h"
	.file	10	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/surface_types.h"
	.file	11	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/texture_types.h"
	.file	12	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/vector_types.h"
	.file	13	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/device_launch_parameters.h"
	.file	14	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/crt/storage_class.h"
	.file	15	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/cuComplex.h"
	.file	16	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/cufft.h"
	.file	17	"util-cuda-bin.cu"
	.file	18	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/common_functions.h"
	.file	19	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/math_functions.h"
	.file	20	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/math_constants.h"
	.file	21	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/device_functions.h"
	.file	22	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/sm_11_atomic_functions.h"
	.file	23	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/sm_12_atomic_functions.h"
	.file	24	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/sm_13_double_functions.h"
	.file	25	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/sm_20_atomic_functions.h"
	.file	26	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/sm_32_atomic_functions.h"
	.file	27	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/sm_35_atomic_functions.h"
	.file	28	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/sm_20_intrinsics.h"
	.file	29	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/sm_30_intrinsics.h"
	.file	30	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/sm_32_intrinsics.h"
	.file	31	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/sm_35_intrinsics.h"
	.file	32	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/surface_functions.h"
	.file	33	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/texture_fetch_functions.h"
	.file	34	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/texture_indirect_functions.h"
	.file	35	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/surface_indirect_functions.h"
	.file	36	"/usr/local/cuda-6.0/bin/../targets/x86_64-linux/include/math_functions_dbl_ptx1.h"


	.entry elt_prod_conj (
		.param .u64 __cudaparm_elt_prod_conj_fc,
		.param .u64 __cudaparm_elt_prod_conj_c1,
		.param .u64 __cudaparm_elt_prod_conj_c2,
		.param .s32 __cudaparm_elt_prod_conj_size)
	{
	.reg .u16 %rh<3>;
	.reg .u32 %r<14>;
	.reg .u64 %rd<18>;
	.reg .f32 %f<16>;
	.reg .f64 %fd<38>;
	.reg .pred %p<10>;
	.shared .align 16 .b8 __cuda___cuda_local_var_17115_45_non_const_sc132[4096];
	.shared .align 16 .b8 __cuda___cuda_local_var_17116_45_non_const_sc24128[4096];
	.shared .align 16 .b8 __cuda___cuda_local_var_17114_45_non_const_sfc8224[4096];
	.loc	17	90	0
$LDWbegin_elt_prod_conj:
	mov.u16 	%rh1, %ctaid.x;
	mul.wide.u16 	%r1, %rh1, 256;
	cvt.u32.u16 	%r2, %tid.x;
	add.u32 	%r3, %r1, %r2;
	ld.param.s32 	%r4, [__cudaparm_elt_prod_conj_size];
	setp.gt.s32 	%p1, %r4, %r3;
	@%p1 bra 	$Lt_0_3842;
	bra.uni 	$LBB10_elt_prod_conj;
$Lt_0_3842:
	.loc	17	101	0
	cvt.u64.u32 	%rd1, %r2;
	mul.wide.u32 	%rd2, %r2, 16;
	mov.u64 	%rd3, __cuda___cuda_local_var_17115_45_non_const_sc132;
	add.u64 	%rd4, %rd2, %rd3;
	cvt.s64.s32 	%rd5, %r3;
	mul.wide.s32 	%rd6, %r3, 16;
	ld.param.u64 	%rd7, [__cudaparm_elt_prod_conj_c1];
	add.u64 	%rd8, %rd7, %rd6;
	ld.global.v2.f64 	{%fd1,%fd2}, [%rd8+0];
	st.shared.f64 	[%rd4+0], %fd1;
	st.shared.f64 	[%rd4+8], %fd2;
	.loc	17	102	0
	mov.u64 	%rd9, __cuda___cuda_local_var_17116_45_non_const_sc24128;
	add.u64 	%rd10, %rd2, %rd9;
	ld.param.u64 	%rd11, [__cudaparm_elt_prod_conj_c2];
	add.u64 	%rd12, %rd11, %rd6;
	ld.global.v2.f64 	{%fd3,%fd4}, [%rd12+0];
	st.shared.f64 	[%rd10+0], %fd3;
	st.shared.f64 	[%rd10+8], %fd4;
	.loc	17	104	0
	bar.sync 	0;
	.loc	17	106	0
	ld.shared.f64 	%fd5, [%rd4+0];
	ld.shared.f64 	%fd6, [%rd4+8];
	ld.shared.f64 	%fd7, [%rd10+0];
	ld.shared.f64 	%fd8, [%rd10+8];
	mov.u64 	%rd13, __cuda___cuda_local_var_17114_45_non_const_sfc8224;
	add.u64 	%rd14, %rd2, %rd13;
	mul.f64 	%fd9, %fd6, %fd8;
	mad.rn.f64 	%fd10, %fd7, %fd5, %fd9;
	mov.f64 	%fd11, %fd10;
	st.shared.f64 	[%rd14+0], %fd10;
	mul.f64 	%fd12, %fd5, %fd8;
	mul.f64 	%fd13, %fd7, %fd6;
	sub.f64 	%fd14, %fd13, %fd12;
	st.shared.f64 	[%rd14+8], %fd14;
	.loc	17	108	0
	abs.f64 	%fd15, %fd10;
	abs.f64 	%fd16, %fd14;
	cvt.rn.f32.f64 	%f1, %fd15;
	cvt.f64.f32 	%fd17, %f1;
	cvt.rn.f32.f64 	%f2, %fd16;
	cvt.f64.f32 	%fd18, %f2;
	setp.gt.f64 	%p2, %fd17, %fd18;
	@!%p2 bra 	$Lt_0_4610;
	mov.f64 	%fd19, %fd17;
	mov.f64 	%fd20, %fd18;
	bra.uni 	$Lt_0_4354;
$Lt_0_4610:
	mov.f64 	%fd19, %fd18;
	mov.f64 	%fd20, %fd17;
$Lt_0_4354:
	mov.f64 	%fd21, 0d7fefffffffffffff;	// 1.79769e+308
	setp.gt.f64 	%p3, %fd20, %fd21;
	mov.f64 	%fd22, 0d0000000000000000;	// 0
	setp.eq.f64 	%p4, %fd19, %fd22;
	mov.f64 	%fd23, 0d7fefffffffffffff;	// 1.79769e+308
	setp.gt.f64 	%p5, %fd19, %fd23;
	cvt.rn.f32.f64 	%f3, %fd20;
	cvt.rn.f32.f64 	%f4, %fd19;
	selp.s32 	%r5, 1, 0, %p4;
	selp.s32 	%r6, 1, 0, %p5;
	div.full.f32 	%f5, %f3, %f4;
	cvt.f64.f32 	%fd24, %f5;
	mov.f64 	%fd25, 0d3ff0000000000000;	// 1
	mad.rn.f64 	%fd26, %fd24, %fd24, %fd25;
	cvt.rn.f32.f64 	%f6, %fd26;
	add.f64 	%fd27, %fd19, %fd20;
	sqrt.approx.f32 	%f7, %f6;
	cvt.f64.f32 	%fd28, %f7;
	mul.f64 	%fd29, %fd28, %fd19;
	selp.s32 	%r7, 1, 0, %p3;
	or.b32 	%r8, %r5, %r6;
	or.b32 	%r9, %r7, %r8;
	neg.s32 	%r10, %r9;
	slct.f64.s32 	%fd30, %fd29, %fd27, %r10;
	mov.f64 	%fd31, %fd30;
	.loc	17	110	0
	mov.f64 	%fd32, 0d0000000000000000;	// 0
	setp.eq.f64 	%p6, %fd30, %fd32;
	@%p6 bra 	$Lt_0_2306;
	.loc	19	9055	0
	abs.f64 	%fd33, %fd30;
	cvt.rn.f32.f64 	%f8, %fd33;
	mov.f32 	%f9, 0f7f800000;     	// ((1.0F)/(0.0F))
	setp.le.f32 	%p7, %f8, %f9;
	selp.s32 	%r11, 0, 1, %p7;
	.loc	17	110	0
	mov.u32 	%r12, 0;
	setp.eq.s32 	%p8, %r11, %r12;
	@%p8 bra 	$Lt_0_2562;
$Lt_0_2306:
	.loc	17	113	0
	mov.f64 	%fd11, 0d3cb0000000000000;	// 2.22045e-16
	mov.f64 	%fd34, 0d3cb0000000000000;	// 2.22045e-16
	st.shared.f64 	[%rd14+0], %fd34;
	mov.f64 	%fd31, 0d3cb0000000000000;	// 2.22045e-16
$Lt_0_2562:
	.loc	17	116	0
	cvt.rn.f32.f64 	%f10, %fd31;
	ld.param.u64 	%rd15, [__cudaparm_elt_prod_conj_fc];
	add.u64 	%rd16, %rd15, %rd6;
	cvt.rn.f32.f64 	%f11, %fd11;
	div.full.f32 	%f12, %f11, %f10;
	cvt.f64.f32 	%fd35, %f12;
	ld.shared.f64 	%fd14, [%rd14+8];
	cvt.rn.f32.f64 	%f13, %fd14;
	div.full.f32 	%f14, %f13, %f10;
	cvt.f64.f32 	%fd36, %f14;
	st.global.v2.f64 	[%rd16+0], {%fd35,%fd36};
$LBB10_elt_prod_conj:
	.loc	17	118	0
	exit;
$LDWend_elt_prod_conj:
	} // elt_prod_conj

	.entry elt_prod_conj_v2 (
		.param .u64 __cudaparm_elt_prod_conj_v2_fc,
		.param .u64 __cudaparm_elt_prod_conj_v2_c1,
		.param .u64 __cudaparm_elt_prod_conj_v2_c2,
		.param .s32 __cudaparm_elt_prod_conj_v2_size)
	{
	.reg .u16 %rh<3>;
	.reg .u32 %r<11>;
	.reg .u64 %rd<14>;
	.reg .f32 %f<11>;
	.reg .f64 %fd<21>;
	.reg .pred %p<5>;
	.shared .align 16 .b8 __cuda___cuda_local_var_17147_45_non_const_sfc12352[4096];
	.loc	17	123	0
$LDWbegin_elt_prod_conj_v2:
	mov.u16 	%rh1, %ctaid.x;
	mul.wide.u16 	%r1, %rh1, 256;
	cvt.u32.u16 	%r2, %tid.x;
	add.u32 	%r3, %r1, %r2;
	ld.param.s32 	%r4, [__cudaparm_elt_prod_conj_v2_size];
	setp.gt.s32 	%p1, %r4, %r3;
	@%p1 bra 	$Lt_1_2050;
	bra.uni 	$LBB6_elt_prod_conj_v2;
$Lt_1_2050:
	.loc	17	135	0
	cvt.s64.s32 	%rd1, %r3;
	mul.wide.s32 	%rd2, %r3, 16;
	ld.param.u64 	%rd3, [__cudaparm_elt_prod_conj_v2_c1];
	add.u64 	%rd4, %rd3, %rd2;
	ld.global.v2.f64 	{%fd1,%fd2}, [%rd4+0];
	ld.param.u64 	%rd5, [__cudaparm_elt_prod_conj_v2_c2];
	add.u64 	%rd6, %rd5, %rd2;
	ld.global.v2.f64 	{%fd3,%fd4}, [%rd6+0];
	mov.u64 	%rd7, __cuda___cuda_local_var_17147_45_non_const_sfc12352;
	cvt.u64.u32 	%rd8, %r2;
	mul.wide.u32 	%rd9, %r2, 16;
	add.u64 	%rd10, %rd7, %rd9;
	mul.f64 	%fd5, %fd2, %fd4;
	mad.rn.f64 	%fd6, %fd3, %fd1, %fd5;
	st.shared.f64 	[%rd10+0], %fd6;
	mul.f64 	%fd7, %fd1, %fd4;
	mul.f64 	%fd8, %fd3, %fd2;
	sub.f64 	%fd9, %fd8, %fd7;
	st.shared.f64 	[%rd10+8], %fd9;
	.loc	17	137	0
	bar.sync 	0;
	.loc	17	142	0
	ld.shared.f64 	%fd10, [%rd10+0];
	ld.shared.f64 	%fd11, [%rd10+8];
	mul.f64 	%fd12, %fd11, %fd11;
	mad.rn.f64 	%fd13, %fd10, %fd10, %fd12;
	cvt.rn.f32.f64 	%f1, %fd13;
	sqrt.approx.f32 	%f2, %f1;
	cvt.f64.f32 	%fd14, %f2;
	mov.f64 	%fd15, %fd14;
	.loc	17	145	0
	abs.f32 	%f3, %f2;
	mov.f32 	%f4, 0f7f800000;     	// ((1.0F)/(0.0F))
	setp.le.f32 	%p2, %f3, %f4;
	selp.s32 	%r5, 0, 1, %p2;
	mov.f64 	%fd16, 0d0000000000000000;	// 0
	set.eq.u32.f64 	%r6, %fd14, %fd16;
	neg.s32 	%r7, %r6;
	or.b32 	%r8, %r5, %r7;
	mov.u32 	%r9, 0;
	setp.eq.s32 	%p3, %r8, %r9;
	@%p3 bra 	$Lt_1_2562;
	.loc	17	148	0
	mov.f64 	%fd10, 0d3cb0000000000000;	// 2.22045e-16
	mov.f64 	%fd17, 0d3cb0000000000000;	// 2.22045e-16
	st.shared.f64 	[%rd10+0], %fd17;
	mov.f64 	%fd15, 0d3cb0000000000000;	// 2.22045e-16
$Lt_1_2562:
	.loc	17	156	0
	cvt.rn.f32.f64 	%f5, %fd15;
	ld.param.u64 	%rd11, [__cudaparm_elt_prod_conj_v2_fc];
	add.u64 	%rd12, %rd11, %rd2;
	cvt.rn.f32.f64 	%f6, %fd10;
	div.full.f32 	%f7, %f6, %f5;
	cvt.f64.f32 	%fd18, %f7;
	.loc	17	142	0
	ld.shared.f64 	%fd11, [%rd10+8];
	.loc	17	156	0
	cvt.rn.f32.f64 	%f8, %fd11;
	div.full.f32 	%f9, %f8, %f5;
	cvt.f64.f32 	%fd19, %f9;
	st.global.v2.f64 	[%rd12+0], {%fd18,%fd19};
$LBB6_elt_prod_conj_v2:
	.loc	17	158	0
	exit;
$LDWend_elt_prod_conj_v2:
	} // elt_prod_conj_v2

	.entry elt_prod_conj_v3 (
		.param .u64 __cudaparm_elt_prod_conj_v3_fc,
		.param .u64 __cudaparm_elt_prod_conj_v3_c1,
		.param .u64 __cudaparm_elt_prod_conj_v3_c2,
		.param .s32 __cudaparm_elt_prod_conj_v3_size)
	{
	.reg .u16 %rh<3>;
	.reg .u32 %r<20>;
	.reg .u64 %rd<10>;
	.reg .f32 %f<18>;
	.reg .f64 %fd<36>;
	.reg .pred %p<7>;
	.loc	17	163	0
$LDWbegin_elt_prod_conj_v3:
	mov.u16 	%rh1, %ctaid.x;
	mul.wide.u16 	%r1, %rh1, 256;
	cvt.u32.u16 	%r2, %tid.x;
	add.u32 	%r3, %r2, %r1;
	ld.param.s32 	%r4, [__cudaparm_elt_prod_conj_v3_size];
	setp.gt.s32 	%p1, %r4, %r3;
	@%p1 bra 	$Lt_2_4866;
	bra.uni 	$LBB9_elt_prod_conj_v3;
$Lt_2_4866:
	.loc	17	170	0
	cvt.s64.s32 	%rd1, %r3;
	mul.wide.s32 	%rd2, %r3, 16;
	ld.param.u64 	%rd3, [__cudaparm_elt_prod_conj_v3_c1];
	add.u64 	%rd4, %rd3, %rd2;
	ld.global.v2.f64 	{%fd1,%fd2}, [%rd4+0];
	.loc	17	171	0
	ld.param.u64 	%rd5, [__cudaparm_elt_prod_conj_v3_c2];
	add.u64 	%rd6, %rd5, %rd2;
	ld.global.v2.f64 	{%fd3,%fd4}, [%rd6+0];
	.loc	17	173	0
	mul.f64 	%fd5, %fd3, %fd2;
	neg.f64 	%fd6, %fd4;
	mad.rn.f64 	%fd7, %fd6, %fd1, %fd5;
	mul.f64 	%fd8, %fd2, %fd4;
	mad.rn.f64 	%fd9, %fd1, %fd3, %fd8;
	mul.f64 	%fd10, %fd7, %fd7;
	mad.rn.f64 	%fd11, %fd9, %fd9, %fd10;
	cvt.rn.f32.f64 	%f1, %fd11;
	sqrt.approx.f32 	%f2, %f1;
	cvt.f64.f32 	%fd12, %f2;
	mov.f64 	%fd13, %fd12;
	.loc	17	176	0
	abs.f32 	%f3, %f2;
	mov.f32 	%f4, 0f7f800000;     	// ((1.0F)/(0.0F))
	setp.le.f32 	%p2, %f3, %f4;
	selp.s32 	%r5, 0, 1, %p2;
	mov.f64 	%fd14, 0d0000000000000000;	// 0
	set.eq.u32.f64 	%r6, %fd12, %fd14;
	neg.s32 	%r7, %r6;
	or.b32 	%r8, %r5, %r7;
	mov.u32 	%r9, 0;
	setp.eq.s32 	%p3, %r8, %r9;
	@%p3 bra 	$Lt_2_5378;
	.loc	17	177	0
	abs.f64 	%fd15, %fd7;
	abs.f64 	%fd16, %fd9;
	cvt.rn.f32.f64 	%f5, %fd15;
	cvt.f64.f32 	%fd17, %f5;
	cvt.rn.f32.f64 	%f6, %fd16;
	cvt.f64.f32 	%fd18, %f6;
	setp.lt.f64 	%p4, %fd17, %fd18;
	@!%p4 bra 	$Lt_2_6146;
	mov.f64 	%fd19, %fd18;
	mov.f64 	%fd20, %fd17;
	bra.uni 	$Lt_2_5890;
$Lt_2_6146:
	mov.f64 	%fd19, %fd17;
	mov.f64 	%fd20, %fd18;
$Lt_2_5890:
	cvt.rn.f32.f64 	%f7, %fd20;
	cvt.rn.f32.f64 	%f8, %fd19;
	div.full.f32 	%f9, %f7, %f8;
	cvt.f64.f32 	%fd21, %f9;
	add.f64 	%fd22, %fd19, %fd20;
	mov.f64 	%fd23, 0d3ff0000000000000;	// 1
	mad.rn.f64 	%fd24, %fd21, %fd21, %fd23;
	cvt.rn.f32.f64 	%f10, %fd24;
	sqrt.approx.f32 	%f11, %f10;
	cvt.f64.f32 	%fd25, %f11;
	mul.f64 	%fd26, %fd25, %fd19;
	mov.f64 	%fd27, 0d7fefffffffffffff;	// 1.79769e+308
	set.gt.u32.f64 	%r10, %fd20, %fd27;
	neg.s32 	%r11, %r10;
	mov.f64 	%fd28, 0d0000000000000000;	// 0
	set.eq.u32.f64 	%r12, %fd19, %fd28;
	neg.s32 	%r13, %r12;
	mov.f64 	%fd29, 0d7fefffffffffffff;	// 1.79769e+308
	set.gt.u32.f64 	%r14, %fd19, %fd29;
	neg.s32 	%r15, %r14;
	or.b32 	%r16, %r13, %r15;
	or.b32 	%r17, %r11, %r16;
	neg.s32 	%r18, %r17;
	slct.f64.s32 	%fd13, %fd26, %fd22, %r18;
$Lt_2_5378:
	.loc	17	182	0
	mov.f64 	%fd30, 0d3cb0000000000000;	// 2.22045e-16
	mov.f64 	%fd31, 0d0000000000000000;	// 0
	setp.eq.f64 	%p5, %fd13, %fd31;
	selp.f64 	%fd32, %fd30, %fd13, %p5;
	cvt.rn.f32.f64 	%f12, %fd32;
	ld.param.u64 	%rd7, [__cudaparm_elt_prod_conj_v3_fc];
	add.u64 	%rd8, %rd7, %rd2;
	cvt.rn.f32.f64 	%f13, %fd9;
	div.full.f32 	%f14, %f13, %f12;
	cvt.f64.f32 	%fd33, %f14;
	cvt.rn.f32.f64 	%f15, %fd7;
	div.full.f32 	%f16, %f15, %f12;
	cvt.f64.f32 	%fd34, %f16;
	st.global.v2.f64 	[%rd8+0], {%fd33,%fd34};
$LBB9_elt_prod_conj_v3:
	.loc	17	183	0
	exit;
$LDWend_elt_prod_conj_v3:
	} // elt_prod_conj_v3

	.entry reduce_max_final (
		.param .u64 __cudaparm_reduce_max_final_g_idata,
		.param .u64 __cudaparm_reduce_max_final_g_odata,
		.param .u64 __cudaparm_reduce_max_final_max_idx,
		.param .u32 __cudaparm_reduce_max_final_n,
		.param .s32 __cudaparm_reduce_max_final_blockSize)
	{
	.reg .u32 %r<42>;
	.reg .u64 %rd<31>;
	.reg .f64 %fd<15>;
	.reg .pred %p<30>;
	.shared .align 8 .b8 __cuda___cuda_local_var_17212_33_non_const_sdata16512[2048];
	.shared .align 4 .b8 __cuda___cuda_local_var_17213_30_non_const_idxData18560[1024];
	.loc	17	188	0
$LDWbegin_reduce_max_final:
	.loc	17	197	0
	ld.param.s32 	%r1, [__cudaparm_reduce_max_final_blockSize];
	cvt.u32.u16 	%r2, %ctaid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	mul.lo.u32 	%r4, %r3, 2;
	cvt.u32.u16 	%r5, %tid.x;
	add.u32 	%r6, %r4, %r5;
	mov.s32 	%r7, %r6;
	ld.param.u32 	%r8, [__cudaparm_reduce_max_final_n];
	setp.ge.u32 	%p1, %r6, %r8;
	@%p1 bra 	$Lt_3_35586;
	cvt.u64.u32 	%rd1, %r6;
	ld.param.s32 	%r1, [__cudaparm_reduce_max_final_blockSize];
	add.u32 	%r9, %r1, %r6;
	cvt.u32.u16 	%r10, %nctaid.x;
	mul.lo.u32 	%r11, %r1, %r10;
	mul.lo.u32 	%r12, %r11, 2;
	ld.param.u64 	%rd2, [__cudaparm_reduce_max_final_g_idata];
	mul.wide.u32 	%rd3, %r6, 8;
	add.u64 	%rd4, %rd2, %rd3;
	cvt.s64.u32 	%rd5, %r12;
	mov.f64 	%fd1, 0d0000000000000000;	// 0
$Lt_3_21762:
 //<loop> Loop body line 197, nesting depth: 1, estimated iterations: unknown
	ld.global.f64 	%fd2, [%rd4+0];
	setp.gt.f64 	%p2, %fd2, %fd1;
	@!%p2 bra 	$Lt_3_22018;
	.loc	17	204	0
	mov.f64 	%fd1, %fd2;
	.loc	17	205	0
	ld.param.u64 	%rd6, [__cudaparm_reduce_max_final_max_idx];
	mul.lo.u64 	%rd7, %rd1, 4;
	add.u64 	%rd8, %rd6, %rd7;
	ld.global.s32 	%r13, [%rd8+0];
	mov.s32 	%r14, %r13;
$Lt_3_22018:
	.loc	17	197	0
	ld.param.u32 	%r8, [__cudaparm_reduce_max_final_n];
	.loc	17	205	0
	setp.ge.u32 	%p3, %r9, %r8;
	@%p3 bra 	$Lt_3_23042;
	cvt.u64.u32 	%rd9, %r9;
	mul.wide.u32 	%rd10, %r9, 8;
	.loc	17	197	0
	ld.param.u64 	%rd2, [__cudaparm_reduce_max_final_g_idata];
	.loc	17	205	0
	add.u64 	%rd11, %rd2, %rd10;
	ld.global.f64 	%fd3, [%rd11+0];
	setp.gt.f64 	%p4, %fd3, %fd1;
	@!%p4 bra 	$Lt_3_23042;
	.loc	17	212	0
	mov.f64 	%fd1, %fd3;
	.loc	17	213	0
	ld.param.u64 	%rd12, [__cudaparm_reduce_max_final_max_idx];
	mul.lo.u64 	%rd13, %rd9, 4;
	add.u64 	%rd14, %rd12, %rd13;
	ld.global.s32 	%r15, [%rd14+0];
	mov.s32 	%r14, %r15;
$Lt_3_23042:
$Lt_3_22530:
	add.u32 	%r7, %r12, %r7;
	add.u32 	%r9, %r9, %r12;
	add.u64 	%rd1, %rd5, %rd1;
	mul.lo.u64 	%rd15, %rd5, 8;
	add.u64 	%rd4, %rd4, %rd15;
	.loc	17	197	0
	ld.param.u32 	%r8, [__cudaparm_reduce_max_final_n];
	.loc	17	213	0
	setp.lt.u32 	%p5, %r7, %r8;
	@%p5 bra 	$Lt_3_21762;
	bra.uni 	$Lt_3_21250;
$Lt_3_35586:
	mov.f64 	%fd1, 0d0000000000000000;	// 0
$Lt_3_21250:
	.loc	17	220	0
	cvt.u64.u32 	%rd16, %r5;
	mov.u64 	%rd17, __cuda___cuda_local_var_17212_33_non_const_sdata16512;
	mul.wide.u32 	%rd18, %r5, 8;
	add.u64 	%rd19, %rd17, %rd18;
	st.shared.f64 	[%rd19+0], %fd1;
	.loc	17	221	0
	mov.u64 	%rd20, __cuda___cuda_local_var_17213_30_non_const_idxData18560;
	mul.wide.u32 	%rd21, %r5, 4;
	add.u64 	%rd22, %rd20, %rd21;
	mov.s32 	%r16, %r14;
	st.shared.s32 	[%rd22+0], %r16;
	.loc	17	223	0
	bar.sync 	0;
	mov.u32 	%r17, 511;
	setp.le.s32 	%p6, %r1, %r17;
	@%p6 bra 	$Lt_3_23810;
	mov.u32 	%r18, 255;
	setp.gt.u32 	%p7, %r5, %r18;
	@%p7 bra 	$Lt_3_24834;
	ld.shared.f64 	%fd4, [%rd19+2048];
	setp.gt.f64 	%p8, %fd4, %fd1;
	@!%p8 bra 	$Lt_3_24834;
	ld.shared.f64 	%fd4, [%rd19+2048];
	.loc	17	231	0
	mov.f64 	%fd1, %fd4;
	st.shared.f64 	[%rd19+0], %fd4;
	.loc	17	232	0
	ld.shared.s32 	%r19, [%rd22+1024];
	st.shared.s32 	[%rd22+0], %r19;
$Lt_3_24834:
$Lt_3_24322:
	.loc	17	235	0
	bar.sync 	0;
$Lt_3_23810:
	mov.u32 	%r20, 255;
	setp.le.s32 	%p9, %r1, %r20;
	@%p9 bra 	$Lt_3_25346;
	mov.u32 	%r21, 127;
	setp.gt.u32 	%p10, %r5, %r21;
	@%p10 bra 	$Lt_3_26370;
	ld.shared.f64 	%fd5, [%rd19+1024];
	setp.gt.f64 	%p11, %fd5, %fd1;
	@!%p11 bra 	$Lt_3_26370;
	ld.shared.f64 	%fd5, [%rd19+1024];
	.loc	17	244	0
	mov.f64 	%fd1, %fd5;
	st.shared.f64 	[%rd19+0], %fd5;
	.loc	17	245	0
	ld.shared.s32 	%r22, [%rd22+512];
	st.shared.s32 	[%rd22+0], %r22;
$Lt_3_26370:
$Lt_3_25858:
	.loc	17	248	0
	bar.sync 	0;
$Lt_3_25346:
	mov.u32 	%r23, 127;
	setp.le.s32 	%p12, %r1, %r23;
	@%p12 bra 	$Lt_3_26882;
	mov.u32 	%r24, 63;
	setp.gt.u32 	%p13, %r5, %r24;
	@%p13 bra 	$Lt_3_27906;
	ld.shared.f64 	%fd6, [%rd19+512];
	setp.gt.f64 	%p14, %fd6, %fd1;
	@!%p14 bra 	$Lt_3_27906;
	ld.shared.f64 	%fd6, [%rd19+512];
	.loc	17	257	0
	mov.f64 	%fd1, %fd6;
	st.shared.f64 	[%rd19+0], %fd6;
	.loc	17	258	0
	ld.shared.s32 	%r25, [%rd22+256];
	st.shared.s32 	[%rd22+0], %r25;
$Lt_3_27906:
$Lt_3_27394:
	.loc	17	261	0
	bar.sync 	0;
$Lt_3_26882:
	mov.u32 	%r26, 31;
	setp.gt.u32 	%p15, %r5, %r26;
	@%p15 bra 	$Lt_3_28418;
	mov.u32 	%r27, 63;
	setp.le.s32 	%p16, %r1, %r27;
	@%p16 bra 	$Lt_3_29442;
	ld.volatile.shared.f64 	%fd7, [%rd19+256];
	setp.gt.f64 	%p17, %fd7, %fd1;
	@!%p17 bra 	$Lt_3_29442;
	.loc	17	272	0
	ld.volatile.shared.f64 	%fd1, [%rd19+256];
	st.volatile.shared.f64 	[%rd19+0], %fd1;
	.loc	17	273	0
	ld.volatile.shared.s32 	%r28, [%rd22+128];
	st.volatile.shared.s32 	[%rd22+0], %r28;
$Lt_3_29442:
$Lt_3_28930:
	mov.u32 	%r29, 31;
	setp.le.s32 	%p18, %r1, %r29;
	@%p18 bra 	$Lt_3_30466;
	ld.volatile.shared.f64 	%fd8, [%rd19+128];
	setp.gt.f64 	%p19, %fd8, %fd1;
	@!%p19 bra 	$Lt_3_30466;
	.loc	17	279	0
	ld.volatile.shared.f64 	%fd1, [%rd19+128];
	st.volatile.shared.f64 	[%rd19+0], %fd1;
	.loc	17	280	0
	ld.volatile.shared.s32 	%r30, [%rd22+64];
	st.volatile.shared.s32 	[%rd22+0], %r30;
$Lt_3_30466:
$Lt_3_29954:
	mov.u32 	%r31, 15;
	setp.le.s32 	%p20, %r1, %r31;
	@%p20 bra 	$Lt_3_31490;
	ld.volatile.shared.f64 	%fd9, [%rd19+64];
	setp.gt.f64 	%p21, %fd9, %fd1;
	@!%p21 bra 	$Lt_3_31490;
	.loc	17	286	0
	ld.volatile.shared.f64 	%fd1, [%rd19+64];
	st.volatile.shared.f64 	[%rd19+0], %fd1;
	.loc	17	287	0
	ld.volatile.shared.s32 	%r32, [%rd22+32];
	st.volatile.shared.s32 	[%rd22+0], %r32;
$Lt_3_31490:
$Lt_3_30978:
	mov.u32 	%r33, 7;
	setp.le.s32 	%p22, %r1, %r33;
	@%p22 bra 	$Lt_3_32514;
	ld.volatile.shared.f64 	%fd10, [%rd19+32];
	setp.gt.f64 	%p23, %fd10, %fd1;
	@!%p23 bra 	$Lt_3_32514;
	.loc	17	293	0
	ld.volatile.shared.f64 	%fd1, [%rd19+32];
	st.volatile.shared.f64 	[%rd19+0], %fd1;
	.loc	17	294	0
	ld.volatile.shared.s32 	%r34, [%rd22+16];
	st.volatile.shared.s32 	[%rd22+0], %r34;
$Lt_3_32514:
$Lt_3_32002:
	mov.u32 	%r35, 3;
	setp.le.s32 	%p24, %r1, %r35;
	@%p24 bra 	$Lt_3_33538;
	ld.volatile.shared.f64 	%fd11, [%rd19+16];
	setp.gt.f64 	%p25, %fd11, %fd1;
	@!%p25 bra 	$Lt_3_33538;
	.loc	17	300	0
	ld.volatile.shared.f64 	%fd1, [%rd19+16];
	st.volatile.shared.f64 	[%rd19+0], %fd1;
	.loc	17	301	0
	ld.volatile.shared.s32 	%r36, [%rd22+8];
	st.volatile.shared.s32 	[%rd22+0], %r36;
$Lt_3_33538:
$Lt_3_33026:
	mov.u32 	%r37, 1;
	setp.le.s32 	%p26, %r1, %r37;
	@%p26 bra 	$Lt_3_34562;
	ld.volatile.shared.f64 	%fd12, [%rd19+8];
	setp.gt.f64 	%p27, %fd12, %fd1;
	@!%p27 bra 	$Lt_3_34562;
	.loc	17	307	0
	ld.volatile.shared.f64 	%fd1, [%rd19+8];
	st.volatile.shared.f64 	[%rd19+0], %fd1;
	.loc	17	308	0
	ld.volatile.shared.s32 	%r38, [%rd22+4];
	st.volatile.shared.s32 	[%rd22+0], %r38;
$Lt_3_34562:
$Lt_3_34050:
	.loc	17	310	0
	bar.sync 	0;
$Lt_3_28418:
	mov.u32 	%r39, 0;
	setp.ne.u32 	%p28, %r5, %r39;
	@%p28 bra 	$Lt_3_35074;
	.loc	17	315	0
	cvt.u64.u32 	%rd23, %r2;
	ld.shared.f64 	%fd13, [__cuda___cuda_local_var_17212_33_non_const_sdata16512+0];
	ld.param.u64 	%rd24, [__cudaparm_reduce_max_final_g_odata];
	mul.wide.u32 	%rd25, %r2, 8;
	add.u64 	%rd26, %rd24, %rd25;
	st.global.f64 	[%rd26+0], %fd13;
	.loc	17	316	0
	ld.shared.s32 	%r40, [__cuda___cuda_local_var_17213_30_non_const_idxData18560+0];
	ld.param.u64 	%rd27, [__cudaparm_reduce_max_final_max_idx];
	mul.wide.u32 	%rd28, %r2, 4;
	add.u64 	%rd29, %rd27, %rd28;
	st.global.s32 	[%rd29+0], %r40;
$Lt_3_35074:
	.loc	17	318	0
	exit;
$LDWend_reduce_max_final:
	} // reduce_max_final

	.entry reduce_max_main (
		.param .u64 __cudaparm_reduce_max_main_g_idata,
		.param .u64 __cudaparm_reduce_max_main_g_odata,
		.param .u64 __cudaparm_reduce_max_main_max_idx,
		.param .u32 __cudaparm_reduce_max_main_n,
		.param .s32 __cudaparm_reduce_max_main_blockSize)
	{
	.reg .u32 %r<38>;
	.reg .u64 %rd<25>;
	.reg .f64 %fd<14>;
	.reg .pred %p<30>;
	.shared .align 8 .b8 __cuda___cuda_local_var_17347_33_non_const_sdata19616[2048];
	.shared .align 4 .b8 __cuda___cuda_local_var_17348_30_non_const_idxData21664[1024];
	.loc	17	323	0
$LDWbegin_reduce_max_main:
	.loc	17	332	0
	ld.param.s32 	%r1, [__cudaparm_reduce_max_main_blockSize];
	cvt.u32.u16 	%r2, %ctaid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	cvt.u32.u16 	%r4, %tid.x;
	add.u32 	%r5, %r3, %r4;
	mov.s32 	%r6, %r5;
	ld.param.u32 	%r7, [__cudaparm_reduce_max_main_n];
	setp.ge.u32 	%p1, %r5, %r7;
	@%p1 bra 	$Lt_4_35586;
	cvt.u32.u16 	%r8, %nctaid.x;
	ld.param.s32 	%r1, [__cudaparm_reduce_max_main_blockSize];
	mul.lo.u32 	%r9, %r1, %r8;
	cvt.s64.u32 	%rd1, %r9;
	ld.param.u64 	%rd2, [__cudaparm_reduce_max_main_g_idata];
	cvt.u64.u32 	%rd3, %r5;
	mul.wide.u32 	%rd4, %r5, 8;
	add.u64 	%rd5, %rd2, %rd4;
	mul.wide.u32 	%rd6, %r9, 8;
	mov.f64 	%fd1, 0d0000000000000000;	// 0
$Lt_4_21762:
 //<loop> Loop body line 332, nesting depth: 1, estimated iterations: unknown
	.loc	17	338	0
	ld.global.f64 	%fd2, [%rd5+0];
	setp.gt.f64 	%p2, %fd2, %fd1;
	@!%p2 bra 	$Lt_4_22018;
	.loc	17	341	0
	mov.f64 	%fd1, %fd2;
	.loc	17	342	0
	mov.s32 	%r10, %r6;
$Lt_4_22018:
	.loc	17	332	0
	ld.param.s32 	%r1, [__cudaparm_reduce_max_main_blockSize];
	.loc	17	342	0
	add.u32 	%r11, %r1, %r6;
	.loc	17	332	0
	ld.param.u32 	%r7, [__cudaparm_reduce_max_main_n];
	.loc	17	342	0
	setp.ge.u32 	%p3, %r11, %r7;
	@%p3 bra 	$Lt_4_23042;
	.loc	17	347	0
	cvt.u64.u32 	%rd7, %r11;
	mul.wide.u32 	%rd8, %r11, 8;
	.loc	17	332	0
	ld.param.u64 	%rd2, [__cudaparm_reduce_max_main_g_idata];
	.loc	17	347	0
	add.u64 	%rd9, %rd2, %rd8;
	ld.global.f64 	%fd2, [%rd9+0];
	setp.gt.f64 	%p4, %fd2, %fd1;
	@!%p4 bra 	$Lt_4_23042;
	.loc	17	350	0
	mov.f64 	%fd1, %fd2;
	.loc	17	351	0
	mov.s32 	%r10, %r11;
$Lt_4_23042:
$Lt_4_22530:
	add.u32 	%r6, %r9, %r6;
	add.u64 	%rd5, %rd5, %rd6;
	.loc	17	332	0
	ld.param.u32 	%r7, [__cudaparm_reduce_max_main_n];
	.loc	17	351	0
	setp.lt.u32 	%p5, %r6, %r7;
	@%p5 bra 	$Lt_4_21762;
	bra.uni 	$Lt_4_21250;
$Lt_4_35586:
	mov.f64 	%fd1, 0d0000000000000000;	// 0
$Lt_4_21250:
	.loc	17	358	0
	cvt.u64.u32 	%rd10, %r4;
	mov.u64 	%rd11, __cuda___cuda_local_var_17347_33_non_const_sdata19616;
	mul.wide.u32 	%rd12, %r4, 8;
	add.u64 	%rd13, %rd11, %rd12;
	st.shared.f64 	[%rd13+0], %fd1;
	.loc	17	359	0
	mov.u64 	%rd14, __cuda___cuda_local_var_17348_30_non_const_idxData21664;
	mul.wide.u32 	%rd15, %r4, 4;
	add.u64 	%rd16, %rd14, %rd15;
	mov.s32 	%r12, %r10;
	st.shared.s32 	[%rd16+0], %r12;
	.loc	17	361	0
	bar.sync 	0;
	mov.u32 	%r13, 511;
	setp.le.s32 	%p6, %r1, %r13;
	@%p6 bra 	$Lt_4_23810;
	mov.u32 	%r14, 255;
	setp.gt.u32 	%p7, %r4, %r14;
	@%p7 bra 	$Lt_4_24834;
	ld.shared.f64 	%fd3, [%rd13+2048];
	setp.gt.f64 	%p8, %fd3, %fd1;
	@!%p8 bra 	$Lt_4_24834;
	ld.shared.f64 	%fd3, [%rd13+2048];
	.loc	17	369	0
	mov.f64 	%fd1, %fd3;
	st.shared.f64 	[%rd13+0], %fd3;
	.loc	17	370	0
	ld.shared.s32 	%r15, [%rd16+1024];
	st.shared.s32 	[%rd16+0], %r15;
$Lt_4_24834:
$Lt_4_24322:
	.loc	17	373	0
	bar.sync 	0;
$Lt_4_23810:
	mov.u32 	%r16, 255;
	setp.le.s32 	%p9, %r1, %r16;
	@%p9 bra 	$Lt_4_25346;
	mov.u32 	%r17, 127;
	setp.gt.u32 	%p10, %r4, %r17;
	@%p10 bra 	$Lt_4_26370;
	ld.shared.f64 	%fd4, [%rd13+1024];
	setp.gt.f64 	%p11, %fd4, %fd1;
	@!%p11 bra 	$Lt_4_26370;
	ld.shared.f64 	%fd4, [%rd13+1024];
	.loc	17	382	0
	mov.f64 	%fd1, %fd4;
	st.shared.f64 	[%rd13+0], %fd4;
	.loc	17	383	0
	ld.shared.s32 	%r18, [%rd16+512];
	st.shared.s32 	[%rd16+0], %r18;
$Lt_4_26370:
$Lt_4_25858:
	.loc	17	386	0
	bar.sync 	0;
$Lt_4_25346:
	mov.u32 	%r19, 127;
	setp.le.s32 	%p12, %r1, %r19;
	@%p12 bra 	$Lt_4_26882;
	mov.u32 	%r20, 63;
	setp.gt.u32 	%p13, %r4, %r20;
	@%p13 bra 	$Lt_4_27906;
	ld.shared.f64 	%fd5, [%rd13+512];
	setp.gt.f64 	%p14, %fd5, %fd1;
	@!%p14 bra 	$Lt_4_27906;
	ld.shared.f64 	%fd5, [%rd13+512];
	.loc	17	395	0
	mov.f64 	%fd1, %fd5;
	st.shared.f64 	[%rd13+0], %fd5;
	.loc	17	396	0
	ld.shared.s32 	%r21, [%rd16+256];
	st.shared.s32 	[%rd16+0], %r21;
$Lt_4_27906:
$Lt_4_27394:
	.loc	17	399	0
	bar.sync 	0;
$Lt_4_26882:
	mov.u32 	%r22, 31;
	setp.gt.u32 	%p15, %r4, %r22;
	@%p15 bra 	$Lt_4_28418;
	mov.u32 	%r23, 63;
	setp.le.s32 	%p16, %r1, %r23;
	@%p16 bra 	$Lt_4_29442;
	ld.volatile.shared.f64 	%fd6, [%rd13+256];
	setp.gt.f64 	%p17, %fd6, %fd1;
	@!%p17 bra 	$Lt_4_29442;
	.loc	17	410	0
	ld.volatile.shared.f64 	%fd1, [%rd13+256];
	st.volatile.shared.f64 	[%rd13+0], %fd1;
	.loc	17	411	0
	ld.volatile.shared.s32 	%r24, [%rd16+128];
	st.volatile.shared.s32 	[%rd16+0], %r24;
$Lt_4_29442:
$Lt_4_28930:
	mov.u32 	%r25, 31;
	setp.le.s32 	%p18, %r1, %r25;
	@%p18 bra 	$Lt_4_30466;
	ld.volatile.shared.f64 	%fd7, [%rd13+128];
	setp.gt.f64 	%p19, %fd7, %fd1;
	@!%p19 bra 	$Lt_4_30466;
	.loc	17	417	0
	ld.volatile.shared.f64 	%fd1, [%rd13+128];
	st.volatile.shared.f64 	[%rd13+0], %fd1;
	.loc	17	418	0
	ld.volatile.shared.s32 	%r26, [%rd16+64];
	st.volatile.shared.s32 	[%rd16+0], %r26;
$Lt_4_30466:
$Lt_4_29954:
	mov.u32 	%r27, 15;
	setp.le.s32 	%p20, %r1, %r27;
	@%p20 bra 	$Lt_4_31490;
	ld.volatile.shared.f64 	%fd8, [%rd13+64];
	setp.gt.f64 	%p21, %fd8, %fd1;
	@!%p21 bra 	$Lt_4_31490;
	.loc	17	424	0
	ld.volatile.shared.f64 	%fd1, [%rd13+64];
	st.volatile.shared.f64 	[%rd13+0], %fd1;
	.loc	17	425	0
	ld.volatile.shared.s32 	%r28, [%rd16+32];
	st.volatile.shared.s32 	[%rd16+0], %r28;
$Lt_4_31490:
$Lt_4_30978:
	mov.u32 	%r29, 7;
	setp.le.s32 	%p22, %r1, %r29;
	@%p22 bra 	$Lt_4_32514;
	ld.volatile.shared.f64 	%fd9, [%rd13+32];
	setp.gt.f64 	%p23, %fd9, %fd1;
	@!%p23 bra 	$Lt_4_32514;
	.loc	17	431	0
	ld.volatile.shared.f64 	%fd1, [%rd13+32];
	st.volatile.shared.f64 	[%rd13+0], %fd1;
	.loc	17	432	0
	ld.volatile.shared.s32 	%r30, [%rd16+16];
	st.volatile.shared.s32 	[%rd16+0], %r30;
$Lt_4_32514:
$Lt_4_32002:
	mov.u32 	%r31, 3;
	setp.le.s32 	%p24, %r1, %r31;
	@%p24 bra 	$Lt_4_33538;
	ld.volatile.shared.f64 	%fd10, [%rd13+16];
	setp.gt.f64 	%p25, %fd10, %fd1;
	@!%p25 bra 	$Lt_4_33538;
	.loc	17	438	0
	ld.volatile.shared.f64 	%fd1, [%rd13+16];
	st.volatile.shared.f64 	[%rd13+0], %fd1;
	.loc	17	439	0
	ld.volatile.shared.s32 	%r32, [%rd16+8];
	st.volatile.shared.s32 	[%rd16+0], %r32;
$Lt_4_33538:
$Lt_4_33026:
	mov.u32 	%r33, 1;
	setp.le.s32 	%p26, %r1, %r33;
	@%p26 bra 	$Lt_4_34562;
	ld.volatile.shared.f64 	%fd11, [%rd13+8];
	setp.gt.f64 	%p27, %fd11, %fd1;
	@!%p27 bra 	$Lt_4_34562;
	.loc	17	445	0
	ld.volatile.shared.f64 	%fd1, [%rd13+8];
	st.volatile.shared.f64 	[%rd13+0], %fd1;
	.loc	17	446	0
	ld.volatile.shared.s32 	%r34, [%rd16+4];
	st.volatile.shared.s32 	[%rd16+0], %r34;
$Lt_4_34562:
$Lt_4_34050:
	.loc	17	448	0
	bar.sync 	0;
$Lt_4_28418:
	mov.u32 	%r35, 0;
	setp.ne.u32 	%p28, %r4, %r35;
	@%p28 bra 	$Lt_4_35074;
	.loc	17	453	0
	cvt.u64.u32 	%rd17, %r2;
	ld.shared.f64 	%fd12, [__cuda___cuda_local_var_17347_33_non_const_sdata19616+0];
	ld.param.u64 	%rd18, [__cudaparm_reduce_max_main_g_odata];
	mul.wide.u32 	%rd19, %r2, 8;
	add.u64 	%rd20, %rd18, %rd19;
	st.global.f64 	[%rd20+0], %fd12;
	.loc	17	454	0
	ld.shared.s32 	%r36, [__cuda___cuda_local_var_17348_30_non_const_idxData21664+0];
	ld.param.u64 	%rd21, [__cudaparm_reduce_max_main_max_idx];
	mul.wide.u32 	%rd22, %r2, 4;
	add.u64 	%rd23, %rd21, %rd22;
	st.global.s32 	[%rd23+0], %r36;
$Lt_4_35074:
	.loc	17	456	0
	exit;
$LDWend_reduce_max_main:
	} // reduce_max_main

	.entry reduce_max_filter_final (
		.param .u64 __cudaparm_reduce_max_filter_final_g_idata,
		.param .u64 __cudaparm_reduce_max_filter_final_g_odata,
		.param .u64 __cudaparm_reduce_max_filter_final_max_idx,
		.param .u32 __cudaparm_reduce_max_filter_final_n,
		.param .u32 __cudaparm_reduce_max_filter_final_width,
		.param .s32 __cudaparm_reduce_max_filter_final_blockSize,
		.param .u64 __cudaparm_reduce_max_filter_final_maxes,
		.param .s32 __cudaparm_reduce_max_filter_final_nMax)
	{
	.reg .u32 %r<160>;
	.reg .u64 %rd<78>;
	.reg .f64 %fd<70>;
	.reg .pred %p<76>;
	.shared .align 4 .b8 __cuda___cuda_local_var_17488_30_non_const_smaxesRow22740[40];
	.shared .align 4 .b8 __cuda___cuda_local_var_17489_30_non_const_smaxesCol22780[40];
	.shared .align 4 .b8 __cuda___cuda_local_var_17490_30_non_const_smaxesVal22820[40];
	.shared .align 8 .b8 __cuda___cuda_local_var_17491_33_non_const_sdata22864[2048];
	.shared .align 4 .b8 __cuda___cuda_local_var_17492_30_non_const_idxData24912[1024];
	.loc	17	464	0
$LDWbegin_reduce_max_filter_final:
	.loc	17	472	0
	ld.param.s32 	%r1, [__cudaparm_reduce_max_filter_final_blockSize];
	cvt.u32.u16 	%r2, %ctaid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	mul.lo.u32 	%r4, %r3, 2;
	cvt.u32.u16 	%r5, %tid.x;
	add.u32 	%r6, %r4, %r5;
	mov.s32 	%r7, %r6;
	cvt.u64.u32 	%rd1, %r5;
	mul.wide.u32 	%rd2, %r5, 4;
	ld.param.s32 	%r8, [__cudaparm_reduce_max_filter_final_nMax];
	setp.le.u32 	%p1, %r8, %r5;
	@%p1 bra 	$Lt_5_56578;
	.loc	17	477	0
	mov.u64 	%rd3, __cuda___cuda_local_var_17488_30_non_const_smaxesRow22740;
	mov.u64 	%rd4, __cuda___cuda_local_var_17489_30_non_const_smaxesCol22780;
	ld.param.u64 	%rd5, [__cudaparm_reduce_max_filter_final_maxes];
	add.u64 	%rd6, %rd5, %rd2;
	ld.global.s32 	%r9, [%rd6+0];
	mov.u64 	%rd7, __cuda___cuda_local_var_17490_30_non_const_smaxesVal22820;
	add.u64 	%rd8, %rd2, %rd7;
	st.shared.s32 	[%rd8+0], %r9;
	.loc	17	478	0
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	div.u32 	%r11, %r9, %r10;
	add.u64 	%rd9, %rd2, %rd3;
	st.shared.s32 	[%rd9+0], %r11;
	.loc	17	479	0
	rem.u32 	%r12, %r9, %r10;
	add.u64 	%rd10, %rd2, %rd4;
	st.shared.s32 	[%rd10+0], %r12;
$Lt_5_56578:
	mov.u64 	%rd3, __cuda___cuda_local_var_17488_30_non_const_smaxesRow22740;
	mov.u64 	%rd4, __cuda___cuda_local_var_17489_30_non_const_smaxesCol22780;
	.loc	17	481	0
	bar.sync 	0;
	ld.param.u32 	%r13, [__cudaparm_reduce_max_filter_final_n];
	setp.ge.u32 	%p2, %r6, %r13;
	@%p2 bra 	$Lt_5_94466;
	cvt.u64.u32 	%rd11, %r6;
	add.u32 	%r14, %r1, %r6;
	cvt.u32.u16 	%r15, %nctaid.x;
	mul.lo.u32 	%r16, %r1, %r15;
	mul.lo.u32 	%r17, %r16, 2;
	ld.param.u64 	%rd12, [__cudaparm_reduce_max_filter_final_g_idata];
	mul.wide.u32 	%rd13, %r6, 8;
	add.u64 	%rd14, %rd12, %rd13;
	cvt.s64.u32 	%rd15, %r17;
	mov.f64 	%fd1, 0d0000000000000000;	// 0
$Lt_5_57602:
 //<loop> Loop body line 481, nesting depth: 1, estimated iterations: unknown
	ld.global.f64 	%fd2, [%rd14+0];
	setp.gt.f64 	%p3, %fd2, %fd1;
	@!%p3 bra 	$Lt_5_59906;
	.loc	17	490	0
	ld.param.u64 	%rd16, [__cudaparm_reduce_max_filter_final_max_idx];
	mul.lo.u64 	%rd17, %rd11, 4;
	add.u64 	%rd18, %rd16, %rd17;
	ld.global.s32 	%r18, [%rd18+0];
	mov.u32 	%r19, 0;
	setp.le.s32 	%p4, %r8, %r19;
	@%p4 bra 	$Lt_5_58370;
	mov.s32 	%r20, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	rem.s32 	%r21, %r18, %r10;
	div.s32 	%r22, %r18, %r10;
	mov.s32 	%r23, 0;
	mov.s32 	%r24, %r20;
$Lt_5_58882:
 //<loop> Loop body line 490, nesting depth: 2, estimated iterations: unknown
	.loc	17	53	0
	cvt.s64.s32 	%rd19, %r23;
	mul.wide.s32 	%rd20, %r23, 4;
	add.u64 	%rd21, %rd20, %rd3;
	ld.shared.s32 	%r25, [%rd21+0];
	add.u64 	%rd22, %rd20, %rd4;
	ld.shared.s32 	%r26, [%rd22+0];
	sub.s32 	%r27, %r26, %r21;
	cvt.rn.f64.s32 	%fd3, %r27;
	sub.s32 	%r28, %r25, %r22;
	cvt.rn.f64.s32 	%fd4, %r28;
	mul.f64 	%fd5, %fd3, %fd3;
	mad.rn.f64 	%fd6, %fd4, %fd4, %fd5;
	mov.f64 	%fd7, 0d4059000000000000;	// 100
	setp.lt.f64 	%p5, %fd6, %fd7;
	@!%p5 bra 	$Lt_5_59138;
	.loc	17	56	0
	mov.s32 	%r29, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_21;
$Lt_5_59138:
	.loc	17	51	0
	add.s32 	%r23, %r23, 1;
	setp.ne.s32 	%p6, %r23, %r8;
	@%p6 bra 	$Lt_5_58882;
$Lt_5_58370:
	.loc	17	61	0
	mov.s32 	%r29, 1;
$LDWendi__Z8distanceiiii_335_21:
	.loc	17	490	0
	mov.u32 	%r30, 0;
	setp.eq.s32 	%p7, %r29, %r30;
	@%p7 bra 	$Lt_5_59906;
	.loc	17	493	0
	mov.f64 	%fd1, %fd2;
	.loc	17	494	0
	mov.s32 	%r31, %r18;
$Lt_5_59906:
$Lt_5_57858:
	.loc	17	481	0
	ld.param.u32 	%r13, [__cudaparm_reduce_max_filter_final_n];
	.loc	17	494	0
	setp.ge.u32 	%p8, %r14, %r13;
	@%p8 bra 	$Lt_5_62978;
	cvt.u64.u32 	%rd23, %r14;
	mul.wide.u32 	%rd24, %r14, 8;
	.loc	17	481	0
	ld.param.u64 	%rd12, [__cudaparm_reduce_max_filter_final_g_idata];
	.loc	17	494	0
	add.u64 	%rd25, %rd12, %rd24;
	ld.global.f64 	%fd8, [%rd25+0];
	setp.gt.f64 	%p9, %fd8, %fd1;
	@!%p9 bra 	$Lt_5_62978;
	.loc	17	502	0
	ld.param.u64 	%rd26, [__cudaparm_reduce_max_filter_final_max_idx];
	mul.lo.u64 	%rd27, %rd23, 4;
	add.u64 	%rd28, %rd26, %rd27;
	ld.global.s32 	%r32, [%rd28+0];
	mov.u32 	%r33, 0;
	setp.le.s32 	%p10, %r8, %r33;
	@%p10 bra 	$Lt_5_61442;
	mov.s32 	%r34, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	rem.s32 	%r35, %r32, %r10;
	div.s32 	%r36, %r32, %r10;
	mov.s32 	%r23, 0;
	mov.s32 	%r37, %r34;
$Lt_5_61954:
 //<loop> Loop body line 502, nesting depth: 2, estimated iterations: unknown
	.loc	17	53	0
	cvt.s64.s32 	%rd29, %r23;
	mul.wide.s32 	%rd20, %r23, 4;
	add.u64 	%rd30, %rd20, %rd3;
	ld.shared.s32 	%r25, [%rd30+0];
	add.u64 	%rd31, %rd20, %rd4;
	ld.shared.s32 	%r26, [%rd31+0];
	sub.s32 	%r38, %r26, %r35;
	cvt.rn.f64.s32 	%fd9, %r38;
	sub.s32 	%r39, %r25, %r36;
	cvt.rn.f64.s32 	%fd10, %r39;
	mul.f64 	%fd11, %fd9, %fd9;
	mad.rn.f64 	%fd12, %fd10, %fd10, %fd11;
	mov.f64 	%fd13, 0d4059000000000000;	// 100
	setp.lt.f64 	%p11, %fd12, %fd13;
	@!%p11 bra 	$Lt_5_62210;
	.loc	17	56	0
	mov.s32 	%r40, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_19;
$Lt_5_62210:
	.loc	17	51	0
	add.s32 	%r23, %r23, 1;
	setp.ne.s32 	%p12, %r23, %r8;
	@%p12 bra 	$Lt_5_61954;
$Lt_5_61442:
	.loc	17	61	0
	mov.s32 	%r40, 1;
$LDWendi__Z8distanceiiii_335_19:
	.loc	17	502	0
	mov.u32 	%r41, 0;
	setp.eq.s32 	%p13, %r40, %r41;
	@%p13 bra 	$Lt_5_62978;
	.loc	17	508	0
	mov.f64 	%fd1, %fd8;
	.loc	17	509	0
	mov.s32 	%r31, %r32;
$Lt_5_62978:
$Lt_5_60930:
$Lt_5_60418:
	.loc	17	514	0
	add.u32 	%r7, %r17, %r7;
	add.u32 	%r14, %r14, %r17;
	add.u64 	%rd11, %rd15, %rd11;
	mul.lo.u64 	%rd32, %rd15, 8;
	add.u64 	%rd14, %rd14, %rd32;
	.loc	17	481	0
	ld.param.u32 	%r13, [__cudaparm_reduce_max_filter_final_n];
	.loc	17	514	0
	setp.lt.u32 	%p14, %r7, %r13;
	@%p14 bra 	$Lt_5_57602;
	bra.uni 	$Lt_5_57090;
$Lt_5_94466:
	mov.f64 	%fd1, 0d0000000000000000;	// 0
$Lt_5_57090:
	.loc	17	517	0
	mov.u64 	%rd33, __cuda___cuda_local_var_17491_33_non_const_sdata22864;
	mul.lo.u64 	%rd34, %rd1, 8;
	add.u64 	%rd35, %rd33, %rd34;
	st.shared.f64 	[%rd35+0], %fd1;
	.loc	17	518	0
	mov.u64 	%rd36, __cuda___cuda_local_var_17492_30_non_const_idxData24912;
	add.u64 	%rd37, %rd2, %rd36;
	mov.s32 	%r42, %r31;
	st.shared.s32 	[%rd37+0], %r42;
	.loc	17	520	0
	bar.sync 	0;
	mov.u32 	%r43, 511;
	setp.le.s32 	%p15, %r1, %r43;
	@%p15 bra 	$Lt_5_63746;
	mov.u32 	%r44, 255;
	setp.gt.u32 	%p16, %r5, %r44;
	@%p16 bra 	$Lt_5_66818;
	ld.shared.f64 	%fd14, [%rd35+2048];
	setp.gt.f64 	%p17, %fd14, %fd1;
	@!%p17 bra 	$Lt_5_66818;
	.loc	17	528	0
	ld.shared.s32 	%r45, [%rd37+1024];
	mov.u32 	%r46, 0;
	setp.le.s32 	%p18, %r8, %r46;
	@%p18 bra 	$Lt_5_65282;
	mov.s32 	%r47, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	ld.shared.s32 	%r45, [%rd37+1024];
	rem.s32 	%r48, %r45, %r10;
	div.s32 	%r49, %r45, %r10;
	mov.s32 	%r23, 0;
	mov.s32 	%r50, %r47;
$Lt_5_65794:
 //<loop> Loop body line 528, nesting depth: 1, estimated iterations: unknown
	.loc	17	53	0
	cvt.s64.s32 	%rd38, %r23;
	mul.wide.s32 	%rd20, %r23, 4;
	add.u64 	%rd39, %rd20, %rd3;
	ld.shared.s32 	%r25, [%rd39+0];
	add.u64 	%rd40, %rd20, %rd4;
	ld.shared.s32 	%r26, [%rd40+0];
	sub.s32 	%r51, %r26, %r48;
	cvt.rn.f64.s32 	%fd15, %r51;
	sub.s32 	%r52, %r25, %r49;
	cvt.rn.f64.s32 	%fd16, %r52;
	mul.f64 	%fd17, %fd15, %fd15;
	mad.rn.f64 	%fd18, %fd16, %fd16, %fd17;
	mov.f64 	%fd19, 0d4059000000000000;	// 100
	setp.lt.f64 	%p19, %fd18, %fd19;
	@!%p19 bra 	$Lt_5_66050;
	.loc	17	56	0
	mov.s32 	%r53, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_17;
$Lt_5_66050:
	.loc	17	51	0
	add.s32 	%r23, %r23, 1;
	setp.ne.s32 	%p20, %r23, %r8;
	@%p20 bra 	$Lt_5_65794;
$Lt_5_65282:
	.loc	17	61	0
	mov.s32 	%r53, 1;
$LDWendi__Z8distanceiiii_335_17:
	.loc	17	528	0
	mov.u32 	%r54, 0;
	setp.eq.s32 	%p21, %r53, %r54;
	@%p21 bra 	$Lt_5_66818;
	.loc	17	520	0
	ld.shared.f64 	%fd14, [%rd35+2048];
	.loc	17	532	0
	mov.f64 	%fd1, %fd14;
	st.shared.f64 	[%rd35+0], %fd14;
	.loc	17	528	0
	ld.shared.s32 	%r45, [%rd37+1024];
	.loc	17	533	0
	st.shared.s32 	[%rd37+0], %r45;
$Lt_5_66818:
$Lt_5_64770:
$Lt_5_64258:
	.loc	17	537	0
	bar.sync 	0;
$Lt_5_63746:
	mov.u32 	%r55, 255;
	setp.le.s32 	%p22, %r1, %r55;
	@%p22 bra 	$Lt_5_67330;
	mov.u32 	%r56, 127;
	setp.gt.u32 	%p23, %r5, %r56;
	@%p23 bra 	$Lt_5_70402;
	ld.shared.f64 	%fd20, [%rd35+1024];
	setp.gt.f64 	%p24, %fd20, %fd1;
	@!%p24 bra 	$Lt_5_70402;
	.loc	17	546	0
	ld.shared.s32 	%r57, [%rd37+512];
	mov.u32 	%r58, 0;
	setp.le.s32 	%p25, %r8, %r58;
	@%p25 bra 	$Lt_5_68866;
	mov.s32 	%r59, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	ld.shared.s32 	%r57, [%rd37+512];
	rem.s32 	%r60, %r57, %r10;
	div.s32 	%r61, %r57, %r10;
	mov.s32 	%r23, 0;
	mov.s32 	%r62, %r59;
$Lt_5_69378:
 //<loop> Loop body line 546, nesting depth: 1, estimated iterations: unknown
	.loc	17	53	0
	cvt.s64.s32 	%rd41, %r23;
	mul.wide.s32 	%rd20, %r23, 4;
	add.u64 	%rd42, %rd20, %rd3;
	ld.shared.s32 	%r25, [%rd42+0];
	add.u64 	%rd43, %rd20, %rd4;
	ld.shared.s32 	%r26, [%rd43+0];
	sub.s32 	%r63, %r26, %r60;
	cvt.rn.f64.s32 	%fd21, %r63;
	sub.s32 	%r64, %r25, %r61;
	cvt.rn.f64.s32 	%fd22, %r64;
	mul.f64 	%fd23, %fd21, %fd21;
	mad.rn.f64 	%fd24, %fd22, %fd22, %fd23;
	mov.f64 	%fd25, 0d4059000000000000;	// 100
	setp.lt.f64 	%p26, %fd24, %fd25;
	@!%p26 bra 	$Lt_5_69634;
	.loc	17	56	0
	mov.s32 	%r65, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_15;
$Lt_5_69634:
	.loc	17	51	0
	add.s32 	%r23, %r23, 1;
	setp.ne.s32 	%p27, %r23, %r8;
	@%p27 bra 	$Lt_5_69378;
$Lt_5_68866:
	.loc	17	61	0
	mov.s32 	%r65, 1;
$LDWendi__Z8distanceiiii_335_15:
	.loc	17	546	0
	mov.u32 	%r66, 0;
	setp.eq.s32 	%p28, %r65, %r66;
	@%p28 bra 	$Lt_5_70402;
	.loc	17	537	0
	ld.shared.f64 	%fd20, [%rd35+1024];
	.loc	17	550	0
	mov.f64 	%fd1, %fd20;
	st.shared.f64 	[%rd35+0], %fd20;
	.loc	17	546	0
	ld.shared.s32 	%r57, [%rd37+512];
	.loc	17	551	0
	st.shared.s32 	[%rd37+0], %r57;
$Lt_5_70402:
$Lt_5_68354:
$Lt_5_67842:
	.loc	17	555	0
	bar.sync 	0;
$Lt_5_67330:
	mov.u32 	%r67, 127;
	setp.le.s32 	%p29, %r1, %r67;
	@%p29 bra 	$Lt_5_70914;
	mov.u32 	%r68, 63;
	setp.gt.u32 	%p30, %r5, %r68;
	@%p30 bra 	$Lt_5_73986;
	ld.shared.f64 	%fd26, [%rd35+512];
	setp.gt.f64 	%p31, %fd26, %fd1;
	@!%p31 bra 	$Lt_5_73986;
	.loc	17	564	0
	ld.shared.s32 	%r69, [%rd37+256];
	mov.u32 	%r70, 0;
	setp.le.s32 	%p32, %r8, %r70;
	@%p32 bra 	$Lt_5_72450;
	mov.s32 	%r71, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	ld.shared.s32 	%r69, [%rd37+256];
	rem.s32 	%r72, %r69, %r10;
	div.s32 	%r73, %r69, %r10;
	mov.s32 	%r23, 0;
	mov.s32 	%r74, %r71;
$Lt_5_72962:
 //<loop> Loop body line 564, nesting depth: 1, estimated iterations: unknown
	.loc	17	53	0
	cvt.s64.s32 	%rd44, %r23;
	mul.wide.s32 	%rd20, %r23, 4;
	add.u64 	%rd45, %rd20, %rd3;
	ld.shared.s32 	%r25, [%rd45+0];
	add.u64 	%rd46, %rd20, %rd4;
	ld.shared.s32 	%r26, [%rd46+0];
	sub.s32 	%r75, %r26, %r72;
	cvt.rn.f64.s32 	%fd27, %r75;
	sub.s32 	%r76, %r25, %r73;
	cvt.rn.f64.s32 	%fd28, %r76;
	mul.f64 	%fd29, %fd27, %fd27;
	mad.rn.f64 	%fd30, %fd28, %fd28, %fd29;
	mov.f64 	%fd31, 0d4059000000000000;	// 100
	setp.lt.f64 	%p33, %fd30, %fd31;
	@!%p33 bra 	$Lt_5_73218;
	.loc	17	56	0
	mov.s32 	%r77, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_13;
$Lt_5_73218:
	.loc	17	51	0
	add.s32 	%r23, %r23, 1;
	setp.ne.s32 	%p34, %r23, %r8;
	@%p34 bra 	$Lt_5_72962;
$Lt_5_72450:
	.loc	17	61	0
	mov.s32 	%r77, 1;
$LDWendi__Z8distanceiiii_335_13:
	.loc	17	564	0
	mov.u32 	%r78, 0;
	setp.eq.s32 	%p35, %r77, %r78;
	@%p35 bra 	$Lt_5_73986;
	.loc	17	555	0
	ld.shared.f64 	%fd26, [%rd35+512];
	.loc	17	568	0
	mov.f64 	%fd1, %fd26;
	st.shared.f64 	[%rd35+0], %fd26;
	.loc	17	564	0
	ld.shared.s32 	%r69, [%rd37+256];
	.loc	17	569	0
	st.shared.s32 	[%rd37+0], %r69;
$Lt_5_73986:
$Lt_5_71938:
$Lt_5_71426:
	.loc	17	573	0
	bar.sync 	0;
$Lt_5_70914:
	mov.u32 	%r79, 31;
	setp.gt.u32 	%p36, %r5, %r79;
	@%p36 bra 	$Lt_5_74498;
	mov.u32 	%r80, 63;
	setp.le.s32 	%p37, %r1, %r80;
	@%p37 bra 	$Lt_5_77570;
	ld.volatile.shared.f64 	%fd32, [%rd35+256];
	setp.gt.f64 	%p38, %fd32, %fd1;
	@!%p38 bra 	$Lt_5_77570;
	.loc	17	587	0
	ld.volatile.shared.s32 	%r81, [%rd37+128];
	mov.u32 	%r82, 0;
	setp.le.s32 	%p39, %r8, %r82;
	@%p39 bra 	$Lt_5_76034;
	mov.s32 	%r83, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	ld.volatile.shared.s32 	%r81, [%rd37+128];
	rem.s32 	%r84, %r81, %r10;
	div.s32 	%r85, %r81, %r10;
	mov.s32 	%r86, 0;
	mov.s32 	%r87, %r83;
$Lt_5_76546:
 //<loop> Loop body line 587, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd47, %r86;
	mul.wide.s32 	%rd48, %r86, 4;
	add.u64 	%rd49, %rd48, %rd3;
	ld.volatile.shared.s32 	%r88, [%rd49+0];
	add.u64 	%rd50, %rd48, %rd4;
	ld.volatile.shared.s32 	%r89, [%rd50+0];
	sub.s32 	%r90, %r89, %r84;
	cvt.rn.f64.s32 	%fd33, %r90;
	sub.s32 	%r91, %r88, %r85;
	cvt.rn.f64.s32 	%fd34, %r91;
	mul.f64 	%fd35, %fd33, %fd33;
	mad.rn.f64 	%fd36, %fd34, %fd34, %fd35;
	mov.f64 	%fd37, 0d4059000000000000;	// 100
	setp.lt.f64 	%p40, %fd36, %fd37;
	@!%p40 bra 	$Lt_5_76802;
	.loc	17	77	0
	mov.s32 	%r92, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_11;
$Lt_5_76802:
	.loc	17	72	0
	add.s32 	%r86, %r86, 1;
	setp.ne.s32 	%p41, %r86, %r8;
	@%p41 bra 	$Lt_5_76546;
$Lt_5_76034:
	.loc	17	82	0
	mov.s32 	%r92, 1;
$LDWendi__Z8distanceiiii_335_11:
	.loc	17	587	0
	mov.u32 	%r93, 0;
	setp.eq.s32 	%p42, %r92, %r93;
	@%p42 bra 	$Lt_5_77570;
	.loc	17	591	0
	ld.volatile.shared.f64 	%fd1, [%rd35+256];
	st.volatile.shared.f64 	[%rd35+0], %fd1;
	.loc	17	592	0
	ld.volatile.shared.s32 	%r94, [%rd37+128];
	st.volatile.shared.s32 	[%rd37+0], %r94;
$Lt_5_77570:
$Lt_5_75522:
$Lt_5_75010:
	mov.u32 	%r95, 31;
	setp.le.s32 	%p43, %r1, %r95;
	@%p43 bra 	$Lt_5_80642;
	ld.volatile.shared.f64 	%fd38, [%rd35+128];
	setp.gt.f64 	%p44, %fd38, %fd1;
	@!%p44 bra 	$Lt_5_80642;
	.loc	17	599	0
	ld.volatile.shared.s32 	%r96, [%rd37+64];
	mov.u32 	%r97, 0;
	setp.le.s32 	%p45, %r8, %r97;
	@%p45 bra 	$Lt_5_79106;
	mov.s32 	%r98, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	ld.volatile.shared.s32 	%r96, [%rd37+64];
	rem.s32 	%r99, %r96, %r10;
	div.s32 	%r100, %r96, %r10;
	mov.s32 	%r86, 0;
	mov.s32 	%r101, %r98;
$Lt_5_79618:
 //<loop> Loop body line 599, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd51, %r86;
	mul.wide.s32 	%rd48, %r86, 4;
	add.u64 	%rd52, %rd48, %rd3;
	ld.volatile.shared.s32 	%r88, [%rd52+0];
	add.u64 	%rd53, %rd48, %rd4;
	ld.volatile.shared.s32 	%r89, [%rd53+0];
	sub.s32 	%r102, %r89, %r99;
	cvt.rn.f64.s32 	%fd39, %r102;
	sub.s32 	%r103, %r88, %r100;
	cvt.rn.f64.s32 	%fd40, %r103;
	mul.f64 	%fd41, %fd39, %fd39;
	mad.rn.f64 	%fd42, %fd40, %fd40, %fd41;
	mov.f64 	%fd43, 0d4059000000000000;	// 100
	setp.lt.f64 	%p46, %fd42, %fd43;
	@!%p46 bra 	$Lt_5_79874;
	.loc	17	77	0
	mov.s32 	%r104, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_9;
$Lt_5_79874:
	.loc	17	72	0
	add.s32 	%r86, %r86, 1;
	setp.ne.s32 	%p47, %r86, %r8;
	@%p47 bra 	$Lt_5_79618;
$Lt_5_79106:
	.loc	17	82	0
	mov.s32 	%r104, 1;
$LDWendi__Z8distanceiiii_335_9:
	.loc	17	599	0
	mov.u32 	%r105, 0;
	setp.eq.s32 	%p48, %r104, %r105;
	@%p48 bra 	$Lt_5_80642;
	.loc	17	603	0
	ld.volatile.shared.f64 	%fd1, [%rd35+128];
	st.volatile.shared.f64 	[%rd35+0], %fd1;
	.loc	17	604	0
	ld.volatile.shared.s32 	%r106, [%rd37+64];
	st.volatile.shared.s32 	[%rd37+0], %r106;
$Lt_5_80642:
$Lt_5_78594:
$Lt_5_78082:
	mov.u32 	%r107, 15;
	setp.le.s32 	%p49, %r1, %r107;
	@%p49 bra 	$Lt_5_83714;
	ld.volatile.shared.f64 	%fd44, [%rd35+64];
	setp.gt.f64 	%p50, %fd44, %fd1;
	@!%p50 bra 	$Lt_5_83714;
	.loc	17	611	0
	ld.volatile.shared.s32 	%r108, [%rd37+32];
	mov.u32 	%r109, 0;
	setp.le.s32 	%p51, %r8, %r109;
	@%p51 bra 	$Lt_5_82178;
	mov.s32 	%r110, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	ld.volatile.shared.s32 	%r108, [%rd37+32];
	rem.s32 	%r111, %r108, %r10;
	div.s32 	%r112, %r108, %r10;
	mov.s32 	%r86, 0;
	mov.s32 	%r113, %r110;
$Lt_5_82690:
 //<loop> Loop body line 611, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd54, %r86;
	mul.wide.s32 	%rd48, %r86, 4;
	add.u64 	%rd55, %rd48, %rd3;
	ld.volatile.shared.s32 	%r88, [%rd55+0];
	add.u64 	%rd56, %rd48, %rd4;
	ld.volatile.shared.s32 	%r89, [%rd56+0];
	sub.s32 	%r114, %r89, %r111;
	cvt.rn.f64.s32 	%fd45, %r114;
	sub.s32 	%r115, %r88, %r112;
	cvt.rn.f64.s32 	%fd46, %r115;
	mul.f64 	%fd47, %fd45, %fd45;
	mad.rn.f64 	%fd48, %fd46, %fd46, %fd47;
	mov.f64 	%fd49, 0d4059000000000000;	// 100
	setp.lt.f64 	%p52, %fd48, %fd49;
	@!%p52 bra 	$Lt_5_82946;
	.loc	17	77	0
	mov.s32 	%r116, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_7;
$Lt_5_82946:
	.loc	17	72	0
	add.s32 	%r86, %r86, 1;
	setp.ne.s32 	%p53, %r86, %r8;
	@%p53 bra 	$Lt_5_82690;
$Lt_5_82178:
	.loc	17	82	0
	mov.s32 	%r116, 1;
$LDWendi__Z8distanceiiii_335_7:
	.loc	17	611	0
	mov.u32 	%r117, 0;
	setp.eq.s32 	%p54, %r116, %r117;
	@%p54 bra 	$Lt_5_83714;
	.loc	17	615	0
	ld.volatile.shared.f64 	%fd1, [%rd35+64];
	st.volatile.shared.f64 	[%rd35+0], %fd1;
	.loc	17	616	0
	ld.volatile.shared.s32 	%r118, [%rd37+32];
	st.volatile.shared.s32 	[%rd37+0], %r118;
$Lt_5_83714:
$Lt_5_81666:
$Lt_5_81154:
	mov.u32 	%r119, 7;
	setp.le.s32 	%p55, %r1, %r119;
	@%p55 bra 	$Lt_5_86786;
	ld.volatile.shared.f64 	%fd50, [%rd35+32];
	setp.gt.f64 	%p56, %fd50, %fd1;
	@!%p56 bra 	$Lt_5_86786;
	.loc	17	623	0
	ld.volatile.shared.s32 	%r120, [%rd37+16];
	mov.u32 	%r121, 0;
	setp.le.s32 	%p57, %r8, %r121;
	@%p57 bra 	$Lt_5_85250;
	mov.s32 	%r122, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	ld.volatile.shared.s32 	%r120, [%rd37+16];
	rem.s32 	%r123, %r120, %r10;
	div.s32 	%r124, %r120, %r10;
	mov.s32 	%r86, 0;
	mov.s32 	%r125, %r122;
$Lt_5_85762:
 //<loop> Loop body line 623, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd57, %r86;
	mul.wide.s32 	%rd48, %r86, 4;
	add.u64 	%rd58, %rd48, %rd3;
	ld.volatile.shared.s32 	%r88, [%rd58+0];
	add.u64 	%rd59, %rd48, %rd4;
	ld.volatile.shared.s32 	%r89, [%rd59+0];
	sub.s32 	%r126, %r89, %r123;
	cvt.rn.f64.s32 	%fd51, %r126;
	sub.s32 	%r127, %r88, %r124;
	cvt.rn.f64.s32 	%fd52, %r127;
	mul.f64 	%fd53, %fd51, %fd51;
	mad.rn.f64 	%fd54, %fd52, %fd52, %fd53;
	mov.f64 	%fd55, 0d4059000000000000;	// 100
	setp.lt.f64 	%p58, %fd54, %fd55;
	@!%p58 bra 	$Lt_5_86018;
	.loc	17	77	0
	mov.s32 	%r128, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_5;
$Lt_5_86018:
	.loc	17	72	0
	add.s32 	%r86, %r86, 1;
	setp.ne.s32 	%p59, %r86, %r8;
	@%p59 bra 	$Lt_5_85762;
$Lt_5_85250:
	.loc	17	82	0
	mov.s32 	%r128, 1;
$LDWendi__Z8distanceiiii_335_5:
	.loc	17	623	0
	mov.u32 	%r129, 0;
	setp.eq.s32 	%p60, %r128, %r129;
	@%p60 bra 	$Lt_5_86786;
	.loc	17	627	0
	ld.volatile.shared.f64 	%fd1, [%rd35+32];
	st.volatile.shared.f64 	[%rd35+0], %fd1;
	.loc	17	628	0
	ld.volatile.shared.s32 	%r130, [%rd37+16];
	st.volatile.shared.s32 	[%rd37+0], %r130;
$Lt_5_86786:
$Lt_5_84738:
$Lt_5_84226:
	mov.u32 	%r131, 3;
	setp.le.s32 	%p61, %r1, %r131;
	@%p61 bra 	$Lt_5_89858;
	ld.volatile.shared.f64 	%fd56, [%rd35+16];
	setp.gt.f64 	%p62, %fd56, %fd1;
	@!%p62 bra 	$Lt_5_89858;
	.loc	17	635	0
	ld.volatile.shared.s32 	%r132, [%rd37+8];
	mov.u32 	%r133, 0;
	setp.le.s32 	%p63, %r8, %r133;
	@%p63 bra 	$Lt_5_88322;
	mov.s32 	%r134, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	ld.volatile.shared.s32 	%r132, [%rd37+8];
	rem.s32 	%r135, %r132, %r10;
	div.s32 	%r136, %r132, %r10;
	mov.s32 	%r86, 0;
	mov.s32 	%r137, %r134;
$Lt_5_88834:
 //<loop> Loop body line 635, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd60, %r86;
	mul.wide.s32 	%rd48, %r86, 4;
	add.u64 	%rd61, %rd48, %rd3;
	ld.volatile.shared.s32 	%r88, [%rd61+0];
	add.u64 	%rd62, %rd48, %rd4;
	ld.volatile.shared.s32 	%r89, [%rd62+0];
	sub.s32 	%r138, %r89, %r135;
	cvt.rn.f64.s32 	%fd57, %r138;
	sub.s32 	%r139, %r88, %r136;
	cvt.rn.f64.s32 	%fd58, %r139;
	mul.f64 	%fd59, %fd57, %fd57;
	mad.rn.f64 	%fd60, %fd58, %fd58, %fd59;
	mov.f64 	%fd61, 0d4059000000000000;	// 100
	setp.lt.f64 	%p64, %fd60, %fd61;
	@!%p64 bra 	$Lt_5_89090;
	.loc	17	77	0
	mov.s32 	%r140, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_3;
$Lt_5_89090:
	.loc	17	72	0
	add.s32 	%r86, %r86, 1;
	setp.ne.s32 	%p65, %r86, %r8;
	@%p65 bra 	$Lt_5_88834;
$Lt_5_88322:
	.loc	17	82	0
	mov.s32 	%r140, 1;
$LDWendi__Z8distanceiiii_335_3:
	.loc	17	635	0
	mov.u32 	%r141, 0;
	setp.eq.s32 	%p66, %r140, %r141;
	@%p66 bra 	$Lt_5_89858;
	.loc	17	639	0
	ld.volatile.shared.f64 	%fd1, [%rd35+16];
	st.volatile.shared.f64 	[%rd35+0], %fd1;
	.loc	17	640	0
	ld.volatile.shared.s32 	%r142, [%rd37+8];
	st.volatile.shared.s32 	[%rd37+0], %r142;
$Lt_5_89858:
$Lt_5_87810:
$Lt_5_87298:
	mov.u32 	%r143, 1;
	setp.le.s32 	%p67, %r1, %r143;
	@%p67 bra 	$Lt_5_92930;
	ld.volatile.shared.f64 	%fd62, [%rd35+8];
	setp.gt.f64 	%p68, %fd62, %fd1;
	@!%p68 bra 	$Lt_5_92930;
	.loc	17	647	0
	ld.volatile.shared.s32 	%r144, [%rd37+4];
	mov.u32 	%r145, 0;
	setp.le.s32 	%p69, %r8, %r145;
	@%p69 bra 	$Lt_5_91394;
	mov.s32 	%r146, %r8;
	ld.param.u32 	%r10, [__cudaparm_reduce_max_filter_final_width];
	ld.volatile.shared.s32 	%r144, [%rd37+4];
	rem.s32 	%r147, %r144, %r10;
	div.s32 	%r148, %r144, %r10;
	mov.s32 	%r86, 0;
	mov.s32 	%r149, %r146;
$Lt_5_91906:
 //<loop> Loop body line 647, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd63, %r86;
	mul.wide.s32 	%rd48, %r86, 4;
	add.u64 	%rd64, %rd48, %rd3;
	ld.volatile.shared.s32 	%r88, [%rd64+0];
	add.u64 	%rd65, %rd48, %rd4;
	ld.volatile.shared.s32 	%r89, [%rd65+0];
	sub.s32 	%r150, %r89, %r147;
	cvt.rn.f64.s32 	%fd63, %r150;
	sub.s32 	%r151, %r88, %r148;
	cvt.rn.f64.s32 	%fd64, %r151;
	mul.f64 	%fd65, %fd63, %fd63;
	mad.rn.f64 	%fd66, %fd64, %fd64, %fd65;
	mov.f64 	%fd67, 0d4059000000000000;	// 100
	setp.lt.f64 	%p70, %fd66, %fd67;
	@!%p70 bra 	$Lt_5_92162;
	.loc	17	77	0
	mov.s32 	%r152, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_335_1;
$Lt_5_92162:
	.loc	17	72	0
	add.s32 	%r86, %r86, 1;
	setp.ne.s32 	%p71, %r86, %r8;
	@%p71 bra 	$Lt_5_91906;
$Lt_5_91394:
	.loc	17	82	0
	mov.s32 	%r152, 1;
$LDWendi__Z8distanceiiii_335_1:
	.loc	17	647	0
	mov.u32 	%r153, 0;
	setp.eq.s32 	%p72, %r152, %r153;
	@%p72 bra 	$Lt_5_92930;
	.loc	17	651	0
	ld.volatile.shared.f64 	%fd1, [%rd35+8];
	st.volatile.shared.f64 	[%rd35+0], %fd1;
	.loc	17	652	0
	ld.volatile.shared.s32 	%r154, [%rd37+4];
	st.volatile.shared.s32 	[%rd37+0], %r154;
$Lt_5_92930:
$Lt_5_90882:
$Lt_5_90370:
	.loc	17	656	0
	bar.sync 	0;
$Lt_5_74498:
	mov.u32 	%r155, 0;
	setp.ne.u32 	%p73, %r5, %r155;
	@%p73 bra 	$Lt_5_93954;
	.loc	17	662	0
	cvt.u64.u32 	%rd66, %r2;
	ld.shared.f64 	%fd68, [__cuda___cuda_local_var_17491_33_non_const_sdata22864+0];
	ld.param.u64 	%rd67, [__cudaparm_reduce_max_filter_final_g_odata];
	mul.wide.u32 	%rd68, %r2, 8;
	add.u64 	%rd69, %rd67, %rd68;
	st.global.f64 	[%rd69+0], %fd68;
	.loc	17	663	0
	ld.shared.s32 	%r156, [__cuda___cuda_local_var_17492_30_non_const_idxData24912+0];
	ld.param.u64 	%rd70, [__cudaparm_reduce_max_filter_final_max_idx];
	mul.wide.u32 	%rd71, %r2, 4;
	add.u64 	%rd72, %rd70, %rd71;
	st.global.s32 	[%rd72+0], %r156;
	cvt.u32.u16 	%r157, %nctaid.x;
	mov.u32 	%r158, 1;
	setp.ne.u32 	%p74, %r157, %r158;
	@%p74 bra 	$Lt_5_93954;
	.loc	17	666	0
	ld.param.u64 	%rd73, [__cudaparm_reduce_max_filter_final_maxes];
	cvt.s64.s32 	%rd74, %r8;
	mul.wide.s32 	%rd75, %r8, 4;
	add.u64 	%rd76, %rd73, %rd75;
	.loc	17	663	0
	ld.shared.s32 	%r156, [__cuda___cuda_local_var_17492_30_non_const_idxData24912+0];
	.loc	17	666	0
	st.global.s32 	[%rd76+0], %r156;
$Lt_5_93954:
$Lt_5_93442:
	.loc	17	668	0
	exit;
$LDWend_reduce_max_filter_final:
	} // reduce_max_filter_final

	.entry reduce_max_filter_main (
		.param .u64 __cudaparm_reduce_max_filter_main_g_idata,
		.param .u64 __cudaparm_reduce_max_filter_main_g_odata,
		.param .u64 __cudaparm_reduce_max_filter_main_max_idx,
		.param .u32 __cudaparm_reduce_max_filter_main_width,
		.param .u32 __cudaparm_reduce_max_filter_main_height,
		.param .s32 __cudaparm_reduce_max_filter_main_blockSize,
		.param .u64 __cudaparm_reduce_max_filter_main_maxes,
		.param .s32 __cudaparm_reduce_max_filter_main_nMax)
	{
	.reg .u32 %r<157>;
	.reg .u64 %rd<72>;
	.reg .f64 %fd<69>;
	.reg .pred %p<76>;
	.shared .align 4 .b8 __cuda___cuda_local_var_17700_30_non_const_smaxesRow25988[40];
	.shared .align 4 .b8 __cuda___cuda_local_var_17701_30_non_const_smaxesCol26028[40];
	.shared .align 4 .b8 __cuda___cuda_local_var_17702_30_non_const_smaxesVal26068[40];
	.shared .align 8 .b8 __cuda___cuda_local_var_17703_33_non_const_sdata26112[2048];
	.shared .align 4 .b8 __cuda___cuda_local_var_17704_30_non_const_idxData28160[1024];
	.loc	17	676	0
$LDWbegin_reduce_max_filter_main:
	.loc	17	684	0
	ld.param.s32 	%r1, [__cudaparm_reduce_max_filter_main_blockSize];
	cvt.u32.u16 	%r2, %ctaid.x;
	mul.lo.u32 	%r3, %r1, %r2;
	cvt.u32.u16 	%r4, %tid.x;
	add.u32 	%r5, %r3, %r4;
	mov.s32 	%r6, %r5;
	cvt.u64.u32 	%rd1, %r4;
	mul.wide.u32 	%rd2, %r4, 4;
	ld.param.u32 	%r7, [__cudaparm_reduce_max_filter_main_width];
	ld.param.s32 	%r8, [__cudaparm_reduce_max_filter_main_nMax];
	setp.le.u32 	%p1, %r8, %r4;
	@%p1 bra 	$Lt_6_56578;
	.loc	17	688	0
	mov.u64 	%rd3, __cuda___cuda_local_var_17700_30_non_const_smaxesRow25988;
	mov.u64 	%rd4, __cuda___cuda_local_var_17701_30_non_const_smaxesCol26028;
	ld.param.u64 	%rd5, [__cudaparm_reduce_max_filter_main_maxes];
	add.u64 	%rd6, %rd5, %rd2;
	ld.global.s32 	%r9, [%rd6+0];
	mov.u64 	%rd7, __cuda___cuda_local_var_17702_30_non_const_smaxesVal26068;
	add.u64 	%rd8, %rd2, %rd7;
	st.shared.s32 	[%rd8+0], %r9;
	.loc	17	684	0
	ld.param.u32 	%r7, [__cudaparm_reduce_max_filter_main_width];
	.loc	17	689	0
	div.u32 	%r10, %r9, %r7;
	add.u64 	%rd9, %rd2, %rd3;
	st.shared.s32 	[%rd9+0], %r10;
	.loc	17	690	0
	rem.u32 	%r11, %r9, %r7;
	add.u64 	%rd10, %rd2, %rd4;
	st.shared.s32 	[%rd10+0], %r11;
$Lt_6_56578:
	mov.u64 	%rd3, __cuda___cuda_local_var_17700_30_non_const_smaxesRow25988;
	mov.u64 	%rd4, __cuda___cuda_local_var_17701_30_non_const_smaxesCol26028;
	.loc	17	692	0
	bar.sync 	0;
	ld.param.u32 	%r12, [__cudaparm_reduce_max_filter_main_height];
	mul.lo.u32 	%r13, %r12, %r7;
	setp.le.u32 	%p2, %r13, %r5;
	@%p2 bra 	$Lt_6_94466;
	cvt.u32.u16 	%r14, %nctaid.x;
	mul.lo.u32 	%r15, %r1, %r14;
	cvt.s64.u32 	%rd11, %r15;
	ld.param.u64 	%rd12, [__cudaparm_reduce_max_filter_main_g_idata];
	cvt.u64.u32 	%rd13, %r5;
	mul.wide.u32 	%rd14, %r5, 8;
	add.u64 	%rd15, %rd12, %rd14;
	mul.wide.u32 	%rd16, %r15, 8;
	mov.f64 	%fd1, 0dfff0000000000000;	// -inf
$Lt_6_57602:
 //<loop> Loop body line 692, nesting depth: 1, estimated iterations: unknown
	.loc	17	700	0
	ld.global.f64 	%fd2, [%rd15+0];
	setp.gt.f64 	%p3, %fd2, %fd1;
	@!%p3 bra 	$Lt_6_59906;
	.loc	17	704	0
	mov.u32 	%r16, 0;
	setp.le.s32 	%p4, %r8, %r16;
	@%p4 bra 	$Lt_6_58370;
	mov.s32 	%r17, %r8;
	rem.s32 	%r18, %r6, %r7;
	div.s32 	%r19, %r6, %r7;
	mov.s32 	%r20, 0;
	mov.s32 	%r21, %r17;
$Lt_6_58882:
 //<loop> Loop body line 704, nesting depth: 2, estimated iterations: unknown
	.loc	17	53	0
	cvt.s64.s32 	%rd17, %r20;
	mul.wide.s32 	%rd18, %r20, 4;
	add.u64 	%rd19, %rd18, %rd3;
	ld.shared.s32 	%r22, [%rd19+0];
	add.u64 	%rd20, %rd18, %rd4;
	ld.shared.s32 	%r23, [%rd20+0];
	sub.s32 	%r24, %r23, %r18;
	cvt.rn.f64.s32 	%fd3, %r24;
	sub.s32 	%r25, %r22, %r19;
	cvt.rn.f64.s32 	%fd4, %r25;
	mul.f64 	%fd5, %fd3, %fd3;
	mad.rn.f64 	%fd6, %fd4, %fd4, %fd5;
	mov.f64 	%fd7, 0d4059000000000000;	// 100
	setp.lt.f64 	%p5, %fd6, %fd7;
	@!%p5 bra 	$Lt_6_59138;
	.loc	17	56	0
	mov.s32 	%r26, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_21;
$Lt_6_59138:
	.loc	17	51	0
	add.s32 	%r20, %r20, 1;
	setp.ne.s32 	%p6, %r20, %r8;
	@%p6 bra 	$Lt_6_58882;
$Lt_6_58370:
	.loc	17	61	0
	mov.s32 	%r26, 1;
$LDWendi__Z8distanceiiii_336_21:
	.loc	17	704	0
	mov.u32 	%r27, 0;
	setp.eq.s32 	%p7, %r26, %r27;
	@%p7 bra 	$Lt_6_59906;
	.loc	17	707	0
	mov.f64 	%fd1, %fd2;
	.loc	17	708	0
	mov.s32 	%r28, %r6;
$Lt_6_59906:
$Lt_6_57858:
	add.u32 	%r29, %r1, %r6;
	setp.le.u32 	%p8, %r13, %r29;
	@%p8 bra 	$Lt_6_62978;
	.loc	17	714	0
	cvt.u64.u32 	%rd21, %r29;
	mul.wide.u32 	%rd22, %r29, 8;
	.loc	17	692	0
	ld.param.u64 	%rd12, [__cudaparm_reduce_max_filter_main_g_idata];
	.loc	17	714	0
	add.u64 	%rd23, %rd12, %rd22;
	ld.global.f64 	%fd2, [%rd23+0];
	setp.gt.f64 	%p9, %fd2, %fd1;
	@!%p9 bra 	$Lt_6_62978;
	.loc	17	718	0
	mov.u32 	%r30, 0;
	setp.le.s32 	%p10, %r8, %r30;
	@%p10 bra 	$Lt_6_61442;
	mov.s32 	%r31, %r8;
	rem.s32 	%r32, %r29, %r7;
	div.s32 	%r33, %r29, %r7;
	mov.s32 	%r20, 0;
	mov.s32 	%r34, %r31;
$Lt_6_61954:
 //<loop> Loop body line 718, nesting depth: 2, estimated iterations: unknown
	.loc	17	53	0
	cvt.s64.s32 	%rd24, %r20;
	mul.wide.s32 	%rd18, %r20, 4;
	add.u64 	%rd25, %rd18, %rd3;
	ld.shared.s32 	%r22, [%rd25+0];
	add.u64 	%rd26, %rd18, %rd4;
	ld.shared.s32 	%r23, [%rd26+0];
	sub.s32 	%r35, %r23, %r32;
	cvt.rn.f64.s32 	%fd8, %r35;
	sub.s32 	%r36, %r22, %r33;
	cvt.rn.f64.s32 	%fd9, %r36;
	mul.f64 	%fd10, %fd8, %fd8;
	mad.rn.f64 	%fd11, %fd9, %fd9, %fd10;
	mov.f64 	%fd12, 0d4059000000000000;	// 100
	setp.lt.f64 	%p11, %fd11, %fd12;
	@!%p11 bra 	$Lt_6_62210;
	.loc	17	56	0
	mov.s32 	%r37, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_19;
$Lt_6_62210:
	.loc	17	51	0
	add.s32 	%r20, %r20, 1;
	setp.ne.s32 	%p12, %r20, %r8;
	@%p12 bra 	$Lt_6_61954;
$Lt_6_61442:
	.loc	17	61	0
	mov.s32 	%r37, 1;
$LDWendi__Z8distanceiiii_336_19:
	.loc	17	718	0
	mov.u32 	%r38, 0;
	setp.eq.s32 	%p13, %r37, %r38;
	@%p13 bra 	$Lt_6_62978;
	.loc	17	721	0
	mov.f64 	%fd1, %fd2;
	.loc	17	722	0
	mov.s32 	%r28, %r29;
$Lt_6_62978:
$Lt_6_60930:
$Lt_6_60418:
	.loc	17	727	0
	add.u32 	%r6, %r15, %r6;
	add.u64 	%rd15, %rd15, %rd16;
	setp.gt.u32 	%p14, %r13, %r6;
	@%p14 bra 	$Lt_6_57602;
	bra.uni 	$Lt_6_57090;
$Lt_6_94466:
	mov.f64 	%fd1, 0dfff0000000000000;	// -inf
$Lt_6_57090:
	.loc	17	730	0
	mov.u64 	%rd27, __cuda___cuda_local_var_17703_33_non_const_sdata26112;
	mul.lo.u64 	%rd28, %rd1, 8;
	add.u64 	%rd29, %rd27, %rd28;
	st.shared.f64 	[%rd29+0], %fd1;
	.loc	17	731	0
	mov.u64 	%rd30, __cuda___cuda_local_var_17704_30_non_const_idxData28160;
	add.u64 	%rd31, %rd2, %rd30;
	mov.s32 	%r39, %r28;
	st.shared.s32 	[%rd31+0], %r39;
	.loc	17	733	0
	bar.sync 	0;
	mov.u32 	%r40, 511;
	setp.le.s32 	%p15, %r1, %r40;
	@%p15 bra 	$Lt_6_63746;
	mov.u32 	%r41, 255;
	setp.gt.u32 	%p16, %r4, %r41;
	@%p16 bra 	$Lt_6_66818;
	ld.shared.f64 	%fd13, [%rd29+2048];
	setp.gt.f64 	%p17, %fd13, %fd1;
	@!%p17 bra 	$Lt_6_66818;
	.loc	17	741	0
	ld.shared.s32 	%r42, [%rd31+1024];
	mov.u32 	%r43, 0;
	setp.le.s32 	%p18, %r8, %r43;
	@%p18 bra 	$Lt_6_65282;
	mov.s32 	%r44, %r8;
	ld.shared.s32 	%r42, [%rd31+1024];
	rem.s32 	%r45, %r42, %r7;
	div.s32 	%r46, %r42, %r7;
	mov.s32 	%r20, 0;
	mov.s32 	%r47, %r44;
$Lt_6_65794:
 //<loop> Loop body line 741, nesting depth: 1, estimated iterations: unknown
	.loc	17	53	0
	cvt.s64.s32 	%rd32, %r20;
	mul.wide.s32 	%rd18, %r20, 4;
	add.u64 	%rd33, %rd18, %rd3;
	ld.shared.s32 	%r22, [%rd33+0];
	add.u64 	%rd34, %rd18, %rd4;
	ld.shared.s32 	%r23, [%rd34+0];
	sub.s32 	%r48, %r23, %r45;
	cvt.rn.f64.s32 	%fd14, %r48;
	sub.s32 	%r49, %r22, %r46;
	cvt.rn.f64.s32 	%fd15, %r49;
	mul.f64 	%fd16, %fd14, %fd14;
	mad.rn.f64 	%fd17, %fd15, %fd15, %fd16;
	mov.f64 	%fd18, 0d4059000000000000;	// 100
	setp.lt.f64 	%p19, %fd17, %fd18;
	@!%p19 bra 	$Lt_6_66050;
	.loc	17	56	0
	mov.s32 	%r50, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_17;
$Lt_6_66050:
	.loc	17	51	0
	add.s32 	%r20, %r20, 1;
	setp.ne.s32 	%p20, %r20, %r8;
	@%p20 bra 	$Lt_6_65794;
$Lt_6_65282:
	.loc	17	61	0
	mov.s32 	%r50, 1;
$LDWendi__Z8distanceiiii_336_17:
	.loc	17	741	0
	mov.u32 	%r51, 0;
	setp.eq.s32 	%p21, %r50, %r51;
	@%p21 bra 	$Lt_6_66818;
	.loc	17	733	0
	ld.shared.f64 	%fd13, [%rd29+2048];
	.loc	17	745	0
	mov.f64 	%fd1, %fd13;
	st.shared.f64 	[%rd29+0], %fd13;
	.loc	17	741	0
	ld.shared.s32 	%r42, [%rd31+1024];
	.loc	17	746	0
	st.shared.s32 	[%rd31+0], %r42;
$Lt_6_66818:
$Lt_6_64770:
$Lt_6_64258:
	.loc	17	750	0
	bar.sync 	0;
$Lt_6_63746:
	mov.u32 	%r52, 255;
	setp.le.s32 	%p22, %r1, %r52;
	@%p22 bra 	$Lt_6_67330;
	mov.u32 	%r53, 127;
	setp.gt.u32 	%p23, %r4, %r53;
	@%p23 bra 	$Lt_6_70402;
	ld.shared.f64 	%fd19, [%rd29+1024];
	setp.gt.f64 	%p24, %fd19, %fd1;
	@!%p24 bra 	$Lt_6_70402;
	.loc	17	759	0
	ld.shared.s32 	%r54, [%rd31+512];
	mov.u32 	%r55, 0;
	setp.le.s32 	%p25, %r8, %r55;
	@%p25 bra 	$Lt_6_68866;
	mov.s32 	%r56, %r8;
	ld.shared.s32 	%r54, [%rd31+512];
	rem.s32 	%r57, %r54, %r7;
	div.s32 	%r58, %r54, %r7;
	mov.s32 	%r20, 0;
	mov.s32 	%r59, %r56;
$Lt_6_69378:
 //<loop> Loop body line 759, nesting depth: 1, estimated iterations: unknown
	.loc	17	53	0
	cvt.s64.s32 	%rd35, %r20;
	mul.wide.s32 	%rd18, %r20, 4;
	add.u64 	%rd36, %rd18, %rd3;
	ld.shared.s32 	%r22, [%rd36+0];
	add.u64 	%rd37, %rd18, %rd4;
	ld.shared.s32 	%r23, [%rd37+0];
	sub.s32 	%r60, %r23, %r57;
	cvt.rn.f64.s32 	%fd20, %r60;
	sub.s32 	%r61, %r22, %r58;
	cvt.rn.f64.s32 	%fd21, %r61;
	mul.f64 	%fd22, %fd20, %fd20;
	mad.rn.f64 	%fd23, %fd21, %fd21, %fd22;
	mov.f64 	%fd24, 0d4059000000000000;	// 100
	setp.lt.f64 	%p26, %fd23, %fd24;
	@!%p26 bra 	$Lt_6_69634;
	.loc	17	56	0
	mov.s32 	%r62, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_15;
$Lt_6_69634:
	.loc	17	51	0
	add.s32 	%r20, %r20, 1;
	setp.ne.s32 	%p27, %r20, %r8;
	@%p27 bra 	$Lt_6_69378;
$Lt_6_68866:
	.loc	17	61	0
	mov.s32 	%r62, 1;
$LDWendi__Z8distanceiiii_336_15:
	.loc	17	759	0
	mov.u32 	%r63, 0;
	setp.eq.s32 	%p28, %r62, %r63;
	@%p28 bra 	$Lt_6_70402;
	.loc	17	750	0
	ld.shared.f64 	%fd19, [%rd29+1024];
	.loc	17	763	0
	mov.f64 	%fd1, %fd19;
	st.shared.f64 	[%rd29+0], %fd19;
	.loc	17	759	0
	ld.shared.s32 	%r54, [%rd31+512];
	.loc	17	764	0
	st.shared.s32 	[%rd31+0], %r54;
$Lt_6_70402:
$Lt_6_68354:
$Lt_6_67842:
	.loc	17	768	0
	bar.sync 	0;
$Lt_6_67330:
	mov.u32 	%r64, 127;
	setp.le.s32 	%p29, %r1, %r64;
	@%p29 bra 	$Lt_6_70914;
	mov.u32 	%r65, 63;
	setp.gt.u32 	%p30, %r4, %r65;
	@%p30 bra 	$Lt_6_73986;
	ld.shared.f64 	%fd25, [%rd29+512];
	setp.gt.f64 	%p31, %fd25, %fd1;
	@!%p31 bra 	$Lt_6_73986;
	.loc	17	777	0
	ld.shared.s32 	%r66, [%rd31+256];
	mov.u32 	%r67, 0;
	setp.le.s32 	%p32, %r8, %r67;
	@%p32 bra 	$Lt_6_72450;
	mov.s32 	%r68, %r8;
	ld.shared.s32 	%r66, [%rd31+256];
	rem.s32 	%r69, %r66, %r7;
	div.s32 	%r70, %r66, %r7;
	mov.s32 	%r20, 0;
	mov.s32 	%r71, %r68;
$Lt_6_72962:
 //<loop> Loop body line 777, nesting depth: 1, estimated iterations: unknown
	.loc	17	53	0
	cvt.s64.s32 	%rd38, %r20;
	mul.wide.s32 	%rd18, %r20, 4;
	add.u64 	%rd39, %rd18, %rd3;
	ld.shared.s32 	%r22, [%rd39+0];
	add.u64 	%rd40, %rd18, %rd4;
	ld.shared.s32 	%r23, [%rd40+0];
	sub.s32 	%r72, %r23, %r69;
	cvt.rn.f64.s32 	%fd26, %r72;
	sub.s32 	%r73, %r22, %r70;
	cvt.rn.f64.s32 	%fd27, %r73;
	mul.f64 	%fd28, %fd26, %fd26;
	mad.rn.f64 	%fd29, %fd27, %fd27, %fd28;
	mov.f64 	%fd30, 0d4059000000000000;	// 100
	setp.lt.f64 	%p33, %fd29, %fd30;
	@!%p33 bra 	$Lt_6_73218;
	.loc	17	56	0
	mov.s32 	%r74, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_13;
$Lt_6_73218:
	.loc	17	51	0
	add.s32 	%r20, %r20, 1;
	setp.ne.s32 	%p34, %r20, %r8;
	@%p34 bra 	$Lt_6_72962;
$Lt_6_72450:
	.loc	17	61	0
	mov.s32 	%r74, 1;
$LDWendi__Z8distanceiiii_336_13:
	.loc	17	777	0
	mov.u32 	%r75, 0;
	setp.eq.s32 	%p35, %r74, %r75;
	@%p35 bra 	$Lt_6_73986;
	.loc	17	768	0
	ld.shared.f64 	%fd25, [%rd29+512];
	.loc	17	781	0
	mov.f64 	%fd1, %fd25;
	st.shared.f64 	[%rd29+0], %fd25;
	.loc	17	777	0
	ld.shared.s32 	%r66, [%rd31+256];
	.loc	17	782	0
	st.shared.s32 	[%rd31+0], %r66;
$Lt_6_73986:
$Lt_6_71938:
$Lt_6_71426:
	.loc	17	786	0
	bar.sync 	0;
$Lt_6_70914:
	mov.u32 	%r76, 31;
	setp.gt.u32 	%p36, %r4, %r76;
	@%p36 bra 	$Lt_6_74498;
	mov.u32 	%r77, 63;
	setp.le.s32 	%p37, %r1, %r77;
	@%p37 bra 	$Lt_6_77570;
	ld.volatile.shared.f64 	%fd31, [%rd29+256];
	setp.gt.f64 	%p38, %fd31, %fd1;
	@!%p38 bra 	$Lt_6_77570;
	.loc	17	800	0
	ld.volatile.shared.s32 	%r78, [%rd31+128];
	mov.u32 	%r79, 0;
	setp.le.s32 	%p39, %r8, %r79;
	@%p39 bra 	$Lt_6_76034;
	mov.s32 	%r80, %r8;
	ld.volatile.shared.s32 	%r78, [%rd31+128];
	rem.s32 	%r81, %r78, %r7;
	div.s32 	%r82, %r78, %r7;
	mov.s32 	%r83, 0;
	mov.s32 	%r84, %r80;
$Lt_6_76546:
 //<loop> Loop body line 800, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd41, %r83;
	mul.wide.s32 	%rd42, %r83, 4;
	add.u64 	%rd43, %rd42, %rd3;
	ld.volatile.shared.s32 	%r85, [%rd43+0];
	add.u64 	%rd44, %rd42, %rd4;
	ld.volatile.shared.s32 	%r86, [%rd44+0];
	sub.s32 	%r87, %r86, %r81;
	cvt.rn.f64.s32 	%fd32, %r87;
	sub.s32 	%r88, %r85, %r82;
	cvt.rn.f64.s32 	%fd33, %r88;
	mul.f64 	%fd34, %fd32, %fd32;
	mad.rn.f64 	%fd35, %fd33, %fd33, %fd34;
	mov.f64 	%fd36, 0d4059000000000000;	// 100
	setp.lt.f64 	%p40, %fd35, %fd36;
	@!%p40 bra 	$Lt_6_76802;
	.loc	17	77	0
	mov.s32 	%r89, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_11;
$Lt_6_76802:
	.loc	17	72	0
	add.s32 	%r83, %r83, 1;
	setp.ne.s32 	%p41, %r83, %r8;
	@%p41 bra 	$Lt_6_76546;
$Lt_6_76034:
	.loc	17	82	0
	mov.s32 	%r89, 1;
$LDWendi__Z8distanceiiii_336_11:
	.loc	17	800	0
	mov.u32 	%r90, 0;
	setp.eq.s32 	%p42, %r89, %r90;
	@%p42 bra 	$Lt_6_77570;
	.loc	17	804	0
	ld.volatile.shared.f64 	%fd1, [%rd29+256];
	st.volatile.shared.f64 	[%rd29+0], %fd1;
	.loc	17	805	0
	ld.volatile.shared.s32 	%r91, [%rd31+128];
	st.volatile.shared.s32 	[%rd31+0], %r91;
$Lt_6_77570:
$Lt_6_75522:
$Lt_6_75010:
	mov.u32 	%r92, 31;
	setp.le.s32 	%p43, %r1, %r92;
	@%p43 bra 	$Lt_6_80642;
	ld.volatile.shared.f64 	%fd37, [%rd29+128];
	setp.gt.f64 	%p44, %fd37, %fd1;
	@!%p44 bra 	$Lt_6_80642;
	.loc	17	813	0
	ld.volatile.shared.s32 	%r93, [%rd31+64];
	mov.u32 	%r94, 0;
	setp.le.s32 	%p45, %r8, %r94;
	@%p45 bra 	$Lt_6_79106;
	mov.s32 	%r95, %r8;
	ld.volatile.shared.s32 	%r93, [%rd31+64];
	rem.s32 	%r96, %r93, %r7;
	div.s32 	%r97, %r93, %r7;
	mov.s32 	%r83, 0;
	mov.s32 	%r98, %r95;
$Lt_6_79618:
 //<loop> Loop body line 813, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd45, %r83;
	mul.wide.s32 	%rd42, %r83, 4;
	add.u64 	%rd46, %rd42, %rd3;
	ld.volatile.shared.s32 	%r85, [%rd46+0];
	add.u64 	%rd47, %rd42, %rd4;
	ld.volatile.shared.s32 	%r86, [%rd47+0];
	sub.s32 	%r99, %r86, %r96;
	cvt.rn.f64.s32 	%fd38, %r99;
	sub.s32 	%r100, %r85, %r97;
	cvt.rn.f64.s32 	%fd39, %r100;
	mul.f64 	%fd40, %fd38, %fd38;
	mad.rn.f64 	%fd41, %fd39, %fd39, %fd40;
	mov.f64 	%fd42, 0d4059000000000000;	// 100
	setp.lt.f64 	%p46, %fd41, %fd42;
	@!%p46 bra 	$Lt_6_79874;
	.loc	17	77	0
	mov.s32 	%r101, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_9;
$Lt_6_79874:
	.loc	17	72	0
	add.s32 	%r83, %r83, 1;
	setp.ne.s32 	%p47, %r83, %r8;
	@%p47 bra 	$Lt_6_79618;
$Lt_6_79106:
	.loc	17	82	0
	mov.s32 	%r101, 1;
$LDWendi__Z8distanceiiii_336_9:
	.loc	17	813	0
	mov.u32 	%r102, 0;
	setp.eq.s32 	%p48, %r101, %r102;
	@%p48 bra 	$Lt_6_80642;
	.loc	17	817	0
	ld.volatile.shared.f64 	%fd1, [%rd29+128];
	st.volatile.shared.f64 	[%rd29+0], %fd1;
	.loc	17	818	0
	ld.volatile.shared.s32 	%r103, [%rd31+64];
	st.volatile.shared.s32 	[%rd31+0], %r103;
$Lt_6_80642:
$Lt_6_78594:
$Lt_6_78082:
	mov.u32 	%r104, 15;
	setp.le.s32 	%p49, %r1, %r104;
	@%p49 bra 	$Lt_6_83714;
	ld.volatile.shared.f64 	%fd43, [%rd29+64];
	setp.gt.f64 	%p50, %fd43, %fd1;
	@!%p50 bra 	$Lt_6_83714;
	.loc	17	825	0
	ld.volatile.shared.s32 	%r105, [%rd31+32];
	mov.u32 	%r106, 0;
	setp.le.s32 	%p51, %r8, %r106;
	@%p51 bra 	$Lt_6_82178;
	mov.s32 	%r107, %r8;
	ld.volatile.shared.s32 	%r105, [%rd31+32];
	rem.s32 	%r108, %r105, %r7;
	div.s32 	%r109, %r105, %r7;
	mov.s32 	%r83, 0;
	mov.s32 	%r110, %r107;
$Lt_6_82690:
 //<loop> Loop body line 825, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd48, %r83;
	mul.wide.s32 	%rd42, %r83, 4;
	add.u64 	%rd49, %rd42, %rd3;
	ld.volatile.shared.s32 	%r85, [%rd49+0];
	add.u64 	%rd50, %rd42, %rd4;
	ld.volatile.shared.s32 	%r86, [%rd50+0];
	sub.s32 	%r111, %r86, %r108;
	cvt.rn.f64.s32 	%fd44, %r111;
	sub.s32 	%r112, %r85, %r109;
	cvt.rn.f64.s32 	%fd45, %r112;
	mul.f64 	%fd46, %fd44, %fd44;
	mad.rn.f64 	%fd47, %fd45, %fd45, %fd46;
	mov.f64 	%fd48, 0d4059000000000000;	// 100
	setp.lt.f64 	%p52, %fd47, %fd48;
	@!%p52 bra 	$Lt_6_82946;
	.loc	17	77	0
	mov.s32 	%r113, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_7;
$Lt_6_82946:
	.loc	17	72	0
	add.s32 	%r83, %r83, 1;
	setp.ne.s32 	%p53, %r83, %r8;
	@%p53 bra 	$Lt_6_82690;
$Lt_6_82178:
	.loc	17	82	0
	mov.s32 	%r113, 1;
$LDWendi__Z8distanceiiii_336_7:
	.loc	17	825	0
	mov.u32 	%r114, 0;
	setp.eq.s32 	%p54, %r113, %r114;
	@%p54 bra 	$Lt_6_83714;
	.loc	17	829	0
	ld.volatile.shared.f64 	%fd1, [%rd29+64];
	st.volatile.shared.f64 	[%rd29+0], %fd1;
	.loc	17	830	0
	ld.volatile.shared.s32 	%r115, [%rd31+32];
	st.volatile.shared.s32 	[%rd31+0], %r115;
$Lt_6_83714:
$Lt_6_81666:
$Lt_6_81154:
	mov.u32 	%r116, 7;
	setp.le.s32 	%p55, %r1, %r116;
	@%p55 bra 	$Lt_6_86786;
	ld.volatile.shared.f64 	%fd49, [%rd29+32];
	setp.gt.f64 	%p56, %fd49, %fd1;
	@!%p56 bra 	$Lt_6_86786;
	.loc	17	837	0
	ld.volatile.shared.s32 	%r117, [%rd31+16];
	mov.u32 	%r118, 0;
	setp.le.s32 	%p57, %r8, %r118;
	@%p57 bra 	$Lt_6_85250;
	mov.s32 	%r119, %r8;
	ld.volatile.shared.s32 	%r117, [%rd31+16];
	rem.s32 	%r120, %r117, %r7;
	div.s32 	%r121, %r117, %r7;
	mov.s32 	%r83, 0;
	mov.s32 	%r122, %r119;
$Lt_6_85762:
 //<loop> Loop body line 837, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd51, %r83;
	mul.wide.s32 	%rd42, %r83, 4;
	add.u64 	%rd52, %rd42, %rd3;
	ld.volatile.shared.s32 	%r85, [%rd52+0];
	add.u64 	%rd53, %rd42, %rd4;
	ld.volatile.shared.s32 	%r86, [%rd53+0];
	sub.s32 	%r123, %r86, %r120;
	cvt.rn.f64.s32 	%fd50, %r123;
	sub.s32 	%r124, %r85, %r121;
	cvt.rn.f64.s32 	%fd51, %r124;
	mul.f64 	%fd52, %fd50, %fd50;
	mad.rn.f64 	%fd53, %fd51, %fd51, %fd52;
	mov.f64 	%fd54, 0d4059000000000000;	// 100
	setp.lt.f64 	%p58, %fd53, %fd54;
	@!%p58 bra 	$Lt_6_86018;
	.loc	17	77	0
	mov.s32 	%r125, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_5;
$Lt_6_86018:
	.loc	17	72	0
	add.s32 	%r83, %r83, 1;
	setp.ne.s32 	%p59, %r83, %r8;
	@%p59 bra 	$Lt_6_85762;
$Lt_6_85250:
	.loc	17	82	0
	mov.s32 	%r125, 1;
$LDWendi__Z8distanceiiii_336_5:
	.loc	17	837	0
	mov.u32 	%r126, 0;
	setp.eq.s32 	%p60, %r125, %r126;
	@%p60 bra 	$Lt_6_86786;
	.loc	17	841	0
	ld.volatile.shared.f64 	%fd1, [%rd29+32];
	st.volatile.shared.f64 	[%rd29+0], %fd1;
	.loc	17	842	0
	ld.volatile.shared.s32 	%r127, [%rd31+16];
	st.volatile.shared.s32 	[%rd31+0], %r127;
$Lt_6_86786:
$Lt_6_84738:
$Lt_6_84226:
	mov.u32 	%r128, 3;
	setp.le.s32 	%p61, %r1, %r128;
	@%p61 bra 	$Lt_6_89858;
	ld.volatile.shared.f64 	%fd55, [%rd29+16];
	setp.gt.f64 	%p62, %fd55, %fd1;
	@!%p62 bra 	$Lt_6_89858;
	.loc	17	849	0
	ld.volatile.shared.s32 	%r129, [%rd31+8];
	mov.u32 	%r130, 0;
	setp.le.s32 	%p63, %r8, %r130;
	@%p63 bra 	$Lt_6_88322;
	mov.s32 	%r131, %r8;
	ld.volatile.shared.s32 	%r129, [%rd31+8];
	rem.s32 	%r132, %r129, %r7;
	div.s32 	%r133, %r129, %r7;
	mov.s32 	%r83, 0;
	mov.s32 	%r134, %r131;
$Lt_6_88834:
 //<loop> Loop body line 849, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd54, %r83;
	mul.wide.s32 	%rd42, %r83, 4;
	add.u64 	%rd55, %rd42, %rd3;
	ld.volatile.shared.s32 	%r85, [%rd55+0];
	add.u64 	%rd56, %rd42, %rd4;
	ld.volatile.shared.s32 	%r86, [%rd56+0];
	sub.s32 	%r135, %r86, %r132;
	cvt.rn.f64.s32 	%fd56, %r135;
	sub.s32 	%r136, %r85, %r133;
	cvt.rn.f64.s32 	%fd57, %r136;
	mul.f64 	%fd58, %fd56, %fd56;
	mad.rn.f64 	%fd59, %fd57, %fd57, %fd58;
	mov.f64 	%fd60, 0d4059000000000000;	// 100
	setp.lt.f64 	%p64, %fd59, %fd60;
	@!%p64 bra 	$Lt_6_89090;
	.loc	17	77	0
	mov.s32 	%r137, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_3;
$Lt_6_89090:
	.loc	17	72	0
	add.s32 	%r83, %r83, 1;
	setp.ne.s32 	%p65, %r83, %r8;
	@%p65 bra 	$Lt_6_88834;
$Lt_6_88322:
	.loc	17	82	0
	mov.s32 	%r137, 1;
$LDWendi__Z8distanceiiii_336_3:
	.loc	17	849	0
	mov.u32 	%r138, 0;
	setp.eq.s32 	%p66, %r137, %r138;
	@%p66 bra 	$Lt_6_89858;
	.loc	17	853	0
	ld.volatile.shared.f64 	%fd1, [%rd29+16];
	st.volatile.shared.f64 	[%rd29+0], %fd1;
	.loc	17	854	0
	ld.volatile.shared.s32 	%r139, [%rd31+8];
	st.volatile.shared.s32 	[%rd31+0], %r139;
$Lt_6_89858:
$Lt_6_87810:
$Lt_6_87298:
	mov.u32 	%r140, 1;
	setp.le.s32 	%p67, %r1, %r140;
	@%p67 bra 	$Lt_6_92930;
	ld.volatile.shared.f64 	%fd61, [%rd29+8];
	setp.gt.f64 	%p68, %fd61, %fd1;
	@!%p68 bra 	$Lt_6_92930;
	.loc	17	861	0
	ld.volatile.shared.s32 	%r141, [%rd31+4];
	mov.u32 	%r142, 0;
	setp.le.s32 	%p69, %r8, %r142;
	@%p69 bra 	$Lt_6_91394;
	mov.s32 	%r143, %r8;
	ld.volatile.shared.s32 	%r141, [%rd31+4];
	rem.s32 	%r144, %r141, %r7;
	div.s32 	%r145, %r141, %r7;
	mov.s32 	%r83, 0;
	mov.s32 	%r146, %r143;
$Lt_6_91906:
 //<loop> Loop body line 861, nesting depth: 1, estimated iterations: unknown
	.loc	17	74	0
	cvt.s64.s32 	%rd57, %r83;
	mul.wide.s32 	%rd42, %r83, 4;
	add.u64 	%rd58, %rd42, %rd3;
	ld.volatile.shared.s32 	%r85, [%rd58+0];
	add.u64 	%rd59, %rd42, %rd4;
	ld.volatile.shared.s32 	%r86, [%rd59+0];
	sub.s32 	%r147, %r86, %r144;
	cvt.rn.f64.s32 	%fd62, %r147;
	sub.s32 	%r148, %r85, %r145;
	cvt.rn.f64.s32 	%fd63, %r148;
	mul.f64 	%fd64, %fd62, %fd62;
	mad.rn.f64 	%fd65, %fd63, %fd63, %fd64;
	mov.f64 	%fd66, 0d4059000000000000;	// 100
	setp.lt.f64 	%p70, %fd65, %fd66;
	@!%p70 bra 	$Lt_6_92162;
	.loc	17	77	0
	mov.s32 	%r149, 0;
	bra.uni 	$LDWendi__Z8distanceiiii_336_1;
$Lt_6_92162:
	.loc	17	72	0
	add.s32 	%r83, %r83, 1;
	setp.ne.s32 	%p71, %r83, %r8;
	@%p71 bra 	$Lt_6_91906;
$Lt_6_91394:
	.loc	17	82	0
	mov.s32 	%r149, 1;
$LDWendi__Z8distanceiiii_336_1:
	.loc	17	861	0
	mov.u32 	%r150, 0;
	setp.eq.s32 	%p72, %r149, %r150;
	@%p72 bra 	$Lt_6_92930;
	.loc	17	865	0
	ld.volatile.shared.f64 	%fd1, [%rd29+8];
	st.volatile.shared.f64 	[%rd29+0], %fd1;
	.loc	17	866	0
	ld.volatile.shared.s32 	%r151, [%rd31+4];
	st.volatile.shared.s32 	[%rd31+0], %r151;
$Lt_6_92930:
$Lt_6_90882:
$Lt_6_90370:
	.loc	17	869	0
	bar.sync 	0;
$Lt_6_74498:
	mov.u32 	%r152, 0;
	setp.ne.u32 	%p73, %r4, %r152;
	@%p73 bra 	$Lt_6_93954;
	.loc	17	874	0
	cvt.u64.u32 	%rd60, %r2;
	ld.shared.f64 	%fd67, [__cuda___cuda_local_var_17703_33_non_const_sdata26112+0];
	ld.param.u64 	%rd61, [__cudaparm_reduce_max_filter_main_g_odata];
	mul.wide.u32 	%rd62, %r2, 8;
	add.u64 	%rd63, %rd61, %rd62;
	st.global.f64 	[%rd63+0], %fd67;
	.loc	17	875	0
	ld.shared.s32 	%r153, [__cuda___cuda_local_var_17704_30_non_const_idxData28160+0];
	ld.param.u64 	%rd64, [__cudaparm_reduce_max_filter_main_max_idx];
	mul.wide.u32 	%rd65, %r2, 4;
	add.u64 	%rd66, %rd64, %rd65;
	st.global.s32 	[%rd66+0], %r153;
	cvt.u32.u16 	%r154, %nctaid.x;
	mov.u32 	%r155, 1;
	setp.ne.u32 	%p74, %r154, %r155;
	@%p74 bra 	$Lt_6_93954;
	.loc	17	878	0
	ld.param.u64 	%rd67, [__cudaparm_reduce_max_filter_main_maxes];
	cvt.s64.s32 	%rd68, %r8;
	mul.wide.s32 	%rd69, %r8, 4;
	add.u64 	%rd70, %rd67, %rd69;
	.loc	17	875	0
	ld.shared.s32 	%r153, [__cuda___cuda_local_var_17704_30_non_const_idxData28160+0];
	.loc	17	878	0
	st.global.s32 	[%rd70+0], %r153;
$Lt_6_93954:
$Lt_6_93442:
	.loc	17	880	0
	exit;
$LDWend_reduce_max_filter_main:
	} // reduce_max_filter_main