.version 2.3
	.target sm_20
	.address_size 64
	// compiled with /usr/local/cuda/open64/lib//be
	// nvopencc 4.0 built on 2011-05-12

	//-----------------------------------------------------------
	// Compiling /tmp/tmpxft_0000bf4b_00000000-9_coul_long_gpu_kernel.cpp3.i (/home/sjplimp/ccBI#.qJTqsI)
	//-----------------------------------------------------------

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

	.file	1	"<command-line>"
	.file	2	"/tmp/tmpxft_0000bf4b_00000000-8_coul_long_gpu_kernel.cudafe2.gpu"
	.file	3	"/usr/lib/gcc/x86_64-redhat-linux/4.4.5/include/stddef.h"
	.file	4	"/usr/local/cuda/include/crt/device_runtime.h"
	.file	5	"/usr/local/cuda/include/host_defines.h"
	.file	6	"/usr/local/cuda/include/builtin_types.h"
	.file	7	"/usr/local/cuda/include/device_types.h"
	.file	8	"/usr/local/cuda/include/driver_types.h"
	.file	9	"/usr/local/cuda/include/surface_types.h"
	.file	10	"/usr/local/cuda/include/texture_types.h"
	.file	11	"/usr/local/cuda/include/vector_types.h"
	.file	12	"/usr/local/cuda/include/device_launch_parameters.h"
	.file	13	"/usr/local/cuda/include/crt/storage_class.h"
	.file	14	"/usr/include/bits/types.h"
	.file	15	"/usr/include/time.h"
	.file	16	"coul_long_gpu_kernel.cu"
	.file	17	"/usr/local/cuda/include/common_functions.h"
	.file	18	"/usr/local/cuda/include/math_functions.h"
	.file	19	"/usr/local/cuda/include/math_constants.h"
	.file	20	"/usr/local/cuda/include/device_functions.h"
	.file	21	"/usr/local/cuda/include/sm_11_atomic_functions.h"
	.file	22	"/usr/local/cuda/include/sm_12_atomic_functions.h"
	.file	23	"/usr/local/cuda/include/sm_13_double_functions.h"
	.file	24	"/usr/local/cuda/include/sm_20_atomic_functions.h"
	.file	25	"/usr/local/cuda/include/sm_20_intrinsics.h"
	.file	26	"/usr/local/cuda/include/surface_functions.h"
	.file	27	"/usr/local/cuda/include/texture_fetch_functions.h"
	.file	28	"/usr/local/cuda/include/math_functions_dbl_ptx3.h"

	.global .texref pos_tex;
	.global .texref q_tex;

	.entry kernel_pair (
		.param .u64 __cudaparm_kernel_pair_x_,
		.param .u64 __cudaparm_kernel_pair_lj1,
		.param .u64 __cudaparm_kernel_pair_lj3,
		.param .s32 __cudaparm_kernel_pair_lj_types,
		.param .u64 __cudaparm_kernel_pair_sp_cl_in,
		.param .u64 __cudaparm_kernel_pair_dev_nbor,
		.param .u64 __cudaparm_kernel_pair_dev_packed,
		.param .u64 __cudaparm_kernel_pair_ans,
		.param .u64 __cudaparm_kernel_pair_engv,
		.param .s32 __cudaparm_kernel_pair_eflag,
		.param .s32 __cudaparm_kernel_pair_vflag,
		.param .s32 __cudaparm_kernel_pair_inum,
		.param .s32 __cudaparm_kernel_pair_nbor_pitch,
		.param .u64 __cudaparm_kernel_pair_q_,
		.param .f32 __cudaparm_kernel_pair_cut_coulsq,
		.param .f32 __cudaparm_kernel_pair_qqrd2e,
		.param .f32 __cudaparm_kernel_pair_g_ewald,
		.param .s32 __cudaparm_kernel_pair_t_per_atom)
	{
	.reg .u32 %r<81>;
	.reg .u64 %rd<57>;
	.reg .f32 %f<132>;
	.reg .pred %p<19>;
	.shared .align 16 .b8 __cuda___cuda_local_var_32498_33_non_const_sp_cl112[16];
	.shared .align 4 .b8 __cuda___cuda_local_var_32585_35_non_const_red_acc128[3072];
	// __cuda_local_var_32505_10_non_const_f = 48
	// __cuda_local_var_32509_9_non_const_virial = 16
	.loc	16	108	0
$LDWbegin_kernel_pair:
	.loc	16	115	0
	ld.param.u64 	%rd1, [__cudaparm_kernel_pair_sp_cl_in];
	ldu.global.f32 	%f1, [%rd1+0];
	.loc	16	116	0
	ld.global.f32 	%f2, [%rd1+4];
	.loc	16	117	0
	ld.global.f32 	%f3, [%rd1+8];
	.loc	16	118	0
	ld.global.f32 	%f4, [%rd1+12];
	st.shared.v4.f32 	[__cuda___cuda_local_var_32498_33_non_const_sp_cl112+0], {%f1,%f2,%f3,%f4};
	.loc	16	127	0
	mov.f32 	%f5, 0f00000000;     	// 0
	mov.f32 	%f6, %f5;
	mov.f32 	%f7, 0f00000000;     	// 0
	mov.f32 	%f8, %f7;
	mov.f32 	%f9, 0f00000000;     	// 0
	mov.f32 	%f10, %f9;
	mov.f32 	%f11, 0f00000000;    	// 0
	mov.f32 	%f12, %f11;
	mov.f32 	%f13, 0f00000000;    	// 0
	mov.f32 	%f14, %f13;
	mov.f32 	%f15, 0f00000000;    	// 0
	mov.f32 	%f16, %f15;
	ld.param.s32 	%r1, [__cudaparm_kernel_pair_t_per_atom];
	cvt.s32.u32 	%r2, %tid.x;
	div.s32 	%r3, %r2, %r1;
	cvt.s32.u32 	%r4, %ntid.x;
	div.s32 	%r5, %r4, %r1;
	rem.s32 	%r6, %r2, %r1;
	cvt.s32.u32 	%r7, %ctaid.x;
	mul.lo.s32 	%r8, %r7, %r5;
	add.s32 	%r9, %r3, %r8;
	ld.param.s32 	%r10, [__cudaparm_kernel_pair_inum];
	setp.lt.s32 	%p1, %r9, %r10;
	@!%p1 bra 	$Lt_0_19202;
	.loc	16	131	0
	cvt.s64.s32 	%rd2, %r9;
	mul.wide.s32 	%rd3, %r9, 4;
	ld.param.u64 	%rd4, [__cudaparm_kernel_pair_dev_nbor];
	add.u64 	%rd5, %rd3, %rd4;
	ld.global.s32 	%r11, [%rd5+0];
	.loc	16	133	0
	ld.param.s32 	%r12, [__cudaparm_kernel_pair_nbor_pitch];
	cvt.s64.s32 	%rd6, %r12;
	mul.wide.s32 	%rd7, %r12, 4;
	add.u64 	%rd8, %rd7, %rd5;
	ld.global.s32 	%r13, [%rd8+0];
	add.u64 	%rd9, %rd7, %rd8;
	ld.param.u64 	%rd10, [__cudaparm_kernel_pair_dev_packed];
	setp.ne.u64 	%p2, %rd10, %rd4;
	@%p2 bra 	$Lt_0_19714;
	.loc	16	139	0
	cvt.s32.s64 	%r14, %rd6;
	mul.lo.s32 	%r15, %r14, %r13;
	cvt.s64.s32 	%rd11, %r15;
	mul.wide.s32 	%rd12, %r15, 4;
	add.u64 	%rd13, %rd9, %rd12;
	.loc	16	140	0
	mul.lo.s32 	%r16, %r6, %r14;
	cvt.s64.s32 	%rd14, %r16;
	mul.wide.s32 	%rd15, %r16, 4;
	add.u64 	%rd16, %rd9, %rd15;
	.loc	16	141	0
	mul.lo.s32 	%r17, %r14, %r1;
	bra.uni 	$Lt_0_19458;
$Lt_0_19714:
	.loc	16	143	0
	ld.global.s32 	%r18, [%rd9+0];
	cvt.s64.s32 	%rd17, %r18;
	mul.wide.s32 	%rd18, %r18, 4;
	add.u64 	%rd19, %rd10, %rd18;
	.loc	16	144	0
	cvt.s64.s32 	%rd20, %r13;
	mul.wide.s32 	%rd21, %r13, 4;
	add.u64 	%rd13, %rd19, %rd21;
	.loc	16	145	0
	mov.s32 	%r17, %r1;
	.loc	16	146	0
	cvt.s64.s32 	%rd22, %r6;
	mul.wide.s32 	%rd23, %r6, 4;
	add.u64 	%rd16, %rd19, %rd23;
$Lt_0_19458:
	.loc	16	149	0
	mov.u32 	%r19, %r11;
	mov.s32 	%r20, 0;
	mov.u32 	%r21, %r20;
	mov.s32 	%r22, 0;
	mov.u32 	%r23, %r22;
	mov.s32 	%r24, 0;
	mov.u32 	%r25, %r24;
	tex.1d.v4.f32.s32 {%f17,%f18,%f19,%f20},[pos_tex,{%r19,%r21,%r23,%r25}];
	mov.f32 	%f21, %f17;
	mov.f32 	%f22, %f18;
	mov.f32 	%f23, %f19;
	.loc	16	150	0
	mov.u32 	%r26, %r11;
	mov.s32 	%r27, 0;
	mov.u32 	%r28, %r27;
	mov.s32 	%r29, 0;
	mov.u32 	%r30, %r29;
	mov.s32 	%r31, 0;
	mov.u32 	%r32, %r31;
	tex.1d.v4.f32.s32 {%f24,%f25,%f26,%f27},[q_tex,{%r26,%r28,%r30,%r32}];
	mov.f32 	%f28, %f24;
	setp.ge.u64 	%p3, %rd16, %rd13;
	@%p3 bra 	$Lt_0_27650;
	cvt.s64.s32 	%rd24, %r17;
	ld.param.f32 	%f29, [__cudaparm_kernel_pair_cut_coulsq];
	mov.f32 	%f30, 0f00000000;    	// 0
	mov.f32 	%f31, 0f00000000;    	// 0
	mov.f32 	%f32, 0f00000000;    	// 0
	mov.f32 	%f33, 0f00000000;    	// 0
	mov.u64 	%rd25, __cuda___cuda_local_var_32498_33_non_const_sp_cl112;
$Lt_0_20482:
 //<loop> Loop body line 150, nesting depth: 1, estimated iterations: unknown
	.loc	16	153	0
	ld.global.s32 	%r33, [%rd16+0];
	.loc	16	156	0
	mov.f32 	%f34, 0f3f800000;    	// 1
	shr.s32 	%r34, %r33, 30;
	and.b32 	%r35, %r34, 3;
	cvt.s64.s32 	%rd26, %r35;
	mul.wide.s32 	%rd27, %r35, 4;
	add.u64 	%rd28, %rd25, %rd27;
	ld.shared.f32 	%f35, [%rd28+0];
	sub.ftz.f32 	%f36, %f34, %f35;
	.loc	16	159	0
	and.b32 	%r36, %r33, 1073741823;
	mov.u32 	%r37, %r36;
	mov.s32 	%r38, 0;
	mov.u32 	%r39, %r38;
	mov.s32 	%r40, 0;
	mov.u32 	%r41, %r40;
	mov.s32 	%r42, 0;
	mov.u32 	%r43, %r42;
	tex.1d.v4.f32.s32 {%f37,%f38,%f39,%f40},[pos_tex,{%r37,%r39,%r41,%r43}];
	mov.f32 	%f41, %f37;
	mov.f32 	%f42, %f38;
	mov.f32 	%f43, %f39;
	sub.ftz.f32 	%f44, %f22, %f42;
	sub.ftz.f32 	%f45, %f21, %f41;
	sub.ftz.f32 	%f46, %f23, %f43;
	mul.ftz.f32 	%f47, %f44, %f44;
	fma.rn.ftz.f32 	%f48, %f45, %f45, %f47;
	fma.rn.ftz.f32 	%f49, %f46, %f46, %f48;
	setp.lt.ftz.f32 	%p4, %f49, %f29;
	@!%p4 bra 	$Lt_0_21250;
	.loc	16	175	0
	sqrt.approx.ftz.f32 	%f50, %f49;
	ld.param.f32 	%f51, [__cudaparm_kernel_pair_g_ewald];
	mul.ftz.f32 	%f52, %f51, %f50;
	mul.ftz.f32 	%f53, %f52, %f52;
	mov.f32 	%f54, 0f3f800000;    	// 1
	mov.f32 	%f55, 0f3ea7ba05;    	// 0.327591
	fma.rn.ftz.f32 	%f56, %f55, %f52, %f54;
	neg.ftz.f32 	%f57, %f53;
	rcp.approx.ftz.f32 	%f58, %f56;
	mov.f32 	%f59, 0f3fb8aa3b;    	// 1.4427
	mul.ftz.f32 	%f60, %f57, %f59;
	ex2.approx.ftz.f32 	%f61, %f60;
	mov.f32 	%f62, 0f3e827906;    	// 0.25483
	mov.f32 	%f63, 0fbe91a98e;    	// -0.284497
	mov.f32 	%f64, 0f3fb5f0e3;    	// 1.42141
	mov.f32 	%f65, 0fbfba00e3;    	// -1.45315
	mov.f32 	%f66, 0f3f87dc22;    	// 1.06141
	fma.rn.ftz.f32 	%f67, %f66, %f58, %f65;
	fma.rn.ftz.f32 	%f68, %f58, %f67, %f64;
	fma.rn.ftz.f32 	%f69, %f58, %f68, %f63;
	fma.rn.ftz.f32 	%f70, %f58, %f69, %f62;
	mul.ftz.f32 	%f71, %f58, %f70;
	mul.ftz.f32 	%f72, %f61, %f71;
	.loc	16	176	0
	mov.u32 	%r44, %r36;
	mov.s32 	%r45, 0;
	mov.u32 	%r46, %r45;
	mov.s32 	%r47, 0;
	mov.u32 	%r48, %r47;
	mov.s32 	%r49, 0;
	mov.u32 	%r50, %r49;
	tex.1d.v4.f32.s32 {%f73,%f74,%f75,%f76},[q_tex,{%r44,%r46,%r48,%r50}];
	mov.f32 	%f77, %f73;
	.loc	16	177	0
	ld.param.f32 	%f78, [__cudaparm_kernel_pair_qqrd2e];
	mul.ftz.f32 	%f79, %f78, %f28;
	mul.ftz.f32 	%f80, %f79, %f77;
	div.approx.ftz.f32 	%f81, %f80, %f50;
	mov.f32 	%f82, 0f3f906ebb;    	// 1.12838
	mul.ftz.f32 	%f83, %f52, %f82;
	fma.rn.ftz.f32 	%f84, %f61, %f83, %f72;
	sub.ftz.f32 	%f85, %f84, %f36;
	mul.ftz.f32 	%f86, %f81, %f85;
	rcp.approx.ftz.f32 	%f87, %f49;
	mul.ftz.f32 	%f88, %f86, %f87;
	.loc	16	179	0
	fma.rn.ftz.f32 	%f32, %f45, %f88, %f32;
	.loc	16	180	0
	fma.rn.ftz.f32 	%f31, %f44, %f88, %f31;
	.loc	16	181	0
	fma.rn.ftz.f32 	%f30, %f46, %f88, %f30;
	.loc	16	168	0
	sub.ftz.f32 	%f89, %f72, %f36;
	fma.rn.ftz.f32 	%f90, %f81, %f89, %f33;
	ld.param.s32 	%r51, [__cudaparm_kernel_pair_eflag];
	mov.s32 	%r52, 0;
	setp.gt.s32 	%p5, %r51, %r52;
	selp.f32 	%f33, %f90, %f33, %p5;
	ld.param.s32 	%r53, [__cudaparm_kernel_pair_vflag];
	mov.u32 	%r54, 0;
	setp.le.s32 	%p6, %r53, %r54;
	@%p6 bra 	$Lt_0_21250;
	.loc	16	187	0
	mov.f32 	%f91, %f6;
	mul.ftz.f32 	%f92, %f45, %f45;
	fma.rn.ftz.f32 	%f93, %f88, %f92, %f91;
	mov.f32 	%f6, %f93;
	.loc	16	188	0
	mov.f32 	%f94, %f8;
	fma.rn.ftz.f32 	%f95, %f88, %f47, %f94;
	mov.f32 	%f8, %f95;
	.loc	16	189	0
	mov.f32 	%f96, %f10;
	mul.ftz.f32 	%f97, %f46, %f46;
	fma.rn.ftz.f32 	%f98, %f88, %f97, %f96;
	mov.f32 	%f10, %f98;
	.loc	16	190	0
	mov.f32 	%f99, %f12;
	mul.ftz.f32 	%f100, %f44, %f45;
	fma.rn.ftz.f32 	%f101, %f88, %f100, %f99;
	mov.f32 	%f12, %f101;
	.loc	16	191	0
	mov.f32 	%f102, %f14;
	mul.ftz.f32 	%f103, %f45, %f46;
	fma.rn.ftz.f32 	%f104, %f88, %f103, %f102;
	mov.f32 	%f14, %f104;
	.loc	16	192	0
	mul.ftz.f32 	%f105, %f44, %f46;
	fma.rn.ftz.f32 	%f15, %f88, %f105, %f15;
	mov.f32 	%f16, %f15;
$Lt_0_21250:
$Lt_0_20738:
	.loc	16	152	0
	mul.lo.u64 	%rd29, %rd24, 4;
	add.u64 	%rd16, %rd16, %rd29;
	setp.lt.u64 	%p7, %rd16, %rd13;
	@%p7 bra 	$Lt_0_20482;
	bra.uni 	$Lt_0_18946;
$Lt_0_27650:
	mov.f32 	%f30, 0f00000000;    	// 0
	mov.f32 	%f31, 0f00000000;    	// 0
	mov.f32 	%f32, 0f00000000;    	// 0
	mov.f32 	%f33, 0f00000000;    	// 0
	bra.uni 	$Lt_0_18946;
$Lt_0_19202:
	mov.f32 	%f30, 0f00000000;    	// 0
	mov.f32 	%f31, 0f00000000;    	// 0
	mov.f32 	%f32, 0f00000000;    	// 0
	mov.f32 	%f33, 0f00000000;    	// 0
$Lt_0_18946:
	mov.u32 	%r55, 1;
	setp.le.s32 	%p8, %r1, %r55;
	@%p8 bra 	$Lt_0_24066;
	.loc	16	203	0
	mov.u64 	%rd30, __cuda___cuda_local_var_32585_35_non_const_red_acc128;
	cvt.s64.s32 	%rd31, %r2;
	mul.wide.s32 	%rd32, %r2, 4;
	add.u64 	%rd33, %rd30, %rd32;
	mov.f32 	%f106, %f32;
	st.shared.f32 	[%rd33+0], %f106;
	.loc	16	204	0
	mov.f32 	%f107, %f31;
	st.shared.f32 	[%rd33+512], %f107;
	.loc	16	205	0
	mov.f32 	%f108, %f30;
	st.shared.f32 	[%rd33+1024], %f108;
	.loc	16	206	0
	mov.f32 	%f109, %f33;
	st.shared.f32 	[%rd33+1536], %f109;
	.loc	16	208	0
	shr.s32 	%r56, %r1, 31;
	mov.s32 	%r57, 1;
	and.b32 	%r58, %r56, %r57;
	add.s32 	%r59, %r58, %r1;
	shr.s32 	%r60, %r59, 1;
	mov.s32 	%r61, %r60;
	mov.u32 	%r62, 0;
	setp.ne.u32 	%p9, %r60, %r62;
	@!%p9 bra 	$Lt_0_22530;
$Lt_0_23042:
	setp.ge.u32 	%p10, %r6, %r61;
	@%p10 bra 	$Lt_0_23298;
	.loc	16	211	0
	add.u32 	%r63, %r2, %r61;
	cvt.u64.u32 	%rd34, %r63;
	mul.wide.u32 	%rd35, %r63, 4;
	add.u64 	%rd36, %rd30, %rd35;
	ld.shared.f32 	%f110, [%rd36+0];
	add.ftz.f32 	%f106, %f110, %f106;
	st.shared.f32 	[%rd33+0], %f106;
	ld.shared.f32 	%f111, [%rd36+512];
	add.ftz.f32 	%f107, %f111, %f107;
	st.shared.f32 	[%rd33+512], %f107;
	ld.shared.f32 	%f112, [%rd36+1024];
	add.ftz.f32 	%f108, %f112, %f108;
	st.shared.f32 	[%rd33+1024], %f108;
	ld.shared.f32 	%f113, [%rd36+1536];
	add.ftz.f32 	%f109, %f113, %f109;
	st.shared.f32 	[%rd33+1536], %f109;
$Lt_0_23298:
	.loc	16	208	0
	shr.u32 	%r61, %r61, 1;
	mov.u32 	%r64, 0;
	setp.ne.u32 	%p11, %r61, %r64;
	@%p11 bra 	$Lt_0_23042;
$Lt_0_22530:
	.loc	16	215	0
	mov.f32 	%f32, %f106;
	.loc	16	216	0
	mov.f32 	%f31, %f107;
	.loc	16	217	0
	mov.f32 	%f30, %f108;
	.loc	16	218	0
	mov.f32 	%f33, %f109;
	ld.param.s32 	%r65, [__cudaparm_kernel_pair_vflag];
	mov.u32 	%r66, 0;
	setp.le.s32 	%p12, %r65, %r66;
	@%p12 bra 	$Lt_0_24066;
	.loc	16	222	0
	mov.f32 	%f106, %f6;
	st.shared.f32 	[%rd33+0], %f106;
	mov.f32 	%f107, %f8;
	st.shared.f32 	[%rd33+512], %f107;
	mov.f32 	%f108, %f10;
	st.shared.f32 	[%rd33+1024], %f108;
	mov.f32 	%f109, %f12;
	st.shared.f32 	[%rd33+1536], %f109;
	mov.f32 	%f114, %f14;
	st.shared.f32 	[%rd33+2048], %f114;
	mov.f32 	%f115, %f16;
	st.shared.f32 	[%rd33+2560], %f115;
	.loc	16	224	0
	mov.s32 	%r67, %r60;
	@!%p9 bra 	$Lt_0_24578;
$Lt_0_25090:
	setp.ge.u32 	%p13, %r6, %r67;
	@%p13 bra 	$Lt_0_25346;
	.loc	16	227	0
	add.u32 	%r68, %r2, %r67;
	cvt.u64.u32 	%rd37, %r68;
	mul.wide.u32 	%rd38, %r68, 4;
	add.u64 	%rd39, %rd30, %rd38;
	ld.shared.f32 	%f116, [%rd39+0];
	add.ftz.f32 	%f106, %f116, %f106;
	st.shared.f32 	[%rd33+0], %f106;
	ld.shared.f32 	%f117, [%rd39+512];
	add.ftz.f32 	%f107, %f117, %f107;
	st.shared.f32 	[%rd33+512], %f107;
	ld.shared.f32 	%f118, [%rd39+1024];
	add.ftz.f32 	%f108, %f118, %f108;
	st.shared.f32 	[%rd33+1024], %f108;
	ld.shared.f32 	%f119, [%rd39+1536];
	add.ftz.f32 	%f109, %f119, %f109;
	st.shared.f32 	[%rd33+1536], %f109;
	ld.shared.f32 	%f120, [%rd39+2048];
	add.ftz.f32 	%f114, %f120, %f114;
	st.shared.f32 	[%rd33+2048], %f114;
	ld.shared.f32 	%f121, [%rd39+2560];
	add.ftz.f32 	%f115, %f121, %f115;
	st.shared.f32 	[%rd33+2560], %f115;
$Lt_0_25346:
	.loc	16	224	0
	shr.u32 	%r67, %r67, 1;
	mov.u32 	%r69, 0;
	setp.ne.u32 	%p14, %r67, %r69;
	@%p14 bra 	$Lt_0_25090;
$Lt_0_24578:
	.loc	16	232	0
	mov.f32 	%f6, %f106;
	mov.f32 	%f8, %f107;
	mov.f32 	%f10, %f108;
	mov.f32 	%f12, %f109;
	mov.f32 	%f14, %f114;
	mov.f32 	%f16, %f115;
$Lt_0_24066:
$Lt_0_22018:
	selp.s32 	%r70, 1, 0, %p1;
	mov.s32 	%r71, 0;
	set.eq.u32.s32 	%r72, %r6, %r71;
	neg.s32 	%r73, %r72;
	and.b32 	%r74, %r70, %r73;
	mov.u32 	%r75, 0;
	setp.eq.s32 	%p15, %r74, %r75;
	@%p15 bra 	$Lt_0_26114;
	.loc	16	238	0
	cvt.s64.s32 	%rd40, %r9;
	ld.param.u64 	%rd41, [__cudaparm_kernel_pair_engv];
	mul.wide.s32 	%rd42, %r9, 4;
	add.u64 	%rd43, %rd41, %rd42;
	ld.param.s32 	%r76, [__cudaparm_kernel_pair_eflag];
	mov.u32 	%r77, 0;
	setp.le.s32 	%p16, %r76, %r77;
	@%p16 bra 	$Lt_0_26626;
	.loc	16	240	0
	mov.f32 	%f122, 0f00000000;   	// 0
	st.global.f32 	[%rd43+0], %f122;
	.loc	16	241	0
	cvt.s64.s32 	%rd44, %r10;
	mul.wide.s32 	%rd45, %r10, 4;
	add.u64 	%rd46, %rd45, %rd43;
	.loc	16	242	0
	st.global.f32 	[%rd46+0], %f33;
	.loc	16	243	0
	add.u64 	%rd43, %rd45, %rd46;
$Lt_0_26626:
	ld.param.s32 	%r78, [__cudaparm_kernel_pair_vflag];
	mov.u32 	%r79, 0;
	setp.le.s32 	%p17, %r78, %r79;
	@%p17 bra 	$Lt_0_27138;
	.loc	16	247	0
	mov.f32 	%f123, %f6;
	st.global.f32 	[%rd43+0], %f123;
	.loc	16	248	0
	cvt.s64.s32 	%rd47, %r10;
	mul.wide.s32 	%rd48, %r10, 4;
	add.u64 	%rd49, %rd48, %rd43;
	.loc	16	247	0
	mov.f32 	%f124, %f8;
	st.global.f32 	[%rd49+0], %f124;
	.loc	16	248	0
	add.u64 	%rd50, %rd48, %rd49;
	.loc	16	247	0
	mov.f32 	%f125, %f10;
	st.global.f32 	[%rd50+0], %f125;
	.loc	16	248	0
	add.u64 	%rd51, %rd48, %rd50;
	.loc	16	247	0
	mov.f32 	%f126, %f12;
	st.global.f32 	[%rd51+0], %f126;
	.loc	16	248	0
	add.u64 	%rd43, %rd48, %rd51;
	.loc	16	247	0
	mov.f32 	%f127, %f14;
	st.global.f32 	[%rd43+0], %f127;
	mov.f32 	%f128, %f16;
	add.u64 	%rd52, %rd48, %rd43;
	st.global.f32 	[%rd52+0], %f128;
$Lt_0_27138:
	.loc	16	251	0
	ld.param.u64 	%rd53, [__cudaparm_kernel_pair_ans];
	mul.lo.u64 	%rd54, %rd40, 16;
	add.u64 	%rd55, %rd53, %rd54;
	mov.f32 	%f129, %f130;
	st.global.v4.f32 	[%rd55+0], {%f32,%f31,%f30,%f129};
$Lt_0_26114:
	.loc	16	253	0
	exit;
$LDWend_kernel_pair:
	} // kernel_pair

	.entry kernel_pair_fast (
		.param .u64 __cudaparm_kernel_pair_fast_x_,
		.param .u64 __cudaparm_kernel_pair_fast_lj1_in,
		.param .u64 __cudaparm_kernel_pair_fast_lj3_in,
		.param .u64 __cudaparm_kernel_pair_fast_sp_cl_in,
		.param .u64 __cudaparm_kernel_pair_fast_dev_nbor,
		.param .u64 __cudaparm_kernel_pair_fast_dev_packed,
		.param .u64 __cudaparm_kernel_pair_fast_ans,
		.param .u64 __cudaparm_kernel_pair_fast_engv,
		.param .s32 __cudaparm_kernel_pair_fast_eflag,
		.param .s32 __cudaparm_kernel_pair_fast_vflag,
		.param .s32 __cudaparm_kernel_pair_fast_inum,
		.param .s32 __cudaparm_kernel_pair_fast_nbor_pitch,
		.param .u64 __cudaparm_kernel_pair_fast_q_,
		.param .f32 __cudaparm_kernel_pair_fast_cut_coulsq,
		.param .f32 __cudaparm_kernel_pair_fast_qqrd2e,
		.param .f32 __cudaparm_kernel_pair_fast_g_ewald,
		.param .s32 __cudaparm_kernel_pair_fast_t_per_atom)
	{
	.reg .u32 %r<82>;
	.reg .u64 %rd<61>;
	.reg .f32 %f<129>;
	.reg .pred %p<20>;
	.shared .align 4 .b8 __cuda___cuda_local_var_32653_33_non_const_sp_cl3304[16];
	.shared .align 4 .b8 __cuda___cuda_local_var_32740_35_non_const_red_acc3320[3072];
	// __cuda_local_var_32658_10_non_const_f = 48
	// __cuda_local_var_32662_9_non_const_virial = 16
	.loc	16	263	0
$LDWbegin_kernel_pair_fast:
	cvt.s32.u32 	%r1, %tid.x;
	mov.u32 	%r2, 3;
	setp.gt.s32 	%p1, %r1, %r2;
	@%p1 bra 	$Lt_1_19714;
	.loc	16	271	0
	mov.u64 	%rd1, __cuda___cuda_local_var_32653_33_non_const_sp_cl3304;
	cvt.s64.s32 	%rd2, %r1;
	mul.wide.s32 	%rd3, %r1, 4;
	ld.param.u64 	%rd4, [__cudaparm_kernel_pair_fast_sp_cl_in];
	add.u64 	%rd5, %rd4, %rd3;
	ld.global.f32 	%f1, [%rd5+0];
	add.u64 	%rd6, %rd3, %rd1;
	st.shared.f32 	[%rd6+0], %f1;
$Lt_1_19714:
	mov.u64 	%rd1, __cuda___cuda_local_var_32653_33_non_const_sp_cl3304;
	.loc	16	280	0
	mov.f32 	%f2, 0f00000000;     	// 0
	mov.f32 	%f3, %f2;
	mov.f32 	%f4, 0f00000000;     	// 0
	mov.f32 	%f5, %f4;
	mov.f32 	%f6, 0f00000000;     	// 0
	mov.f32 	%f7, %f6;
	mov.f32 	%f8, 0f00000000;     	// 0
	mov.f32 	%f9, %f8;
	mov.f32 	%f10, 0f00000000;    	// 0
	mov.f32 	%f11, %f10;
	mov.f32 	%f12, 0f00000000;    	// 0
	mov.f32 	%f13, %f12;
	.loc	16	282	0
	bar.sync 	0;
	ld.param.s32 	%r3, [__cudaparm_kernel_pair_fast_t_per_atom];
	div.s32 	%r4, %r1, %r3;
	cvt.s32.u32 	%r5, %ntid.x;
	div.s32 	%r6, %r5, %r3;
	rem.s32 	%r7, %r1, %r3;
	cvt.s32.u32 	%r8, %ctaid.x;
	mul.lo.s32 	%r9, %r8, %r6;
	add.s32 	%r10, %r4, %r9;
	ld.param.s32 	%r11, [__cudaparm_kernel_pair_fast_inum];
	setp.lt.s32 	%p2, %r10, %r11;
	@!%p2 bra 	$Lt_1_20482;
	.loc	16	286	0
	cvt.s64.s32 	%rd7, %r10;
	mul.wide.s32 	%rd8, %r10, 4;
	ld.param.u64 	%rd9, [__cudaparm_kernel_pair_fast_dev_nbor];
	add.u64 	%rd10, %rd8, %rd9;
	ld.global.s32 	%r12, [%rd10+0];
	.loc	16	288	0
	ld.param.s32 	%r13, [__cudaparm_kernel_pair_fast_nbor_pitch];
	cvt.s64.s32 	%rd11, %r13;
	mul.wide.s32 	%rd12, %r13, 4;
	add.u64 	%rd13, %rd12, %rd10;
	ld.global.s32 	%r14, [%rd13+0];
	add.u64 	%rd14, %rd12, %rd13;
	ld.param.u64 	%rd15, [__cudaparm_kernel_pair_fast_dev_packed];
	setp.ne.u64 	%p3, %rd15, %rd9;
	@%p3 bra 	$Lt_1_20994;
	.loc	16	294	0
	cvt.s32.s64 	%r15, %rd11;
	mul.lo.s32 	%r16, %r15, %r14;
	cvt.s64.s32 	%rd16, %r16;
	mul.wide.s32 	%rd17, %r16, 4;
	add.u64 	%rd18, %rd14, %rd17;
	.loc	16	295	0
	mul.lo.s32 	%r17, %r7, %r15;
	cvt.s64.s32 	%rd19, %r17;
	mul.wide.s32 	%rd20, %r17, 4;
	add.u64 	%rd21, %rd14, %rd20;
	.loc	16	296	0
	mul.lo.s32 	%r18, %r15, %r3;
	bra.uni 	$Lt_1_20738;
$Lt_1_20994:
	.loc	16	298	0
	ld.global.s32 	%r19, [%rd14+0];
	cvt.s64.s32 	%rd22, %r19;
	mul.wide.s32 	%rd23, %r19, 4;
	add.u64 	%rd24, %rd15, %rd23;
	.loc	16	299	0
	cvt.s64.s32 	%rd25, %r14;
	mul.wide.s32 	%rd26, %r14, 4;
	add.u64 	%rd18, %rd24, %rd26;
	.loc	16	300	0
	mov.s32 	%r18, %r3;
	.loc	16	301	0
	cvt.s64.s32 	%rd27, %r7;
	mul.wide.s32 	%rd28, %r7, 4;
	add.u64 	%rd21, %rd24, %rd28;
$Lt_1_20738:
	.loc	16	304	0
	mov.u32 	%r20, %r12;
	mov.s32 	%r21, 0;
	mov.u32 	%r22, %r21;
	mov.s32 	%r23, 0;
	mov.u32 	%r24, %r23;
	mov.s32 	%r25, 0;
	mov.u32 	%r26, %r25;
	tex.1d.v4.f32.s32 {%f14,%f15,%f16,%f17},[pos_tex,{%r20,%r22,%r24,%r26}];
	mov.f32 	%f18, %f14;
	mov.f32 	%f19, %f15;
	mov.f32 	%f20, %f16;
	.loc	16	305	0
	mov.u32 	%r27, %r12;
	mov.s32 	%r28, 0;
	mov.u32 	%r29, %r28;
	mov.s32 	%r30, 0;
	mov.u32 	%r31, %r30;
	mov.s32 	%r32, 0;
	mov.u32 	%r33, %r32;
	tex.1d.v4.f32.s32 {%f21,%f22,%f23,%f24},[q_tex,{%r27,%r29,%r31,%r33}];
	mov.f32 	%f25, %f21;
	setp.ge.u64 	%p4, %rd21, %rd18;
	@%p4 bra 	$Lt_1_28930;
	cvt.s64.s32 	%rd29, %r18;
	ld.param.f32 	%f26, [__cudaparm_kernel_pair_fast_cut_coulsq];
	mov.f32 	%f27, 0f00000000;    	// 0
	mov.f32 	%f28, 0f00000000;    	// 0
	mov.f32 	%f29, 0f00000000;    	// 0
	mov.f32 	%f30, 0f00000000;    	// 0
$Lt_1_21762:
 //<loop> Loop body line 305, nesting depth: 1, estimated iterations: unknown
	.loc	16	308	0
	ld.global.s32 	%r34, [%rd21+0];
	.loc	16	311	0
	mov.f32 	%f31, 0f3f800000;    	// 1
	shr.s32 	%r35, %r34, 30;
	and.b32 	%r36, %r35, 3;
	cvt.s64.s32 	%rd30, %r36;
	mul.wide.s32 	%rd31, %r36, 4;
	add.u64 	%rd32, %rd1, %rd31;
	ld.shared.f32 	%f32, [%rd32+0];
	sub.ftz.f32 	%f33, %f31, %f32;
	.loc	16	314	0
	and.b32 	%r37, %r34, 1073741823;
	mov.u32 	%r38, %r37;
	mov.s32 	%r39, 0;
	mov.u32 	%r40, %r39;
	mov.s32 	%r41, 0;
	mov.u32 	%r42, %r41;
	mov.s32 	%r43, 0;
	mov.u32 	%r44, %r43;
	tex.1d.v4.f32.s32 {%f34,%f35,%f36,%f37},[pos_tex,{%r38,%r40,%r42,%r44}];
	mov.f32 	%f38, %f34;
	mov.f32 	%f39, %f35;
	mov.f32 	%f40, %f36;
	sub.ftz.f32 	%f41, %f19, %f39;
	sub.ftz.f32 	%f42, %f18, %f38;
	sub.ftz.f32 	%f43, %f20, %f40;
	mul.ftz.f32 	%f44, %f41, %f41;
	fma.rn.ftz.f32 	%f45, %f42, %f42, %f44;
	fma.rn.ftz.f32 	%f46, %f43, %f43, %f45;
	setp.lt.ftz.f32 	%p5, %f46, %f26;
	@!%p5 bra 	$Lt_1_22530;
	.loc	16	330	0
	sqrt.approx.ftz.f32 	%f47, %f46;
	ld.param.f32 	%f48, [__cudaparm_kernel_pair_fast_g_ewald];
	mul.ftz.f32 	%f49, %f48, %f47;
	mul.ftz.f32 	%f50, %f49, %f49;
	mov.f32 	%f51, 0f3f800000;    	// 1
	mov.f32 	%f52, 0f3ea7ba05;    	// 0.327591
	fma.rn.ftz.f32 	%f53, %f52, %f49, %f51;
	neg.ftz.f32 	%f54, %f50;
	rcp.approx.ftz.f32 	%f55, %f53;
	mov.f32 	%f56, 0f3fb8aa3b;    	// 1.4427
	mul.ftz.f32 	%f57, %f54, %f56;
	ex2.approx.ftz.f32 	%f58, %f57;
	mov.f32 	%f59, 0f3e827906;    	// 0.25483
	mov.f32 	%f60, 0fbe91a98e;    	// -0.284497
	mov.f32 	%f61, 0f3fb5f0e3;    	// 1.42141
	mov.f32 	%f62, 0fbfba00e3;    	// -1.45315
	mov.f32 	%f63, 0f3f87dc22;    	// 1.06141
	fma.rn.ftz.f32 	%f64, %f63, %f55, %f62;
	fma.rn.ftz.f32 	%f65, %f55, %f64, %f61;
	fma.rn.ftz.f32 	%f66, %f55, %f65, %f60;
	fma.rn.ftz.f32 	%f67, %f55, %f66, %f59;
	mul.ftz.f32 	%f68, %f55, %f67;
	mul.ftz.f32 	%f69, %f58, %f68;
	.loc	16	331	0
	mov.u32 	%r45, %r37;
	mov.s32 	%r46, 0;
	mov.u32 	%r47, %r46;
	mov.s32 	%r48, 0;
	mov.u32 	%r49, %r48;
	mov.s32 	%r50, 0;
	mov.u32 	%r51, %r50;
	tex.1d.v4.f32.s32 {%f70,%f71,%f72,%f73},[q_tex,{%r45,%r47,%r49,%r51}];
	mov.f32 	%f74, %f70;
	.loc	16	332	0
	ld.param.f32 	%f75, [__cudaparm_kernel_pair_fast_qqrd2e];
	mul.ftz.f32 	%f76, %f75, %f25;
	mul.ftz.f32 	%f77, %f76, %f74;
	div.approx.ftz.f32 	%f78, %f77, %f47;
	mov.f32 	%f79, 0f3f906ebb;    	// 1.12838
	mul.ftz.f32 	%f80, %f49, %f79;
	fma.rn.ftz.f32 	%f81, %f58, %f80, %f69;
	sub.ftz.f32 	%f82, %f81, %f33;
	mul.ftz.f32 	%f83, %f78, %f82;
	rcp.approx.ftz.f32 	%f84, %f46;
	mul.ftz.f32 	%f85, %f83, %f84;
	.loc	16	334	0
	fma.rn.ftz.f32 	%f29, %f42, %f85, %f29;
	.loc	16	335	0
	fma.rn.ftz.f32 	%f28, %f41, %f85, %f28;
	.loc	16	336	0
	fma.rn.ftz.f32 	%f27, %f43, %f85, %f27;
	.loc	16	323	0
	sub.ftz.f32 	%f86, %f69, %f33;
	fma.rn.ftz.f32 	%f87, %f78, %f86, %f30;
	ld.param.s32 	%r52, [__cudaparm_kernel_pair_fast_eflag];
	mov.s32 	%r53, 0;
	setp.gt.s32 	%p6, %r52, %r53;
	selp.f32 	%f30, %f87, %f30, %p6;
	ld.param.s32 	%r54, [__cudaparm_kernel_pair_fast_vflag];
	mov.u32 	%r55, 0;
	setp.le.s32 	%p7, %r54, %r55;
	@%p7 bra 	$Lt_1_22530;
	.loc	16	342	0
	mov.f32 	%f88, %f3;
	mul.ftz.f32 	%f89, %f42, %f42;
	fma.rn.ftz.f32 	%f90, %f85, %f89, %f88;
	mov.f32 	%f3, %f90;
	.loc	16	343	0
	mov.f32 	%f91, %f5;
	fma.rn.ftz.f32 	%f92, %f85, %f44, %f91;
	mov.f32 	%f5, %f92;
	.loc	16	344	0
	mov.f32 	%f93, %f7;
	mul.ftz.f32 	%f94, %f43, %f43;
	fma.rn.ftz.f32 	%f95, %f85, %f94, %f93;
	mov.f32 	%f7, %f95;
	.loc	16	345	0
	mov.f32 	%f96, %f9;
	mul.ftz.f32 	%f97, %f41, %f42;
	fma.rn.ftz.f32 	%f98, %f85, %f97, %f96;
	mov.f32 	%f9, %f98;
	.loc	16	346	0
	mov.f32 	%f99, %f11;
	mul.ftz.f32 	%f100, %f42, %f43;
	fma.rn.ftz.f32 	%f101, %f85, %f100, %f99;
	mov.f32 	%f11, %f101;
	.loc	16	347	0
	mul.ftz.f32 	%f102, %f41, %f43;
	fma.rn.ftz.f32 	%f12, %f85, %f102, %f12;
	mov.f32 	%f13, %f12;
$Lt_1_22530:
$Lt_1_22018:
	.loc	16	307	0
	mul.lo.u64 	%rd33, %rd29, 4;
	add.u64 	%rd21, %rd21, %rd33;
	setp.lt.u64 	%p8, %rd21, %rd18;
	@%p8 bra 	$Lt_1_21762;
	bra.uni 	$Lt_1_20226;
$Lt_1_28930:
	mov.f32 	%f27, 0f00000000;    	// 0
	mov.f32 	%f28, 0f00000000;    	// 0
	mov.f32 	%f29, 0f00000000;    	// 0
	mov.f32 	%f30, 0f00000000;    	// 0
	bra.uni 	$Lt_1_20226;
$Lt_1_20482:
	mov.f32 	%f27, 0f00000000;    	// 0
	mov.f32 	%f28, 0f00000000;    	// 0
	mov.f32 	%f29, 0f00000000;    	// 0
	mov.f32 	%f30, 0f00000000;    	// 0
$Lt_1_20226:
	mov.u32 	%r56, 1;
	setp.le.s32 	%p9, %r3, %r56;
	@%p9 bra 	$Lt_1_25346;
	.loc	16	358	0
	mov.u64 	%rd34, __cuda___cuda_local_var_32740_35_non_const_red_acc3320;
	cvt.s64.s32 	%rd35, %r1;
	mul.wide.s32 	%rd36, %r1, 4;
	add.u64 	%rd37, %rd34, %rd36;
	mov.f32 	%f103, %f29;
	st.shared.f32 	[%rd37+0], %f103;
	.loc	16	359	0
	mov.f32 	%f104, %f28;
	st.shared.f32 	[%rd37+512], %f104;
	.loc	16	360	0
	mov.f32 	%f105, %f27;
	st.shared.f32 	[%rd37+1024], %f105;
	.loc	16	361	0
	mov.f32 	%f106, %f30;
	st.shared.f32 	[%rd37+1536], %f106;
	.loc	16	363	0
	shr.s32 	%r57, %r3, 31;
	mov.s32 	%r58, 1;
	and.b32 	%r59, %r57, %r58;
	add.s32 	%r60, %r59, %r3;
	shr.s32 	%r61, %r60, 1;
	mov.s32 	%r62, %r61;
	mov.u32 	%r63, 0;
	setp.ne.u32 	%p10, %r61, %r63;
	@!%p10 bra 	$Lt_1_23810;
$Lt_1_24322:
	setp.ge.u32 	%p11, %r7, %r62;
	@%p11 bra 	$Lt_1_24578;
	.loc	16	366	0
	add.u32 	%r64, %r1, %r62;
	cvt.u64.u32 	%rd38, %r64;
	mul.wide.u32 	%rd39, %r64, 4;
	add.u64 	%rd40, %rd34, %rd39;
	ld.shared.f32 	%f107, [%rd40+0];
	add.ftz.f32 	%f103, %f107, %f103;
	st.shared.f32 	[%rd37+0], %f103;
	ld.shared.f32 	%f108, [%rd40+512];
	add.ftz.f32 	%f104, %f108, %f104;
	st.shared.f32 	[%rd37+512], %f104;
	ld.shared.f32 	%f109, [%rd40+1024];
	add.ftz.f32 	%f105, %f109, %f105;
	st.shared.f32 	[%rd37+1024], %f105;
	ld.shared.f32 	%f110, [%rd40+1536];
	add.ftz.f32 	%f106, %f110, %f106;
	st.shared.f32 	[%rd37+1536], %f106;
$Lt_1_24578:
	.loc	16	363	0
	shr.u32 	%r62, %r62, 1;
	mov.u32 	%r65, 0;
	setp.ne.u32 	%p12, %r62, %r65;
	@%p12 bra 	$Lt_1_24322;
$Lt_1_23810:
	.loc	16	370	0
	mov.f32 	%f29, %f103;
	.loc	16	371	0
	mov.f32 	%f28, %f104;
	.loc	16	372	0
	mov.f32 	%f27, %f105;
	.loc	16	373	0
	mov.f32 	%f30, %f106;
	ld.param.s32 	%r66, [__cudaparm_kernel_pair_fast_vflag];
	mov.u32 	%r67, 0;
	setp.le.s32 	%p13, %r66, %r67;
	@%p13 bra 	$Lt_1_25346;
	.loc	16	377	0
	mov.f32 	%f103, %f3;
	st.shared.f32 	[%rd37+0], %f103;
	mov.f32 	%f104, %f5;
	st.shared.f32 	[%rd37+512], %f104;
	mov.f32 	%f105, %f7;
	st.shared.f32 	[%rd37+1024], %f105;
	mov.f32 	%f106, %f9;
	st.shared.f32 	[%rd37+1536], %f106;
	mov.f32 	%f111, %f11;
	st.shared.f32 	[%rd37+2048], %f111;
	mov.f32 	%f112, %f13;
	st.shared.f32 	[%rd37+2560], %f112;
	.loc	16	379	0
	mov.s32 	%r68, %r61;
	@!%p10 bra 	$Lt_1_25858;
$Lt_1_26370:
	setp.ge.u32 	%p14, %r7, %r68;
	@%p14 bra 	$Lt_1_26626;
	.loc	16	382	0
	add.u32 	%r69, %r1, %r68;
	cvt.u64.u32 	%rd41, %r69;
	mul.wide.u32 	%rd42, %r69, 4;
	add.u64 	%rd43, %rd34, %rd42;
	ld.shared.f32 	%f113, [%rd43+0];
	add.ftz.f32 	%f103, %f113, %f103;
	st.shared.f32 	[%rd37+0], %f103;
	ld.shared.f32 	%f114, [%rd43+512];
	add.ftz.f32 	%f104, %f114, %f104;
	st.shared.f32 	[%rd37+512], %f104;
	ld.shared.f32 	%f115, [%rd43+1024];
	add.ftz.f32 	%f105, %f115, %f105;
	st.shared.f32 	[%rd37+1024], %f105;
	ld.shared.f32 	%f116, [%rd43+1536];
	add.ftz.f32 	%f106, %f116, %f106;
	st.shared.f32 	[%rd37+1536], %f106;
	ld.shared.f32 	%f117, [%rd43+2048];
	add.ftz.f32 	%f111, %f117, %f111;
	st.shared.f32 	[%rd37+2048], %f111;
	ld.shared.f32 	%f118, [%rd43+2560];
	add.ftz.f32 	%f112, %f118, %f112;
	st.shared.f32 	[%rd37+2560], %f112;
$Lt_1_26626:
	.loc	16	379	0
	shr.u32 	%r68, %r68, 1;
	mov.u32 	%r70, 0;
	setp.ne.u32 	%p15, %r68, %r70;
	@%p15 bra 	$Lt_1_26370;
$Lt_1_25858:
	.loc	16	387	0
	mov.f32 	%f3, %f103;
	mov.f32 	%f5, %f104;
	mov.f32 	%f7, %f105;
	mov.f32 	%f9, %f106;
	mov.f32 	%f11, %f111;
	mov.f32 	%f13, %f112;
$Lt_1_25346:
$Lt_1_23298:
	selp.s32 	%r71, 1, 0, %p2;
	mov.s32 	%r72, 0;
	set.eq.u32.s32 	%r73, %r7, %r72;
	neg.s32 	%r74, %r73;
	and.b32 	%r75, %r71, %r74;
	mov.u32 	%r76, 0;
	setp.eq.s32 	%p16, %r75, %r76;
	@%p16 bra 	$Lt_1_27394;
	.loc	16	393	0
	cvt.s64.s32 	%rd44, %r10;
	ld.param.u64 	%rd45, [__cudaparm_kernel_pair_fast_engv];
	mul.wide.s32 	%rd46, %r10, 4;
	add.u64 	%rd47, %rd45, %rd46;
	ld.param.s32 	%r77, [__cudaparm_kernel_pair_fast_eflag];
	mov.u32 	%r78, 0;
	setp.le.s32 	%p17, %r77, %r78;
	@%p17 bra 	$Lt_1_27906;
	.loc	16	395	0
	mov.f32 	%f119, 0f00000000;   	// 0
	st.global.f32 	[%rd47+0], %f119;
	.loc	16	396	0
	cvt.s64.s32 	%rd48, %r11;
	mul.wide.s32 	%rd49, %r11, 4;
	add.u64 	%rd50, %rd49, %rd47;
	.loc	16	397	0
	st.global.f32 	[%rd50+0], %f30;
	.loc	16	398	0
	add.u64 	%rd47, %rd49, %rd50;
$Lt_1_27906:
	ld.param.s32 	%r79, [__cudaparm_kernel_pair_fast_vflag];
	mov.u32 	%r80, 0;
	setp.le.s32 	%p18, %r79, %r80;
	@%p18 bra 	$Lt_1_28418;
	.loc	16	402	0
	mov.f32 	%f120, %f3;
	st.global.f32 	[%rd47+0], %f120;
	.loc	16	403	0
	cvt.s64.s32 	%rd51, %r11;
	mul.wide.s32 	%rd52, %r11, 4;
	add.u64 	%rd53, %rd52, %rd47;
	.loc	16	402	0
	mov.f32 	%f121, %f5;
	st.global.f32 	[%rd53+0], %f121;
	.loc	16	403	0
	add.u64 	%rd54, %rd52, %rd53;
	.loc	16	402	0
	mov.f32 	%f122, %f7;
	st.global.f32 	[%rd54+0], %f122;
	.loc	16	403	0
	add.u64 	%rd55, %rd52, %rd54;
	.loc	16	402	0
	mov.f32 	%f123, %f9;
	st.global.f32 	[%rd55+0], %f123;
	.loc	16	403	0
	add.u64 	%rd47, %rd52, %rd55;
	.loc	16	402	0
	mov.f32 	%f124, %f11;
	st.global.f32 	[%rd47+0], %f124;
	mov.f32 	%f125, %f13;
	add.u64 	%rd56, %rd52, %rd47;
	st.global.f32 	[%rd56+0], %f125;
$Lt_1_28418:
	.loc	16	406	0
	ld.param.u64 	%rd57, [__cudaparm_kernel_pair_fast_ans];
	mul.lo.u64 	%rd58, %rd44, 16;
	add.u64 	%rd59, %rd57, %rd58;
	mov.f32 	%f126, %f127;
	st.global.v4.f32 	[%rd59+0], {%f29,%f28,%f27,%f126};
$Lt_1_27394:
	.loc	16	408	0
	exit;
$LDWend_kernel_pair_fast:
	} // kernel_pair_fast