	.version 2.1
	.target sm_20
	// compiled with /usr/local/cuda/open64/lib//be
	// nvopencc 3.2 built on 2010-09-08

	//-----------------------------------------------------------
	// Compiling /tmp/tmpxft_000041ca_00000000-9_kernels.cpp3.i (/tmp/ccBI#.H4WGVs)
	//-----------------------------------------------------------

	//-----------------------------------------------------------
	// 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_000041ca_00000000-8_kernels.cudafe2.gpu"
	.file	3	"/usr/lib/gcc/x86_64-linux-gnu/4.4.3/include/stddef.h"
	.file	4	"/usr/local/cuda/bin/../include/crt/device_runtime.h"
	.file	5	"/usr/local/cuda/bin/../include/host_defines.h"
	.file	6	"/usr/local/cuda/bin/../include/builtin_types.h"
	.file	7	"/usr/local/cuda/bin/../include/device_types.h"
	.file	8	"/usr/local/cuda/bin/../include/driver_types.h"
	.file	9	"/usr/local/cuda/bin/../include/surface_types.h"
	.file	10	"/usr/local/cuda/bin/../include/texture_types.h"
	.file	11	"/usr/local/cuda/bin/../include/vector_types.h"
	.file	12	"/usr/local/cuda/bin/../include/device_launch_parameters.h"
	.file	13	"/usr/local/cuda/bin/../include/crt/storage_class.h"
	.file	14	"/usr/include/bits/types.h"
	.file	15	"/usr/include/time.h"
	.file	16	"/usr/local/cuda/bin/../include/texture_fetch_functions.h"
	.file	17	"../../tests/ptx/synthetic/kernels.cu"
	.file	18	"/usr/local/cuda/bin/../include/common_functions.h"
	.file	19	"/usr/local/cuda/bin/../include/math_functions.h"
	.file	20	"/usr/local/cuda/bin/../include/math_constants.h"
	.file	21	"/usr/local/cuda/bin/../include/device_functions.h"
	.file	22	"/usr/local/cuda/bin/../include/sm_11_atomic_functions.h"
	.file	23	"/usr/local/cuda/bin/../include/sm_12_atomic_functions.h"
	.file	24	"/usr/local/cuda/bin/../include/sm_13_double_functions.h"
	.file	25	"/usr/local/cuda/bin/../include/sm_20_atomic_functions.h"
	.file	26	"/usr/local/cuda/bin/../include/sm_20_intrinsics.h"
	.file	27	"/usr/local/cuda/bin/../include/surface_functions.h"
	.file	28	"/usr/local/cuda/bin/../include/math_functions_dbl_ptx3.h"

	.const .align 4 .b8 __cudart_i2opi_f[24] = {65,144,67,60,153,149,98,219,192,221,52,245,209,87,39,252,41,21,68,78,110,131,249,162};

	.entry _Z17k_sequenceLoopingPfi (
		.param .u64 __cudaparm__Z17k_sequenceLoopingPfi_ptr,
		.param .s32 __cudaparm__Z17k_sequenceLoopingPfi_N)
	{
	.reg .u32 %r<79>;
	.reg .u64 %rd<12>;
	.reg .f32 %f<37>;
	.reg .pred %p<14>;
	.local .align 4 .b8 __cuda___cuda_result_3212[28];
	.loc	17	6	0
$LDWbegin__Z17k_sequenceLoopingPfi:
	.loc	19	1702	0
	cvt.s32.u32 	%r1, %tid.x;
	cvt.rn.f32.s32 	%f1, %r1;
	abs.f32 	%f2, %f1;
	mov.f32 	%f3, 0f7f800000;     	// ((1.0F)/(0.0F))
	setp.eq.f32 	%p1, %f2, %f3;
	@!%p1 bra 	$Lt_0_11778;
	.loc	19	1703	0
	neg.f32 	%f4, %f1;
	add.rn.f32 	%f5, %f1, %f4;
	bra.uni 	$LDWendi___isinff_177_1;
$Lt_0_11778:
	mov.f32 	%f6, 0f473ba700;     	// 48039
	setp.gt.f32 	%p2, %f2, %f6;
	@!%p2 bra 	$Lt_0_12290;
	.loc	19	1396	0
	mov.b32 	%r2, %f1;
	and.b32 	%r3, %r2, -2147483648;
	mov.s32 	%r4, %r3;
	.loc	19	24	0
	shl.b32 	%r5, %r2, 8;
	mov.u64 	%rd1, __cudart_i2opi_f;
	mov.u64 	%rd2, __cuda___cuda_result_3212;
	or.b32 	%r6, %r5, -2147483648;
	mov.s32 	%r7, 0;
	mov.u32 	%r8, 0;
$Lt_0_13314:
	.pragma "nounroll";
 //<loop> Loop body line 24, nesting depth: 1, iterations: 6
	.loc	19	1410	0
	ld.const.u32 	%r9, [%rd1+0];
	mul.lo.u32 	%r10, %r6, %r9;
	add.u32 	%r11, %r10, %r8;
	.loc	19	1411	0
	set.gt.u32.u32 	%r12, %r10, %r11;
	neg.s32 	%r13, %r12;
	mul.hi.u32 	%r14, %r9, %r6;
	add.u32 	%r8, %r13, %r14;
	.loc	19	1412	0
	st.local.u32 	[%rd2+0], %r11;
	add.s32 	%r7, %r7, 1;
	add.u64 	%rd2, %rd2, 4;
	add.u64 	%rd1, %rd1, 4;
	mov.u32 	%r15, 6;
	setp.ne.s32 	%p3, %r7, %r15;
	@%p3 bra 	$Lt_0_13314;
	.loc	19	1414	0
	st.local.u32 	[__cuda___cuda_result_3212+24], %r8;
	.loc	19	1419	0
	shl.b32 	%r16, %r2, 1;
	shr.u32 	%r17, %r16, 24;
	sub.u32 	%r18, %r17, 128;
	mov.u64 	%rd3, __cuda___cuda_result_3212;
	shr.u32 	%r19, %r18, 5;
	mov.s32 	%r20, 4;
	sub.s32 	%r21, %r20, %r19;
	cvt.s64.s32 	%rd4, %r21;
	mul.wide.s32 	%rd5, %r21, 4;
	add.u64 	%rd6, %rd3, %rd5;
	ld.local.u32 	%r8, [%rd6+8];
	.loc	19	1420	0
	ld.local.u32 	%r22, [%rd6+4];
	and.b32 	%r23, %r18, 31;
	mov.u32 	%r24, 0;
	setp.eq.u32 	%p4, %r23, %r24;
	@%p4 bra 	$Lt_0_13826;
	.loc	19	1423	0
	mov.s32 	%r25, 32;
	sub.s32 	%r26, %r25, %r23;
	shr.u32 	%r27, %r22, %r26;
	shl.b32 	%r28, %r8, %r23;
	add.u32 	%r8, %r27, %r28;
	.loc	19	1424	0
	ld.local.u32 	%r29, [%rd6+0];
	shr.u32 	%r30, %r29, %r26;
	shl.b32 	%r31, %r22, %r23;
	add.u32 	%r22, %r30, %r31;
$Lt_0_13826:
	.loc	19	1426	0
	shr.u32 	%r7, %r8, 30;
	.loc	19	1428	0
	shr.u32 	%r32, %r22, 30;
	shl.b32 	%r33, %r8, 2;
	add.u32 	%r8, %r32, %r33;
	.loc	19	1429	0
	shl.b32 	%r22, %r22, 2;
	mov.u32 	%r34, 0;
	setp.eq.u32 	%p5, %r22, %r34;
	@%p5 bra 	$Lt_0_14594;
	.loc	19	1430	0
	add.u32 	%r35, %r8, 1;
	mov.u32 	%r36, -2147483648;
	set.gt.u32.u32 	%r37, %r35, %r36;
	neg.s32 	%r38, %r37;
	bra.uni 	$Lt_0_14338;
$Lt_0_14594:
	mov.u32 	%r39, -2147483648;
	set.gt.u32.u32 	%r40, %r8, %r39;
	neg.s32 	%r38, %r40;
$Lt_0_14338:
	.loc	19	1431	0
	add.u32 	%r41, %r7, %r38;
	.loc	19	1430	0
	neg.s32 	%r42, %r41;
	mov.u32 	%r43, 0;
	setp.ne.u32 	%p6, %r3, %r43;
	selp.s32 	%r7, %r42, %r41, %p6;
	mov.u32 	%r44, 0;
	setp.eq.u32 	%p7, %r38, %r44;
	@%p7 bra 	$Lt_0_14850;
	.loc	19	1436	0
	neg.s32 	%r22, %r22;
	.loc	19	1438	0
	mov.u32 	%r45, 0;
	set.eq.u32.u32 	%r46, %r22, %r45;
	neg.s32 	%r47, %r46;
	not.b32 	%r48, %r8;
	add.u32 	%r8, %r47, %r48;
	.loc	19	1439	0
	xor.b32 	%r4, %r3, -2147483648;
$Lt_0_14850:
	.loc	19	1441	0
	mov.s32 	%r49, %r7;
	mov.u32 	%r50, 0;
	setp.le.s32 	%p8, %r8, %r50;
	@%p8 bra 	$Lt_0_17410;
	mov.u32 	%r51, 0;
$Lt_0_15874:
 //<loop> Loop body line 1441, nesting depth: 1, estimated iterations: unknown
	.loc	19	1445	0
	shr.u32 	%r52, %r22, 31;
	shl.b32 	%r53, %r8, 1;
	add.u32 	%r8, %r52, %r53;
	.loc	19	1446	0
	shl.b32 	%r22, %r22, 1;
	.loc	19	1447	0
	sub.u32 	%r51, %r51, 1;
	mov.u32 	%r54, 0;
	setp.gt.s32 	%p9, %r8, %r54;
	@%p9 bra 	$Lt_0_15874;
	bra.uni 	$Lt_0_15362;
$Lt_0_17410:
	mov.u32 	%r51, 0;
$Lt_0_15362:
	.loc	19	1449	0
	mul.lo.u32 	%r22, %r8, -921707870;
	.loc	19	1450	0
	mov.u32 	%r55, -921707870;
	mul.hi.u32 	%r8, %r8, %r55;
	mov.u32 	%r56, 0;
	setp.le.s32 	%p10, %r8, %r56;
	@%p10 bra 	$Lt_0_16386;
	.loc	19	1452	0
	shr.u32 	%r57, %r22, 31;
	shl.b32 	%r58, %r8, 1;
	add.u32 	%r8, %r57, %r58;
	.loc	19	1453	0
	shl.b32 	%r22, %r22, 1;
	.loc	19	1454	0
	sub.u32 	%r51, %r51, 1;
$Lt_0_16386:
	.loc	19	1456	0
	mov.u32 	%r59, 0;
	set.ne.u32.u32 	%r60, %r22, %r59;
	neg.s32 	%r61, %r60;
	add.u32 	%r8, %r61, %r8;
	.loc	19	1458	0
	shl.b32 	%r62, %r8, 24;
	mov.s32 	%r63, 0;
	set.lt.u32.s32 	%r64, %r62, %r63;
	neg.s32 	%r65, %r64;
	shr.u32 	%r66, %r8, 8;
	add.u32 	%r67, %r51, 126;
	shl.b32 	%r68, %r67, 23;
	add.u32 	%r69, %r66, %r68;
	add.u32 	%r70, %r65, %r69;
	or.b32 	%r71, %r4, %r70;
	mov.b32 	%f7, %r71;
	bra.uni 	$LDWendi___internal_fmad_177_2;
$Lt_0_12290:
	.loc	19	1173	0
	mov.f32 	%f8, 0f3f22f983;     	// 0.63662
	mul.f32 	%f9, %f1, %f8;
	cvt.rni.s32.f32 	%r72, %f9;
	cvt.rn.f32.s32 	%f10, %r72;
	neg.f32 	%f11, %f10;
	mov.f32 	%f12, 0f2e85a309;    	// 6.0771e-11
	mov.f32 	%f13, 0f39fdaa20;    	// 0.000483827
	mov.f32 	%f14, 0f3fc90000;    	// 1.57031
	fma.rn.f32 	%f15, %f11, %f14, %f1;
	fma.rn.f32 	%f16, %f11, %f13, %f15;
	fma.rn.f32 	%f17, %f11, %f12, %f16;
	.loc	19	1471	0
	mov.s32 	%r49, %r72;
	.loc	19	1472	0
	mov.f32 	%f7, %f17;
$LDWendi___internal_fmad_177_2:
	.loc	19	1705	0
	add.s32 	%r73, %r49, 1;
	mov.f32 	%f18, %f7;
	mul.f32 	%f19, %f18, %f18;
	and.b32 	%r74, %r73, 1;
	mov.u32 	%r75, 0;
	setp.eq.s32 	%p11, %r74, %r75;
	@%p11 bra 	$Lt_0_17154;
	.loc	19	1709	0
	mov.f32 	%f20, 0f37ccf5ce;    	// 2.44332e-05
	mov.f32 	%f21, 0fbab6061a;    	// -0.00138873
	fma.rn.f32 	%f22, %f20, %f19, %f21;
	mov.f32 	%f23, 0f3d2aaaa5;    	// 0.0416666
	fma.rn.f32 	%f24, %f22, %f19, %f23;
	mov.f32 	%f25, 0fbf000000;    	// -0.5
	fma.rn.f32 	%f26, %f24, %f19, %f25;
	mov.f32 	%f27, 0f3f800000;    	// 1
	fma.rn.f32 	%f28, %f26, %f19, %f27;
	bra.uni 	$Lt_0_16898;
$Lt_0_17154:
	.loc	19	1711	0
	mov.f32 	%f29, 0fb94ca1f9;    	// -0.000195153
	mov.f32 	%f30, 0f3c08839e;    	// 0.00833216
	fma.rn.f32 	%f31, %f29, %f19, %f30;
	mov.f32 	%f32, 0fbe2aaaa3;    	// -0.166667
	fma.rn.f32 	%f33, %f31, %f19, %f32;
	mul.f32 	%f34, %f19, %f33;
	fma.rn.f32 	%f28, %f34, %f18, %f18;
$Lt_0_16898:
	.loc	19	1713	0
	neg.f32 	%f35, %f28;
	and.b32 	%r76, %r73, 2;
	mov.s32 	%r77, 0;
	setp.ne.s32 	%p12, %r76, %r77;
	selp.f32 	%f28, %f35, %f28, %p12;
	.loc	19	1716	0
	mov.f32 	%f5, %f28;
$LDWendi___isinff_177_1:
	.loc	17	8	0
	ld.param.u64 	%rd7, [__cudaparm__Z17k_sequenceLoopingPfi_ptr];
	cvt.s64.s32 	%rd8, %r1;
	mul.wide.s32 	%rd9, %r1, 4;
	add.u64 	%rd10, %rd7, %rd9;
	st.global.f32 	[%rd10+0], %f5;
	.loc	17	10	0
	exit;
$LDWend__Z17k_sequenceLoopingPfi:
	} // _Z17k_sequenceLoopingPfi

	.entry _Z21k_matrixVectorProductPKfS0_Pfii (
		.param .u64 __cudaparm__Z21k_matrixVectorProductPKfS0_Pfii___val_paramA,
		.param .u64 __cudaparm__Z21k_matrixVectorProductPKfS0_Pfii___val_paramV,
		.param .u64 __cudaparm__Z21k_matrixVectorProductPKfS0_Pfii_R,
		.param .s32 __cudaparm__Z21k_matrixVectorProductPKfS0_Pfii_M,
		.param .s32 __cudaparm__Z21k_matrixVectorProductPKfS0_Pfii_N)
	{
	.reg .u32 %r<15>;
	.reg .u64 %rd<22>;
	.reg .f32 %f<6>;
	.reg .pred %p<6>;
	.shared .align 4 .b8 __cuda___cuda_local_var_27983_32_non_const_V_buffer72[32];
	.loc	17	21	0
$LDWbegin__Z21k_matrixVectorProductPKfS0_Pfii:
	.loc	17	26	0
	mov.u32 	%r1, %ntid.x;
	mov.u32 	%r2, %ctaid.x;
	mul.lo.u32 	%r3, %r2, %r1;
	mov.u32 	%r4, %tid.x;
	add.u32 	%r5, %r3, %r4;
	cvt.u64.u32 	%rd1, %r5;
	mul.wide.u32 	%rd2, %r5, 4;
	ld.param.u64 	%rd3, [__cudaparm__Z21k_matrixVectorProductPKfS0_Pfii___val_paramA];
	add.u64 	%rd4, %rd3, %rd2;
	.loc	17	27	0
	cvt.u64.u32 	%rd5, %r4;
	mul.wide.u32 	%rd6, %r4, 4;
	ld.param.u64 	%rd7, [__cudaparm__Z21k_matrixVectorProductPKfS0_Pfii___val_paramV];
	add.u64 	%rd8, %rd7, %rd6;
	ld.param.s32 	%r6, [__cudaparm__Z21k_matrixVectorProductPKfS0_Pfii_N];
	mov.u32 	%r7, 0;
	setp.le.s32 	%p1, %r6, %r7;
	@%p1 bra 	$Lt_1_4354;
	mov.u64 	%rd9, __cuda___cuda_local_var_27983_32_non_const_V_buffer72;
	mov.u32 	%r8, 0;
	setp.gt.u32 	%p2, %r1, %r8;
	cvt.u64.u32 	%rd10, %r1;
	add.u64 	%rd11, %rd6, %rd9;
	mov.s32 	%r9, 0;
	mov.f32 	%f1, 0f00000000;     	// 0
$Lt_1_2818:
 //<loop> Loop body line 27, nesting depth: 1, estimated iterations: unknown
	.loc	17	32	0
	ld.global.f32 	%f2, [%rd8+0];
	st.shared.f32 	[%rd11+0], %f2;
	.loc	17	33	0
	bar.sync 	0;
	@!%p2 bra 	$Lt_1_3074;
 //<loop> Part of loop body line 27, head labeled $Lt_1_2818
	mov.s32 	%r10, %r1;
	mov.s64 	%rd12, %rd9;
	ld.param.s32 	%r11, [__cudaparm__Z21k_matrixVectorProductPKfS0_Pfii_M];
	cvt.s64.s32 	%rd13, %r11;
	mul.wide.s32 	%rd14, %r11, 4;
	mov.s64 	%rd15, %rd4;
	mov.s32 	%r12, 0;
 //<loop> Part of loop body line 27, head labeled $Lt_1_2818
	mov.s32 	%r13, %r10;
$Lt_1_3586:
 //<loop> Loop body line 33, nesting depth: 2, estimated iterations: unknown
	.loc	17	37	0
	ld.shared.f32 	%f3, [%rd12+0];
	ld.global.f32 	%f4, [%rd15+0];
	fma.rn.f32 	%f1, %f3, %f4, %f1;
	add.s32 	%r12, %r12, 1;
	add.u64 	%rd15, %rd14, %rd15;
	add.u64 	%rd12, %rd12, 4;
	setp.ne.u32 	%p3, %r12, %r1;
	@%p3 bra 	$Lt_1_3586;
 //<loop> Part of loop body line 27, head labeled $Lt_1_2818
	mul.lo.u64 	%rd16, %rd13, %rd10;
	mul.lo.u64 	%rd17, %rd16, 4;
	add.u64 	%rd4, %rd4, %rd17;
$Lt_1_3074:
 //<loop> Part of loop body line 27, head labeled $Lt_1_2818
	.loc	17	40	0
	mul.lo.u64 	%rd18, %rd10, 4;
	add.u64 	%rd8, %rd8, %rd18;
	.loc	17	42	0
	bar.sync 	0;
	add.u32 	%r9, %r9, %r1;
	setp.gt.s32 	%p4, %r6, %r9;
	@%p4 bra 	$Lt_1_2818;
	bra.uni 	$Lt_1_2306;
$Lt_1_4354:
	mov.f32 	%f1, 0f00000000;     	// 0
$Lt_1_2306:
	.loc	17	45	0
	ld.param.u64 	%rd19, [__cudaparm__Z21k_matrixVectorProductPKfS0_Pfii_R];
	add.u64 	%rd20, %rd19, %rd2;
	st.global.f32 	[%rd20+0], %f1;
	.loc	17	46	0
	exit;
$LDWend__Z21k_matrixVectorProductPKfS0_Pfii:
	} // _Z21k_matrixVectorProductPKfS0_Pfii

