.version 1.3
	.target sm_10, map_f64_to_f32
	// compiled with /usr/local/cuda/open64/lib//be
	// nvopencc built on 2008-12-03

	.reg .u32 %ra<17>;
	.reg .u64 %rda<17>;
	.reg .f32 %fa<17>;
	.reg .f64 %fda<17>;
	.reg .u32 %rv<5>;
	.reg .u64 %rdv<5>;
	.reg .f32 %fv<5>;
	.reg .f64 %fdv<5>;


	//-----------------------------------------------------------
	// Compiling /tmp/tmpxft_00006f95_00000000-7_balloon-kernel.cpp3.i (/tmp/ccBI#.38KdMm)
	//-----------------------------------------------------------

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

	.file	1	"/tmp/tmpxft_00006f95_00000000-6_balloon-kernel.cudafe2.gpu"
	.file	2	"balloon.cuh"
	.file	3	"balloon-kernel.cu"
	.file	4	"/usr/lib/gcc/i386-redhat-linux/4.1.2/include/stddef.h"
	.file	5	"/usr/local/cuda/bin/../include/crt/device_runtime.h"
	.file	6	"/usr/local/cuda/bin/../include/host_defines.h"
	.file	7	"/usr/local/cuda/bin/../include/builtin_types.h"
	.file	8	"/usr/local/cuda/bin/../include/device_types.h"
	.file	9	"/usr/local/cuda/bin/../include/driver_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/include/stdint.h"
	.file	17	"/usr/local/cuda/bin/../include/common_functions.h"
	.file	18	"/usr/local/cuda/bin/../include/crt/func_macro.h"
	.file	19	"/usr/local/cuda/bin/../include/math_functions.h"
	.file	20	"/usr/local/cuda/bin/../include/device_functions.h"
	.file	21	"/usr/local/cuda/bin/../include/math_constants.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/texture_fetch_functions.h"
	.file	26	"/usr/local/cuda/bin/../include/math_functions_dbl_ptx1.h"

	.const .u32 tri_strc;
	.const .u32 vtx_strc;
	.const .u32 tri_data;
	.const .u32 tower_volumes;
	.const .u32 centroid_parts;
	.tex .u32 vtx_data_tex;
	.tex .u32 tri_data_tex;
	.const .u32 tri_work_strc;
	.const .s32 tri_work_per_vtx;
	.const .s32 tri_work_per_vtx_lg;
	.const .f32 volume_cpu;
	.const .s32 tri_count;
	.const .s32 point_count;
	.const .s8 opt_gravity;
	.const .f32 spring_constant;
	.const .f32 damping_v;
	.const .f32 pressure_factor_coeff;
	.const .f32 gas_m_over_temp;
	.const .f32 air_resistance;
	.const .f32 gas_mass_per_vertex;
	.const .f32 air_particle_mass;
	.const .f32 gravity_mag;
	.const .f32 delta_t;
	.const .f32 rep_constant;
	.const .f32 point_mass;
	.const .f32 point_mass_inv;
	.const .f32 platform_xmin;
	.const .f32 platform_xmax;
	.const .f32 platform_zmin;
	.const .f32 platform_zmax;

	.entry _Z14pass_trianglesv
	{
	.reg .u16 %rh<4>;
	.reg .u32 %r<80>;
	.reg .f32 %f<198>;
	.reg .pred %p<9>;
	.shared .align 4 .b8 __cuda_volumes0[256];
	.loc	3	247	0
$LBB1__Z14pass_trianglesv:
	mov.u16 	%rh1, %ctaid.x;      	// 
	mov.u16 	%rh2, %ntid.x;       	// 
	mul.wide.u16 	%r1, %rh2, %rh1;	// 
	cvt.s32.u16 	%r2, %tid.x;     	// 
	add.u32 	%r3, %r2, %r1;       	// 
	ld.const.s32 	%r4, [tri_count];	// id:2274 tri_count+0x0
	setp.le.s32 	%p1, %r4, %r3;   	// 
	@!%p1 bra 	$Lt_0_63;          	// 
	.loc	3	252	0
	mov.u32 	%r5, __cuda_volumes0;	// 
	mov.f32 	%f1, 0f00000000;     	// 0
	mul24.lo.u32 	%r6, %r2, 4;    	// 
	add.u32 	%r7, %r5, %r6;       	// 
	st.shared.f32 	[%r7+0], %f1;  	// id:2275 __cuda_volumes0+0x0
$Lt_0_63:
	mov.u32 	%r5, __cuda_volumes0;	// 
	.loc	3	253	0
	bar.sync 	0;                  	// 
	@!%p1 bra 	$Lt_0_65;          	// 
	bra.uni 	$LBB18__Z14pass_trianglesv;	// 
$Lt_0_65:
	.loc	3	256	0
	mul.lo.u32 	%r8, %r3, 16;     	// 
	ld.const.u32 	%r9, [tri_strc];	// id:2276 tri_strc+0x0
	add.u32 	%r10, %r9, %r8;      	// 
	ld.global.v4.s16 	{%r11,%r12,%r13,%r14}, [%r10+0];	// 
	ld.global.v2.s16 	{%r15,%r16}, [%r10+8];	// 
	ld.global.f32 	%f2, [%r10+12];	// id:2283
	.loc	3	258	0
	mul.lo.s32 	%r17, %r11, 3;    	// 
	add.s32 	%r18, %r17, 2;       	// 
	mov.s32 	%r19, 0;             	// 
	mov.s32 	%r20, 0;             	// 
	mov.s32 	%r21, 0;             	// 
	tex.1d.v4.f32.s32 {%f3,%f4,%f5,%f6},[vtx_data_tex,{%r18,%r19,%r20,%r21}];
	.loc	3	171	0
	mov.f32 	%f7, %f3;            	// 
	mov.f32 	%f8, %f4;            	// 
	mov.f32 	%f9, %f5;            	// 
	.loc	3	259	0
	mul.lo.s32 	%r22, %r12, 3;    	// 
	add.s32 	%r23, %r22, 2;       	// 
	mov.s32 	%r24, 0;             	// 
	mov.s32 	%r25, 0;             	// 
	mov.s32 	%r26, 0;             	// 
	tex.1d.v4.f32.s32 {%f10,%f11,%f12,%f13},[vtx_data_tex,{%r23,%r24,%r25,%r26}];
	.loc	3	171	0
	mov.f32 	%f14, %f10;          	// 
	mov.f32 	%f15, %f11;          	// 
	mov.f32 	%f16, %f12;          	// 
	.loc	3	260	0
	mul.lo.s32 	%r27, %r13, 3;    	// 
	add.s32 	%r28, %r27, 2;       	// 
	mov.s32 	%r29, 0;             	// 
	mov.s32 	%r30, 0;             	// 
	mov.s32 	%r31, 0;             	// 
	tex.1d.v4.f32.s32 {%f17,%f18,%f19,%f20},[vtx_data_tex,{%r28,%r29,%r30,%r31}];
	.loc	3	171	0
	mov.f32 	%f21, %f17;          	// 
	mov.f32 	%f22, %f18;          	// 
	mov.f32 	%f23, %f19;          	// 
	.loc	3	220	0
	mul.lo.s32 	%r32, %r14, 3;    	// 
	add.s32 	%r33, %r32, 2;       	// 
	mov.s32 	%r34, 0;             	// 
	mov.s32 	%r35, 0;             	// 
	mov.s32 	%r36, 0;             	// 
	tex.1d.v4.f32.s32 {%f24,%f25,%f26,%f27},[vtx_data_tex,{%r33,%r34,%r35,%r36}];
	.loc	3	171	0
	mov.f32 	%f28, %f24;          	// 
	mov.f32 	%f29, %f25;          	// 
	mov.f32 	%f30, %f26;          	// 
	.loc	20	1328	0
	ld.const.f32 	%f31, [rep_constant];	// id:2284 rep_constant+0x0
	mov.f32 	%f32, %f31;          	// 
	sub.f32 	%f33, %f7, %f28;     	// 
	sub.f32 	%f34, %f8, %f29;     	// 
	sub.f32 	%f35, %f9, %f30;     	// 
	mul.f32 	%f36, %f33, %f33;    	// 
	mul.f32 	%f37, %f34, %f34;    	// 
	mul.f32 	%f38, %f35, %f35;    	// 
	add.f32 	%f39, %f36, %f37;    	// 
	add.f32 	%f40, %f38, %f39;    	// 
	mov.f32 	%f41, 0f3a83126f;    	// 0.001
	max.f32 	%f42, %f40, %f41;    	// 
	mov.f32 	%f43, %f42;          	// 
	.loc	20	1250	0
	abs.f32 	%f44, %f42;          	// 
	mov.f32 	%f45, 0f7e800000;    	// 8.50706e+37
	setp.gt.f32 	%p2, %f44, %f45; 	// 
	@!%p2 bra 	$Lt_0_67;          	// 
	.loc	20	1251	0
	mov.f32 	%f46, 0f3e800000;    	// 0.25
	mul.f32 	%f32, %f31, %f46;    	// 
	.loc	20	1252	0
	mov.f32 	%f47, 0f3e800000;    	// 0.25
	mul.f32 	%f43, %f42, %f47;    	// 
$Lt_0_67:
	.loc	3	268	0
	div.f32 	%f48, %f32, %f43;    	// 
	rsqrt.f32 	%f49, %f40;        	// 
	mul.f32 	%f50, %f33, %f49;    	// 
	mul.f32 	%f51, %f48, %f50;    	// 
	mul.f32 	%f52, %f34, %f49;    	// 
	mul.f32 	%f53, %f48, %f52;    	// 
	mul.f32 	%f54, %f35, %f49;    	// 
	mul.f32 	%f55, %f48, %f54;    	// 
	.loc	3	220	0
	mul.lo.s32 	%r37, %r15, 3;    	// 
	add.s32 	%r38, %r37, 2;       	// 
	mov.s32 	%r39, 0;             	// 
	mov.s32 	%r40, 0;             	// 
	mov.s32 	%r41, 0;             	// 
	tex.1d.v4.f32.s32 {%f56,%f57,%f58,%f59},[vtx_data_tex,{%r38,%r39,%r40,%r41}];
	.loc	3	171	0
	mov.f32 	%f60, %f56;          	// 
	mov.f32 	%f61, %f57;          	// 
	mov.f32 	%f62, %f58;          	// 
	.loc	20	1328	0
	mov.f32 	%f32, %f31;          	// 
	sub.f32 	%f63, %f14, %f60;    	// 
	sub.f32 	%f64, %f15, %f61;    	// 
	sub.f32 	%f65, %f16, %f62;    	// 
	mul.f32 	%f66, %f63, %f63;    	// 
	mul.f32 	%f67, %f64, %f64;    	// 
	mul.f32 	%f68, %f65, %f65;    	// 
	add.f32 	%f69, %f66, %f67;    	// 
	add.f32 	%f70, %f68, %f69;    	// 
	mov.f32 	%f71, 0f3a83126f;    	// 0.001
	max.f32 	%f72, %f70, %f71;    	// 
	mov.f32 	%f43, %f72;          	// 
	.loc	20	1250	0
	abs.f32 	%f73, %f72;          	// 
	mov.f32 	%f74, 0f7e800000;    	// 8.50706e+37
	setp.gt.f32 	%p3, %f73, %f74; 	// 
	@!%p3 bra 	$Lt_0_69;          	// 
	.loc	20	1251	0
	mov.f32 	%f75, 0f3e800000;    	// 0.25
	mul.f32 	%f32, %f31, %f75;    	// 
	.loc	20	1252	0
	mov.f32 	%f76, 0f3e800000;    	// 0.25
	mul.f32 	%f43, %f72, %f76;    	// 
$Lt_0_69:
	.loc	3	269	0
	div.f32 	%f77, %f32, %f43;    	// 
	rsqrt.f32 	%f78, %f70;        	// 
	mul.f32 	%f79, %f63, %f78;    	// 
	mul.f32 	%f80, %f77, %f79;    	// 
	mul.f32 	%f81, %f64, %f78;    	// 
	mul.f32 	%f82, %f77, %f81;    	// 
	mul.f32 	%f83, %f65, %f78;    	// 
	mul.f32 	%f84, %f77, %f83;    	// 
	.loc	3	220	0
	mul.lo.s32 	%r42, %r16, 3;    	// 
	add.s32 	%r43, %r42, 2;       	// 
	mov.s32 	%r44, 0;             	// 
	mov.s32 	%r45, 0;             	// 
	mov.s32 	%r46, 0;             	// 
	tex.1d.v4.f32.s32 {%f85,%f86,%f87,%f88},[vtx_data_tex,{%r43,%r44,%r45,%r46}];
	.loc	3	171	0
	mov.f32 	%f89, %f85;          	// 
	mov.f32 	%f90, %f86;          	// 
	mov.f32 	%f91, %f87;          	// 
	.loc	20	1328	0
	mov.f32 	%f32, %f31;          	// 
	sub.f32 	%f92, %f21, %f89;    	// 
	sub.f32 	%f93, %f22, %f90;    	// 
	sub.f32 	%f94, %f23, %f91;    	// 
	mul.f32 	%f95, %f92, %f92;    	// 
	mul.f32 	%f96, %f93, %f93;    	// 
	mul.f32 	%f97, %f94, %f94;    	// 
	add.f32 	%f98, %f95, %f96;    	// 
	add.f32 	%f99, %f97, %f98;    	// 
	mov.f32 	%f100, 0f3a83126f;   	// 0.001
	max.f32 	%f101, %f99, %f100;  	// 
	mov.f32 	%f43, %f101;         	// 
	.loc	20	1250	0
	abs.f32 	%f102, %f101;        	// 
	mov.f32 	%f103, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p4, %f102, %f103;	// 
	@!%p4 bra 	$Lt_0_71;          	// 
	.loc	20	1251	0
	mov.f32 	%f104, 0f3e800000;   	// 0.25
	mul.f32 	%f32, %f31, %f104;   	// 
	.loc	20	1252	0
	mov.f32 	%f105, 0f3e800000;   	// 0.25
	mul.f32 	%f43, %f101, %f105;  	// 
$Lt_0_71:
	.loc	3	278	0
	sub.f32 	%f106, %f15, %f8;    	// 
	sub.f32 	%f107, %f23, %f9;    	// 
	sub.f32 	%f108, %f22, %f8;    	// 
	sub.f32 	%f109, %f16, %f9;    	// 
	mul.lo.u32 	%r47, %r3, 48;    	// 
	ld.const.u32 	%r48, [tri_data];	// id:2285 tri_data+0x0
	add.u32 	%r49, %r48, %r47;    	// 
	mul.f32 	%f110, %f108, %f109; 	// 
	mul.f32 	%f111, %f106, %f107; 	// 
	sub.f32 	%f112, %f111, %f110; 	// 
	sub.f32 	%f113, %f21, %f7;    	// 
	sub.f32 	%f114, %f14, %f7;    	// 
	mul.f32 	%f115, %f109, %f113; 	// 
	mul.f32 	%f116, %f107, %f114; 	// 
	sub.f32 	%f117, %f115, %f116; 	// 
	mul.f32 	%f118, %f106, %f113; 	// 
	mul.f32 	%f119, %f108, %f114; 	// 
	sub.f32 	%f120, %f119, %f118; 	// 
	.loc	3	281	0
	add.f32 	%f121, %f14, %f21;   	// 
	add.f32 	%f122, %f15, %f22;   	// 
	add.f32 	%f123, %f16, %f23;   	// 
	add.f32 	%f124, %f121, %f7;   	// 
	add.f32 	%f125, %f122, %f8;   	// 
	add.f32 	%f126, %f123, %f9;   	// 
	mov.f32 	%f127, 0f3eaaaaab;   	// 0.333333
	mul.f32 	%f128, %f124, %f127; 	// 
	mov.f32 	%f129, 0f3eaaaaab;   	// 0.333333
	mul.f32 	%f130, %f125, %f129; 	// 
	mov.f32 	%f131, 0f3eaaaaab;   	// 0.333333
	mul.f32 	%f132, %f126, %f131; 	// 
	sub.f32 	%f133, %f128, %f7;   	// 
	sub.f32 	%f134, %f128, %f14;  	// 
	sub.f32 	%f135, %f128, %f21;  	// 
	sub.f32 	%f136, %f130, %f8;   	// 
	sub.f32 	%f137, %f130, %f15;  	// 
	sub.f32 	%f138, %f130, %f22;  	// 
	sub.f32 	%f139, %f132, %f9;   	// 
	sub.f32 	%f140, %f132, %f16;  	// 
	sub.f32 	%f141, %f132, %f23;  	// 
	mul.f32 	%f142, %f136, %f136; 	// 
	mad.f32 	%f143, %f133, %f133, %f142;	// 
	mad.f32 	%f144, %f139, %f139, %f143;	// 
	sqrt.f32 	%f145, %f144;       	// 
	mul.f32 	%f146, %f137, %f137; 	// 
	mad.f32 	%f147, %f134, %f134, %f146;	// 
	mad.f32 	%f148, %f140, %f140, %f147;	// 
	sqrt.f32 	%f149, %f148;       	// 
	add.f32 	%f150, %f145, %f149; 	// 
	mul.f32 	%f151, %f138, %f138; 	// 
	mad.f32 	%f152, %f135, %f135, %f151;	// 
	mad.f32 	%f153, %f141, %f141, %f152;	// 
	sqrt.f32 	%f154, %f153;       	// 
	add.f32 	%f155, %f150, %f154; 	// 
	sub.f32 	%f156, %f155, %f2;   	// 
	mov.f32 	%f157, 0f00000000;   	// 0
	max.f32 	%f158, %f156, %f157; 	// 
	.loc	3	284	0
	ld.const.f32 	%f159, [spring_constant];	// id:2289 spring_constant+0x0
	mul.f32 	%f160, %f159, %f158; 	// 
	mad.f32 	%f161, %f160, %f133, %f51;	// 
	st.global.v4.f32 	[%r49+0], {%f112,%f117,%f120,%f161};	// 
	mad.f32 	%f162, %f160, %f136, %f53;	// 
	mad.f32 	%f163, %f160, %f139, %f55;	// 
	st.global.v2.f32 	[%r49+16], {%f162,%f163};	// 
	.loc	3	285	0
	mad.f32 	%f164, %f160, %f134, %f80;	// 
	st.global.f32 	[%r49+24], %f164;	// id:2293
	mad.f32 	%f165, %f160, %f137, %f82;	// 
	st.global.f32 	[%r49+28], %f165;	// id:2294
	mad.f32 	%f166, %f160, %f140, %f84;	// 
	st.global.f32 	[%r49+32], %f166;	// id:2295
	.loc	3	286	0
	div.f32 	%f167, %f32, %f43;   	// 
	rsqrt.f32 	%f168, %f99;       	// 
	mul.f32 	%f169, %f160, %f135; 	// 
	mul.f32 	%f170, %f92, %f168;  	// 
	mad.f32 	%f171, %f167, %f170, %f169;	// 
	st.global.f32 	[%r49+36], %f171;	// id:2296
	mul.f32 	%f172, %f160, %f138; 	// 
	mul.f32 	%f173, %f93, %f168;  	// 
	mad.f32 	%f174, %f167, %f173, %f172;	// 
	mul.f32 	%f175, %f160, %f141; 	// 
	mul.f32 	%f176, %f94, %f168;  	// 
	mad.f32 	%f177, %f167, %f176, %f175;	// 
	st.global.v2.f32 	[%r49+40], {%f174,%f177};	// 
	.loc	3	112	0
	mul24.lo.u32 	%r50, %r2, 4;   	// 
	add.u32 	%r51, %r50, %r5;     	// 
	mul.f32 	%f178, %f117, %f130; 	// 
	neg.f32 	%f179, %f178;        	// 
	mov.f32 	%f180, 0f3f000000;   	// 0.5
	mul.f32 	%f181, %f179, %f180; 	// 
	st.shared.f32 	[%r51+0], %f181;	// id:2299 __cuda_volumes0+0x0
	mov.f32 	%f182, %f181;        	// 
	.loc	3	113	0
	bar.sync 	0;                  	// 
	mov.u32 	%r52, 8;             	// 
	setp.ge.s32 	%p5, %r2, %r52;  	// 
	@%p5 bra 	$Lt_0_73;           	// 
	.loc	3	128	0
	add.s32 	%r53, %r2, 8;        	// 
	mul.lo.u32 	%r54, %r53, 4;    	// 
	add.u32 	%r55, %r5, %r54;     	// 
	ld.shared.f32 	%f183, [%r55+0];	// id:2300 __cuda_volumes0+0x0
	add.f32 	%f182, %f183, %f181; 	// 
	.loc	3	129	0
	add.s32 	%r56, %r2, 16;       	// 
	mul.lo.u32 	%r57, %r56, 4;    	// 
	add.u32 	%r58, %r5, %r57;     	// 
	ld.shared.f32 	%f184, [%r58+0];	// id:2301 __cuda_volumes0+0x0
	add.f32 	%f182, %f184, %f182; 	// 
	add.s32 	%r59, %r2, 24;       	// 
	mul.lo.u32 	%r60, %r59, 4;    	// 
	add.u32 	%r61, %r5, %r60;     	// 
	ld.shared.f32 	%f185, [%r61+0];	// id:2302 __cuda_volumes0+0x0
	add.f32 	%f182, %f185, %f182; 	// 
	.loc	3	130	0
	add.s32 	%r62, %r2, 32;       	// 
	mul.lo.u32 	%r63, %r62, 4;    	// 
	add.u32 	%r64, %r5, %r63;     	// 
	ld.shared.f32 	%f186, [%r64+0];	// id:2303 __cuda_volumes0+0x0
	add.f32 	%f182, %f186, %f182; 	// 
	add.s32 	%r65, %r2, 40;       	// 
	mul.lo.u32 	%r66, %r65, 4;    	// 
	add.u32 	%r67, %r5, %r66;     	// 
	ld.shared.f32 	%f187, [%r67+0];	// id:2304 __cuda_volumes0+0x0
	add.f32 	%f182, %f187, %f182; 	// 
	add.s32 	%r68, %r2, 48;       	// 
	mul.lo.u32 	%r69, %r68, 4;    	// 
	add.u32 	%r70, %r5, %r69;     	// 
	ld.shared.f32 	%f188, [%r70+0];	// id:2305 __cuda_volumes0+0x0
	add.f32 	%f182, %f188, %f182; 	// 
	add.s32 	%r71, %r2, 56;       	// 
	mul.lo.u32 	%r72, %r71, 4;    	// 
	add.u32 	%r73, %r5, %r72;     	// 
	ld.shared.f32 	%f189, [%r73+0];	// id:2306 __cuda_volumes0+0x0
	add.f32 	%f182, %f189, %f182; 	// 
	.loc	3	135	0
	st.shared.f32 	[%r51+0], %f182;	// id:2307 __cuda_volumes0+0x0
$Lt_0_73:
	.loc	3	140	0
	bar.sync 	0;                  	// 
	mov.u32 	%r74, 0;             	// 
	setp.ne.s32 	%p6, %r2, %r74;  	// 
	@%p6 bra 	$Lt_0_85;           	// 
	.loc	3	145	0
	ld.shared.f32 	%f190, [__cuda_volumes0+4];	// id:2308 __cuda_volumes0+0x4
	add.f32 	%f182, %f190, %f182; 	// 
	.loc	3	146	0
	ld.shared.f32 	%f191, [__cuda_volumes0+8];	// id:2309 __cuda_volumes0+0x8
	add.f32 	%f182, %f191, %f182; 	// 
	ld.shared.f32 	%f192, [__cuda_volumes0+12];	// id:2310 __cuda_volumes0+0xc
	add.f32 	%f182, %f192, %f182; 	// 
	.loc	3	147	0
	ld.shared.f32 	%f193, [__cuda_volumes0+16];	// id:2311 __cuda_volumes0+0x10
	add.f32 	%f182, %f193, %f182; 	// 
	ld.shared.f32 	%f194, [__cuda_volumes0+20];	// id:2312 __cuda_volumes0+0x14
	add.f32 	%f182, %f194, %f182; 	// 
	ld.shared.f32 	%f195, [__cuda_volumes0+24];	// id:2313 __cuda_volumes0+0x18
	add.f32 	%f182, %f195, %f182; 	// 
	ld.shared.f32 	%f196, [__cuda_volumes0+28];	// id:2314 __cuda_volumes0+0x1c
	add.f32 	%f182, %f196, %f182; 	// 
$Lt_0_85:
	mov.u32 	%r75, 0;             	// 
	setp.ne.u32 	%p7, %r2, %r75;  	// 
	@%p7 bra 	$LBB18__Z14pass_trianglesv;	// 
	.loc	3	290	0
	ld.const.u32 	%r76, [tower_volumes];	// id:2315 tower_volumes+0x0
	mul.wide.u16 	%r77, %rh1, 4;  	// 
	add.u32 	%r78, %r76, %r77;    	// 
	st.global.f32 	[%r78+0], %f182;	// id:2316
$LBB18__Z14pass_trianglesv:
	.loc	3	291	0
	exit;                         	// 
$LDWend__Z14pass_trianglesv:
	} // _Z14pass_trianglesv

	.entry _Z13pass_verticesP13CUDA_Vtx_Data
	{
	.reg .u16 %rh<6>;
	.reg .u32 %r<344>;
	.reg .f32 %f<430>;
	.reg .pred %p<51>;
	.param .u32 __cudaparm__Z13pass_verticesP13CUDA_Vtx_Data_vtx_data_out;
	.shared .align 4 .b8 __cuda_volumes260[256];
	.loc	3	306	0
$LBB1__Z13pass_verticesP13CUDA_Vtx_Data:
	.loc	3	318	0
	ld.const.s32 	%r1, [tri_count];	// id:5056 tri_count+0x0
	and.b32 	%r2, %r1, -64;       	// 
	shr.s32 	%r3, %r1, 6;         	// 
	setp.ne.s32 	%p1, %r2, %r1;   	// 
	selp.s32 	%r4, 1, 0, %p1;     	// 
	add.s32 	%r5, %r3, %r4;       	// 
	and.b32 	%r6, %r5, -64;       	// 
	shr.s32 	%r7, %r5, 6;         	// 
	setp.ne.s32 	%p2, %r5, %r6;   	// 
	selp.s32 	%r8, 1, 0, %p2;     	// 
	add.s32 	%r9, %r7, %r8;       	// 
	cvt.s32.u16 	%r10, %tid.x;    	// 
	mul.lo.s32 	%r11, %r9, %r10;  	// 
	add.s32 	%r12, %r9, %r11;     	// 
	min.s32 	%r13, %r5, %r12;     	// 
	setp.le.s32 	%p3, %r13, %r11; 	// 
	mov.f32 	%f1, 0f00000000;     	// 0
	@%p3 bra 	$Lt_1_249;          	// 
	sub.s32 	%r14, %r13, %r11;    	// 
	mul.lo.u32 	%r15, %r11, 4;    	// 
	ld.const.u32 	%r16, [tower_volumes];	// id:5017 tower_volumes+0x0
	add.u32 	%r17, %r15, %r16;    	// 
	mul.lo.u32 	%r18, %r13, 4;    	// 
	add.u32 	%r19, %r18, %r16;    	// 
	mov.s32 	%r20, %r14;          	// 
$Lt_1_185:
 //<loop> Loop body line 318, nesting depth: 1, estimated iterations: unknown
	ld.global.f32 	%f2, [%r17+0]; 	// id:5058
	add.f32 	%f1, %f2, %f1;       	// 
	add.u32 	%r17, %r17, 4;       	// 
	setp.ne.u32 	%p4, %r17, %r19; 	// 
	@%p4 bra 	$Lt_1_185;          	// 
	bra.uni 	$Lt_1_183;           	// 
$Lt_1_249:
$Lt_1_183:
	.loc	3	319	0
	mov.u32 	%r21, __cuda_volumes260;	// 
	.loc	3	112	0
	mul24.lo.u32 	%r22, %r10, 4;  	// 
	add.u32 	%r23, %r22, %r21;    	// 
	mov.f32 	%f3, %f1;            	// 
	st.shared.f32 	[%r23+0], %f3; 	// id:5059 __cuda_volumes260+0x0
	mov.f32 	%f4, %f3;            	// 
	.loc	3	113	0
	bar.sync 	0;                  	// 
	mov.u32 	%r24, 8;             	// 
	setp.ge.s32 	%p5, %r10, %r24; 	// 
	@%p5 bra 	$Lt_1_187;          	// 
	.loc	3	128	0
	add.s32 	%r25, %r10, 8;       	// 
	mul.lo.u32 	%r26, %r25, 4;    	// 
	add.u32 	%r27, %r21, %r26;    	// 
	ld.shared.f32 	%f5, [%r27+0]; 	// id:5060 __cuda_volumes260+0x0
	add.f32 	%f4, %f5, %f3;       	// 
	.loc	3	129	0
	add.s32 	%r28, %r10, 16;      	// 
	mul.lo.u32 	%r29, %r28, 4;    	// 
	add.u32 	%r30, %r21, %r29;    	// 
	ld.shared.f32 	%f6, [%r30+0]; 	// id:5061 __cuda_volumes260+0x0
	add.f32 	%f4, %f6, %f4;       	// 
	add.s32 	%r31, %r10, 24;      	// 
	mul.lo.u32 	%r32, %r31, 4;    	// 
	add.u32 	%r33, %r21, %r32;    	// 
	ld.shared.f32 	%f7, [%r33+0]; 	// id:5062 __cuda_volumes260+0x0
	add.f32 	%f4, %f7, %f4;       	// 
	.loc	3	130	0
	add.s32 	%r34, %r10, 32;      	// 
	mul.lo.u32 	%r35, %r34, 4;    	// 
	add.u32 	%r36, %r21, %r35;    	// 
	ld.shared.f32 	%f8, [%r36+0]; 	// id:5063 __cuda_volumes260+0x0
	add.f32 	%f4, %f8, %f4;       	// 
	add.s32 	%r37, %r10, 40;      	// 
	mul.lo.u32 	%r38, %r37, 4;    	// 
	add.u32 	%r39, %r21, %r38;    	// 
	ld.shared.f32 	%f9, [%r39+0]; 	// id:5064 __cuda_volumes260+0x0
	add.f32 	%f4, %f9, %f4;       	// 
	add.s32 	%r40, %r10, 48;      	// 
	mul.lo.u32 	%r41, %r40, 4;    	// 
	add.u32 	%r42, %r21, %r41;    	// 
	ld.shared.f32 	%f10, [%r42+0];	// id:5065 __cuda_volumes260+0x0
	add.f32 	%f4, %f10, %f4;      	// 
	add.s32 	%r43, %r10, 56;      	// 
	mul.lo.u32 	%r44, %r43, 4;    	// 
	add.u32 	%r45, %r21, %r44;    	// 
	ld.shared.f32 	%f11, [%r45+0];	// id:5066 __cuda_volumes260+0x0
	add.f32 	%f4, %f11, %f4;      	// 
	.loc	3	135	0
	st.shared.f32 	[%r23+0], %f4; 	// id:5067 __cuda_volumes260+0x0
$Lt_1_187:
	.loc	3	140	0
	bar.sync 	0;                  	// 
	mov.s32 	%r46, 0;             	// 
	setp.eq.s32 	%p6, %r10, %r46; 	// 
	@!%p6 bra 	$Lt_1_199;         	// 
	.loc	3	145	0
	ld.shared.f32 	%f12, [__cuda_volumes260+4];	// id:5068 __cuda_volumes260+0x4
	add.f32 	%f4, %f12, %f4;      	// 
	.loc	3	146	0
	ld.shared.f32 	%f13, [__cuda_volumes260+8];	// id:5069 __cuda_volumes260+0x8
	add.f32 	%f4, %f13, %f4;      	// 
	ld.shared.f32 	%f14, [__cuda_volumes260+12];	// id:5070 __cuda_volumes260+0xc
	add.f32 	%f4, %f14, %f4;      	// 
	.loc	3	147	0
	ld.shared.f32 	%f15, [__cuda_volumes260+16];	// id:5071 __cuda_volumes260+0x10
	add.f32 	%f4, %f15, %f4;      	// 
	ld.shared.f32 	%f16, [__cuda_volumes260+20];	// id:5072 __cuda_volumes260+0x14
	add.f32 	%f4, %f16, %f4;      	// 
	ld.shared.f32 	%f17, [__cuda_volumes260+24];	// id:5073 __cuda_volumes260+0x18
	add.f32 	%f4, %f17, %f4;      	// 
	ld.shared.f32 	%f18, [__cuda_volumes260+28];	// id:5074 __cuda_volumes260+0x1c
	add.f32 	%f4, %f18, %f4;      	// 
$Lt_1_199:
	@!%p6 bra 	$Lt_1_213;         	// 
	.loc	3	154	0
	st.shared.f32 	[__cuda_volumes260+0], %f4;	// id:5075 __cuda_volumes260+0x0
$Lt_1_213:
	.loc	3	155	0
	bar.sync 	0;                  	// 
	.loc	3	327	0
	mov.u16 	%rh1, %ctaid.x;      	// 
	mov.u16 	%rh2, %ntid.x;       	// 
	mul.wide.u16 	%r47, %rh1, %rh2;	// 
	add.u32 	%r48, %r10, %r47;    	// 
	mul.lo.s32 	%r49, %r48, 3;    	// 
	add.s32 	%r50, %r49, 2;       	// 
	mov.s32 	%r51, 0;             	// 
	mov.s32 	%r52, 0;             	// 
	mov.s32 	%r53, 0;             	// 
	tex.1d.v4.f32.s32 {%f19,%f20,%f21,%f22},[vtx_data_tex,{%r50,%r51,%r52,%r53}];
	.loc	3	171	0
	mov.f32 	%f23, %f19;          	// 
	mov.f32 	%f24, %f20;          	// 
	mov.f32 	%f25, %f21;          	// 
	.loc	3	328	0
	add.s32 	%r54, %r49, 1;       	// 
	mov.s32 	%r55, 0;             	// 
	mov.s32 	%r56, 0;             	// 
	mov.s32 	%r57, 0;             	// 
	tex.1d.v4.f32.s32 {%f26,%f27,%f28,%f29},[vtx_data_tex,{%r54,%r55,%r56,%r57}];
	.loc	3	178	0
	mov.f32 	%f30, %f26;          	// 
	mov.f32 	%f31, %f27;          	// 
	mov.f32 	%f32, %f28;          	// 
	.loc	3	329	0
	mul.lo.u32 	%r58, %r48, 16;   	// 
	ld.const.u32 	%r59, [vtx_strc];	// id:5078 vtx_strc+0x0
	add.u32 	%r60, %r59, %r58;    	// 
	ld.global.v4.u16 	{%r61,%r62,%r63,%r64}, [%r60+0];	// 
	ld.global.v4.u16 	{%r65,%r66,%r67,%r68}, [%r60+8];	// 
	mov.u32 	%r69, -1;            	// 
	setp.eq.s32 	%p7, %r61, %r69; 	// 
	@%p7 bra 	$Lt_1_216;          	// 
	.loc	3	347	0
	shr.s32 	%r70, %r61, 2;       	// 
	mul.lo.s32 	%r71, %r70, 3;    	// 
	mov.s32 	%r72, %r71;          	// 
	mov.s32 	%r73, 0;             	// 
	mov.s32 	%r74, 0;             	// 
	mov.s32 	%r75, 0;             	// 
	tex.1d.v4.f32.s32 {%f33,%f34,%f35,%f36},[tri_data_tex,{%r72,%r73,%r74,%r75}];
	.loc	3	185	0
	mov.f32 	%f37, %f33;          	// 
	mov.f32 	%f38, %f34;          	// 
	mov.f32 	%f39, %f35;          	// 
	.loc	3	67	0
	mov.f32 	%f40, %f37;          	// 
	mov.f32 	%f41, %f38;          	// 
	mov.f32 	%f42, %f39;          	// 
	.loc	3	193	0
	and.b32 	%r76, %r61, 3;       	// 
	mov.u32 	%r77, 0;             	// 
	setp.eq.s32 	%p8, %r76, %r77; 	// 
	@%p8 bra 	$Lt_1_58;           	// 
	mov.u32 	%r78, 1;             	// 
	setp.eq.s32 	%p9, %r76, %r78; 	// 
	@%p9 bra 	$Lt_1_59;           	// 
	mov.u32 	%r79, 2;             	// 
	setp.eq.s32 	%p10, %r76, %r79;	// 
	@%p10 bra 	$Lt_1_60;          	// 
	bra.uni 	$Lt_1_61;            	// 
$Lt_1_58:
	mov.s32 	%r80, %r71;          	// 
	mov.s32 	%r81, 0;             	// 
	mov.s32 	%r82, 0;             	// 
	mov.s32 	%r83, 0;             	// 
	tex.1d.v4.f32.s32 {%f43,%f44,%f45,%f46},[tri_data_tex,{%r80,%r81,%r82,%r83}];
	.loc	3	195	0
	mov.f32 	%f47, %f46;          	// 
	add.s32 	%r84, %r71, 1;       	// 
	mov.s32 	%r85, 0;             	// 
	mov.s32 	%r86, 0;             	// 
	mov.s32 	%r87, 0;             	// 
	tex.1d.v4.f32.s32 {%f48,%f49,%f50,%f51},[tri_data_tex,{%r84,%r85,%r86,%r87}];
	.loc	3	196	0
	mov.f32 	%f52, %f48;          	// 
	mov.f32 	%f53, %f49;          	// 
	mov.f32 	%f54, %f47;          	// 
	mov.f32 	%f55, %f52;          	// 
	mov.f32 	%f56, %f53;          	// 
	bra.uni 	$Lt_1_57;            	// 
$Lt_1_59:
	.loc	3	198	0
	add.s32 	%r88, %r71, 1;       	// 
	mov.s32 	%r89, 0;             	// 
	mov.s32 	%r90, 0;             	// 
	mov.s32 	%r91, 0;             	// 
	tex.1d.v4.f32.s32 {%f57,%f58,%f59,%f60},[tri_data_tex,{%r88,%r89,%r90,%r91}];
	.loc	3	199	0
	mov.f32 	%f61, %f59;          	// 
	mov.f32 	%f62, %f60;          	// 
	add.s32 	%r92, %r71, 2;       	// 
	mov.s32 	%r93, 0;             	// 
	mov.s32 	%r94, 0;             	// 
	mov.s32 	%r95, 0;             	// 
	tex.1d.v4.f32.s32 {%f63,%f64,%f65,%f66},[tri_data_tex,{%r92,%r93,%r94,%r95}];
	.loc	3	200	0
	mov.f32 	%f67, %f63;          	// 
	mov.f32 	%f54, %f61;          	// 
	mov.f32 	%f55, %f62;          	// 
	mov.f32 	%f56, %f67;          	// 
	bra.uni 	$Lt_1_57;            	// 
$Lt_1_60:
	.loc	3	202	0
	add.s32 	%r96, %r71, 2;       	// 
	mov.s32 	%r97, 0;             	// 
	mov.s32 	%r98, 0;             	// 
	mov.s32 	%r99, 0;             	// 
	tex.1d.v4.f32.s32 {%f68,%f69,%f70,%f71},[tri_data_tex,{%r96,%r97,%r98,%r99}];
	.loc	3	203	0
	mov.f32 	%f72, %f69;          	// 
	mov.f32 	%f73, %f70;          	// 
	mov.f32 	%f74, %f71;          	// 
	mov.f32 	%f54, %f72;          	// 
	mov.f32 	%f55, %f73;          	// 
	mov.f32 	%f56, %f74;          	// 
	bra.uni 	$Lt_1_57;            	// 
$Lt_1_61:
	.loc	3	205	0
	mov.f32 	%f56, 0f00000000;    	// 0
	mov.f32 	%f55, 0f00000000;    	// 0
	mov.f32 	%f54, 0f00000000;    	// 0
$Lt_1_57:
	.loc	3	67	0
	mov.f32 	%f75, %f54;          	// 
	mov.f32 	%f76, %f55;          	// 
	mov.f32 	%f77, %f56;          	// 
	.loc	3	347	0
	bra.uni 	$Lt_1_215;           	// 
$Lt_1_216:
	mov.f32 	%f42, 0f00000000;    	// 0
	mov.f32 	%f41, 0f00000000;    	// 0
	mov.f32 	%f40, 0f00000000;    	// 0
	mov.f32 	%f77, 0f00000000;    	// 0
	mov.f32 	%f76, 0f00000000;    	// 0
	mov.f32 	%f75, 0f00000000;    	// 0
$Lt_1_215:
	mov.u32 	%r100, -1;           	// 
	setp.eq.s32 	%p11, %r62, %r100;	// 
	@%p11 bra 	$Lt_1_217;         	// 
	shr.s32 	%r101, %r62, 2;      	// 
	mul.lo.s32 	%r102, %r101, 3;  	// 
	mov.s32 	%r103, %r102;        	// 
	mov.s32 	%r104, 0;            	// 
	mov.s32 	%r105, 0;            	// 
	mov.s32 	%r106, 0;            	// 
	tex.1d.v4.f32.s32 {%f78,%f79,%f80,%f81},[tri_data_tex,{%r103,%r104,%r105,%r106}];
	.loc	3	185	0
	mov.f32 	%f37, %f78;          	// 
	mov.f32 	%f38, %f79;          	// 
	mov.f32 	%f39, %f80;          	// 
	.loc	3	67	0
	add.f32 	%f40, %f37, %f40;    	// 
	add.f32 	%f41, %f38, %f41;    	// 
	add.f32 	%f42, %f39, %f42;    	// 
	.loc	3	193	0
	and.b32 	%r107, %r62, 3;      	// 
	mov.u32 	%r108, 0;            	// 
	setp.eq.s32 	%p12, %r107, %r108;	// 
	@%p12 bra 	$Lt_1_50;          	// 
	mov.u32 	%r109, 1;            	// 
	setp.eq.s32 	%p13, %r107, %r109;	// 
	@%p13 bra 	$Lt_1_51;          	// 
	mov.u32 	%r110, 2;            	// 
	setp.eq.s32 	%p14, %r107, %r110;	// 
	@%p14 bra 	$Lt_1_52;          	// 
	bra.uni 	$Lt_1_53;            	// 
$Lt_1_50:
	mov.s32 	%r111, %r102;        	// 
	mov.s32 	%r112, 0;            	// 
	mov.s32 	%r113, 0;            	// 
	mov.s32 	%r114, 0;            	// 
	tex.1d.v4.f32.s32 {%f82,%f83,%f84,%f85},[tri_data_tex,{%r111,%r112,%r113,%r114}];
	.loc	3	195	0
	mov.f32 	%f47, %f85;          	// 
	add.s32 	%r115, %r102, 1;     	// 
	mov.s32 	%r116, 0;            	// 
	mov.s32 	%r117, 0;            	// 
	mov.s32 	%r118, 0;            	// 
	tex.1d.v4.f32.s32 {%f86,%f87,%f88,%f89},[tri_data_tex,{%r115,%r116,%r117,%r118}];
	.loc	3	196	0
	mov.f32 	%f52, %f86;          	// 
	mov.f32 	%f53, %f87;          	// 
	mov.f32 	%f90, %f47;          	// 
	mov.f32 	%f91, %f52;          	// 
	mov.f32 	%f92, %f53;          	// 
	bra.uni 	$Lt_1_49;            	// 
$Lt_1_51:
	.loc	3	198	0
	add.s32 	%r119, %r102, 1;     	// 
	mov.s32 	%r120, 0;            	// 
	mov.s32 	%r121, 0;            	// 
	mov.s32 	%r122, 0;            	// 
	tex.1d.v4.f32.s32 {%f93,%f94,%f95,%f96},[tri_data_tex,{%r119,%r120,%r121,%r122}];
	.loc	3	199	0
	mov.f32 	%f61, %f95;          	// 
	mov.f32 	%f62, %f96;          	// 
	add.s32 	%r123, %r102, 2;     	// 
	mov.s32 	%r124, 0;            	// 
	mov.s32 	%r125, 0;            	// 
	mov.s32 	%r126, 0;            	// 
	tex.1d.v4.f32.s32 {%f97,%f98,%f99,%f100},[tri_data_tex,{%r123,%r124,%r125,%r126}];
	.loc	3	200	0
	mov.f32 	%f67, %f97;          	// 
	mov.f32 	%f90, %f61;          	// 
	mov.f32 	%f91, %f62;          	// 
	mov.f32 	%f92, %f67;          	// 
	bra.uni 	$Lt_1_49;            	// 
$Lt_1_52:
	.loc	3	202	0
	add.s32 	%r127, %r102, 2;     	// 
	mov.s32 	%r128, 0;            	// 
	mov.s32 	%r129, 0;            	// 
	mov.s32 	%r130, 0;            	// 
	tex.1d.v4.f32.s32 {%f101,%f102,%f103,%f104},[tri_data_tex,{%r127,%r128,%r129,%r130}];
	.loc	3	203	0
	mov.f32 	%f72, %f102;         	// 
	mov.f32 	%f73, %f103;         	// 
	mov.f32 	%f74, %f104;         	// 
	mov.f32 	%f90, %f72;          	// 
	mov.f32 	%f91, %f73;          	// 
	mov.f32 	%f92, %f74;          	// 
	bra.uni 	$Lt_1_49;            	// 
$Lt_1_53:
	.loc	3	205	0
	mov.f32 	%f92, 0f00000000;    	// 0
	mov.f32 	%f91, 0f00000000;    	// 0
	mov.f32 	%f90, 0f00000000;    	// 0
$Lt_1_49:
	.loc	3	67	0
	add.f32 	%f75, %f90, %f75;    	// 
	add.f32 	%f76, %f91, %f76;    	// 
	add.f32 	%f77, %f92, %f77;    	// 
$Lt_1_217:
	.loc	3	347	0
	mov.u32 	%r131, -1;           	// 
	setp.eq.s32 	%p15, %r63, %r131;	// 
	@%p15 bra 	$Lt_1_219;         	// 
	shr.s32 	%r132, %r63, 2;      	// 
	mul.lo.s32 	%r133, %r132, 3;  	// 
	mov.s32 	%r134, %r133;        	// 
	mov.s32 	%r135, 0;            	// 
	mov.s32 	%r136, 0;            	// 
	mov.s32 	%r137, 0;            	// 
	tex.1d.v4.f32.s32 {%f105,%f106,%f107,%f108},[tri_data_tex,{%r134,%r135,%r136,%r137}];
	.loc	3	185	0
	mov.f32 	%f37, %f105;         	// 
	mov.f32 	%f38, %f106;         	// 
	mov.f32 	%f39, %f107;         	// 
	.loc	3	67	0
	add.f32 	%f40, %f37, %f40;    	// 
	add.f32 	%f41, %f38, %f41;    	// 
	add.f32 	%f42, %f39, %f42;    	// 
	.loc	3	193	0
	and.b32 	%r138, %r63, 3;      	// 
	mov.u32 	%r139, 0;            	// 
	setp.eq.s32 	%p16, %r138, %r139;	// 
	@%p16 bra 	$Lt_1_42;          	// 
	mov.u32 	%r140, 1;            	// 
	setp.eq.s32 	%p17, %r138, %r140;	// 
	@%p17 bra 	$Lt_1_43;          	// 
	mov.u32 	%r141, 2;            	// 
	setp.eq.s32 	%p18, %r138, %r141;	// 
	@%p18 bra 	$Lt_1_44;          	// 
	bra.uni 	$Lt_1_45;            	// 
$Lt_1_42:
	mov.s32 	%r142, %r133;        	// 
	mov.s32 	%r143, 0;            	// 
	mov.s32 	%r144, 0;            	// 
	mov.s32 	%r145, 0;            	// 
	tex.1d.v4.f32.s32 {%f109,%f110,%f111,%f112},[tri_data_tex,{%r142,%r143,%r144,%r145}];
	.loc	3	195	0
	mov.f32 	%f47, %f112;         	// 
	add.s32 	%r146, %r133, 1;     	// 
	mov.s32 	%r147, 0;            	// 
	mov.s32 	%r148, 0;            	// 
	mov.s32 	%r149, 0;            	// 
	tex.1d.v4.f32.s32 {%f113,%f114,%f115,%f116},[tri_data_tex,{%r146,%r147,%r148,%r149}];
	.loc	3	196	0
	mov.f32 	%f52, %f113;         	// 
	mov.f32 	%f53, %f114;         	// 
	mov.f32 	%f117, %f47;         	// 
	mov.f32 	%f118, %f52;         	// 
	mov.f32 	%f119, %f53;         	// 
	bra.uni 	$Lt_1_41;            	// 
$Lt_1_43:
	.loc	3	198	0
	add.s32 	%r150, %r133, 1;     	// 
	mov.s32 	%r151, 0;            	// 
	mov.s32 	%r152, 0;            	// 
	mov.s32 	%r153, 0;            	// 
	tex.1d.v4.f32.s32 {%f120,%f121,%f122,%f123},[tri_data_tex,{%r150,%r151,%r152,%r153}];
	.loc	3	199	0
	mov.f32 	%f61, %f122;         	// 
	mov.f32 	%f62, %f123;         	// 
	add.s32 	%r154, %r133, 2;     	// 
	mov.s32 	%r155, 0;            	// 
	mov.s32 	%r156, 0;            	// 
	mov.s32 	%r157, 0;            	// 
	tex.1d.v4.f32.s32 {%f124,%f125,%f126,%f127},[tri_data_tex,{%r154,%r155,%r156,%r157}];
	.loc	3	200	0
	mov.f32 	%f67, %f124;         	// 
	mov.f32 	%f117, %f61;         	// 
	mov.f32 	%f118, %f62;         	// 
	mov.f32 	%f119, %f67;         	// 
	bra.uni 	$Lt_1_41;            	// 
$Lt_1_44:
	.loc	3	202	0
	add.s32 	%r158, %r133, 2;     	// 
	mov.s32 	%r159, 0;            	// 
	mov.s32 	%r160, 0;            	// 
	mov.s32 	%r161, 0;            	// 
	tex.1d.v4.f32.s32 {%f128,%f129,%f130,%f131},[tri_data_tex,{%r158,%r159,%r160,%r161}];
	.loc	3	203	0
	mov.f32 	%f72, %f129;         	// 
	mov.f32 	%f73, %f130;         	// 
	mov.f32 	%f74, %f131;         	// 
	mov.f32 	%f117, %f72;         	// 
	mov.f32 	%f118, %f73;         	// 
	mov.f32 	%f119, %f74;         	// 
	bra.uni 	$Lt_1_41;            	// 
$Lt_1_45:
	.loc	3	205	0
	mov.f32 	%f119, 0f00000000;   	// 0
	mov.f32 	%f118, 0f00000000;   	// 0
	mov.f32 	%f117, 0f00000000;   	// 0
$Lt_1_41:
	.loc	3	67	0
	add.f32 	%f75, %f117, %f75;   	// 
	add.f32 	%f76, %f118, %f76;   	// 
	add.f32 	%f77, %f119, %f77;   	// 
$Lt_1_219:
	.loc	3	347	0
	mov.u32 	%r162, -1;           	// 
	setp.eq.s32 	%p19, %r64, %r162;	// 
	@%p19 bra 	$Lt_1_221;         	// 
	shr.s32 	%r163, %r64, 2;      	// 
	mul.lo.s32 	%r164, %r163, 3;  	// 
	mov.s32 	%r165, %r164;        	// 
	mov.s32 	%r166, 0;            	// 
	mov.s32 	%r167, 0;            	// 
	mov.s32 	%r168, 0;            	// 
	tex.1d.v4.f32.s32 {%f132,%f133,%f134,%f135},[tri_data_tex,{%r165,%r166,%r167,%r168}];
	.loc	3	185	0
	mov.f32 	%f37, %f132;         	// 
	mov.f32 	%f38, %f133;         	// 
	mov.f32 	%f39, %f134;         	// 
	.loc	3	67	0
	add.f32 	%f40, %f37, %f40;    	// 
	add.f32 	%f41, %f38, %f41;    	// 
	add.f32 	%f42, %f39, %f42;    	// 
	.loc	3	193	0
	and.b32 	%r169, %r64, 3;      	// 
	mov.u32 	%r170, 0;            	// 
	setp.eq.s32 	%p20, %r169, %r170;	// 
	@%p20 bra 	$Lt_1_34;          	// 
	mov.u32 	%r171, 1;            	// 
	setp.eq.s32 	%p21, %r169, %r171;	// 
	@%p21 bra 	$Lt_1_35;          	// 
	mov.u32 	%r172, 2;            	// 
	setp.eq.s32 	%p22, %r169, %r172;	// 
	@%p22 bra 	$Lt_1_36;          	// 
	bra.uni 	$Lt_1_37;            	// 
$Lt_1_34:
	mov.s32 	%r173, %r164;        	// 
	mov.s32 	%r174, 0;            	// 
	mov.s32 	%r175, 0;            	// 
	mov.s32 	%r176, 0;            	// 
	tex.1d.v4.f32.s32 {%f136,%f137,%f138,%f139},[tri_data_tex,{%r173,%r174,%r175,%r176}];
	.loc	3	195	0
	mov.f32 	%f47, %f139;         	// 
	add.s32 	%r177, %r164, 1;     	// 
	mov.s32 	%r178, 0;            	// 
	mov.s32 	%r179, 0;            	// 
	mov.s32 	%r180, 0;            	// 
	tex.1d.v4.f32.s32 {%f140,%f141,%f142,%f143},[tri_data_tex,{%r177,%r178,%r179,%r180}];
	.loc	3	196	0
	mov.f32 	%f52, %f140;         	// 
	mov.f32 	%f53, %f141;         	// 
	mov.f32 	%f144, %f47;         	// 
	mov.f32 	%f145, %f52;         	// 
	mov.f32 	%f146, %f53;         	// 
	bra.uni 	$Lt_1_33;            	// 
$Lt_1_35:
	.loc	3	198	0
	add.s32 	%r181, %r164, 1;     	// 
	mov.s32 	%r182, 0;            	// 
	mov.s32 	%r183, 0;            	// 
	mov.s32 	%r184, 0;            	// 
	tex.1d.v4.f32.s32 {%f147,%f148,%f149,%f150},[tri_data_tex,{%r181,%r182,%r183,%r184}];
	.loc	3	199	0
	mov.f32 	%f61, %f149;         	// 
	mov.f32 	%f62, %f150;         	// 
	add.s32 	%r185, %r164, 2;     	// 
	mov.s32 	%r186, 0;            	// 
	mov.s32 	%r187, 0;            	// 
	mov.s32 	%r188, 0;            	// 
	tex.1d.v4.f32.s32 {%f151,%f152,%f153,%f154},[tri_data_tex,{%r185,%r186,%r187,%r188}];
	.loc	3	200	0
	mov.f32 	%f67, %f151;         	// 
	mov.f32 	%f144, %f61;         	// 
	mov.f32 	%f145, %f62;         	// 
	mov.f32 	%f146, %f67;         	// 
	bra.uni 	$Lt_1_33;            	// 
$Lt_1_36:
	.loc	3	202	0
	add.s32 	%r189, %r164, 2;     	// 
	mov.s32 	%r190, 0;            	// 
	mov.s32 	%r191, 0;            	// 
	mov.s32 	%r192, 0;            	// 
	tex.1d.v4.f32.s32 {%f155,%f156,%f157,%f158},[tri_data_tex,{%r189,%r190,%r191,%r192}];
	.loc	3	203	0
	mov.f32 	%f72, %f156;         	// 
	mov.f32 	%f73, %f157;         	// 
	mov.f32 	%f74, %f158;         	// 
	mov.f32 	%f144, %f72;         	// 
	mov.f32 	%f145, %f73;         	// 
	mov.f32 	%f146, %f74;         	// 
	bra.uni 	$Lt_1_33;            	// 
$Lt_1_37:
	.loc	3	205	0
	mov.f32 	%f146, 0f00000000;   	// 0
	mov.f32 	%f145, 0f00000000;   	// 0
	mov.f32 	%f144, 0f00000000;   	// 0
$Lt_1_33:
	.loc	3	67	0
	add.f32 	%f75, %f144, %f75;   	// 
	add.f32 	%f76, %f145, %f76;   	// 
	add.f32 	%f77, %f146, %f77;   	// 
$Lt_1_221:
	.loc	3	347	0
	mov.u32 	%r193, -1;           	// 
	setp.eq.s32 	%p23, %r65, %r193;	// 
	@%p23 bra 	$Lt_1_223;         	// 
	.loc	3	348	0
	shr.s32 	%r194, %r65, 2;      	// 
	mul.lo.s32 	%r195, %r194, 3;  	// 
	mov.s32 	%r196, %r195;        	// 
	mov.s32 	%r197, 0;            	// 
	mov.s32 	%r198, 0;            	// 
	mov.s32 	%r199, 0;            	// 
	tex.1d.v4.f32.s32 {%f159,%f160,%f161,%f162},[tri_data_tex,{%r196,%r197,%r198,%r199}];
	.loc	3	185	0
	mov.f32 	%f37, %f159;         	// 
	mov.f32 	%f38, %f160;         	// 
	mov.f32 	%f39, %f161;         	// 
	.loc	3	67	0
	add.f32 	%f40, %f37, %f40;    	// 
	add.f32 	%f41, %f38, %f41;    	// 
	add.f32 	%f42, %f39, %f42;    	// 
	.loc	3	193	0
	and.b32 	%r200, %r65, 3;      	// 
	mov.u32 	%r201, 0;            	// 
	setp.eq.s32 	%p24, %r200, %r201;	// 
	@%p24 bra 	$Lt_1_26;          	// 
	mov.u32 	%r202, 1;            	// 
	setp.eq.s32 	%p25, %r200, %r202;	// 
	@%p25 bra 	$Lt_1_27;          	// 
	mov.u32 	%r203, 2;            	// 
	setp.eq.s32 	%p26, %r200, %r203;	// 
	@%p26 bra 	$Lt_1_28;          	// 
	bra.uni 	$Lt_1_29;            	// 
$Lt_1_26:
	mov.s32 	%r204, %r195;        	// 
	mov.s32 	%r205, 0;            	// 
	mov.s32 	%r206, 0;            	// 
	mov.s32 	%r207, 0;            	// 
	tex.1d.v4.f32.s32 {%f163,%f164,%f165,%f166},[tri_data_tex,{%r204,%r205,%r206,%r207}];
	.loc	3	195	0
	mov.f32 	%f47, %f166;         	// 
	add.s32 	%r208, %r195, 1;     	// 
	mov.s32 	%r209, 0;            	// 
	mov.s32 	%r210, 0;            	// 
	mov.s32 	%r211, 0;            	// 
	tex.1d.v4.f32.s32 {%f167,%f168,%f169,%f170},[tri_data_tex,{%r208,%r209,%r210,%r211}];
	.loc	3	196	0
	mov.f32 	%f52, %f167;         	// 
	mov.f32 	%f53, %f168;         	// 
	mov.f32 	%f171, %f47;         	// 
	mov.f32 	%f172, %f52;         	// 
	mov.f32 	%f173, %f53;         	// 
	bra.uni 	$Lt_1_25;            	// 
$Lt_1_27:
	.loc	3	198	0
	add.s32 	%r212, %r195, 1;     	// 
	mov.s32 	%r213, 0;            	// 
	mov.s32 	%r214, 0;            	// 
	mov.s32 	%r215, 0;            	// 
	tex.1d.v4.f32.s32 {%f174,%f175,%f176,%f177},[tri_data_tex,{%r212,%r213,%r214,%r215}];
	.loc	3	199	0
	mov.f32 	%f61, %f176;         	// 
	mov.f32 	%f62, %f177;         	// 
	add.s32 	%r216, %r195, 2;     	// 
	mov.s32 	%r217, 0;            	// 
	mov.s32 	%r218, 0;            	// 
	mov.s32 	%r219, 0;            	// 
	tex.1d.v4.f32.s32 {%f178,%f179,%f180,%f181},[tri_data_tex,{%r216,%r217,%r218,%r219}];
	.loc	3	200	0
	mov.f32 	%f67, %f178;         	// 
	mov.f32 	%f171, %f61;         	// 
	mov.f32 	%f172, %f62;         	// 
	mov.f32 	%f173, %f67;         	// 
	bra.uni 	$Lt_1_25;            	// 
$Lt_1_28:
	.loc	3	202	0
	add.s32 	%r220, %r195, 2;     	// 
	mov.s32 	%r221, 0;            	// 
	mov.s32 	%r222, 0;            	// 
	mov.s32 	%r223, 0;            	// 
	tex.1d.v4.f32.s32 {%f182,%f183,%f184,%f185},[tri_data_tex,{%r220,%r221,%r222,%r223}];
	.loc	3	203	0
	mov.f32 	%f72, %f183;         	// 
	mov.f32 	%f73, %f184;         	// 
	mov.f32 	%f74, %f185;         	// 
	mov.f32 	%f171, %f72;         	// 
	mov.f32 	%f172, %f73;         	// 
	mov.f32 	%f173, %f74;         	// 
	bra.uni 	$Lt_1_25;            	// 
$Lt_1_29:
	.loc	3	205	0
	mov.f32 	%f173, 0f00000000;   	// 0
	mov.f32 	%f172, 0f00000000;   	// 0
	mov.f32 	%f171, 0f00000000;   	// 0
$Lt_1_25:
	.loc	3	67	0
	add.f32 	%f75, %f171, %f75;   	// 
	add.f32 	%f76, %f172, %f76;   	// 
	add.f32 	%f77, %f173, %f77;   	// 
$Lt_1_223:
	.loc	3	348	0
	mov.u32 	%r224, -1;           	// 
	setp.eq.s32 	%p27, %r66, %r224;	// 
	@%p27 bra 	$Lt_1_225;         	// 
	shr.s32 	%r225, %r66, 2;      	// 
	mul.lo.s32 	%r226, %r225, 3;  	// 
	mov.s32 	%r227, %r226;        	// 
	mov.s32 	%r228, 0;            	// 
	mov.s32 	%r229, 0;            	// 
	mov.s32 	%r230, 0;            	// 
	tex.1d.v4.f32.s32 {%f186,%f187,%f188,%f189},[tri_data_tex,{%r227,%r228,%r229,%r230}];
	.loc	3	185	0
	mov.f32 	%f37, %f186;         	// 
	mov.f32 	%f38, %f187;         	// 
	mov.f32 	%f39, %f188;         	// 
	.loc	3	67	0
	add.f32 	%f40, %f37, %f40;    	// 
	add.f32 	%f41, %f38, %f41;    	// 
	add.f32 	%f42, %f39, %f42;    	// 
	.loc	3	193	0
	and.b32 	%r231, %r66, 3;      	// 
	mov.u32 	%r232, 0;            	// 
	setp.eq.s32 	%p28, %r231, %r232;	// 
	@%p28 bra 	$Lt_1_18;          	// 
	mov.u32 	%r233, 1;            	// 
	setp.eq.s32 	%p29, %r231, %r233;	// 
	@%p29 bra 	$Lt_1_19;          	// 
	mov.u32 	%r234, 2;            	// 
	setp.eq.s32 	%p30, %r231, %r234;	// 
	@%p30 bra 	$Lt_1_20;          	// 
	bra.uni 	$Lt_1_21;            	// 
$Lt_1_18:
	mov.s32 	%r235, %r226;        	// 
	mov.s32 	%r236, 0;            	// 
	mov.s32 	%r237, 0;            	// 
	mov.s32 	%r238, 0;            	// 
	tex.1d.v4.f32.s32 {%f190,%f191,%f192,%f193},[tri_data_tex,{%r235,%r236,%r237,%r238}];
	.loc	3	195	0
	mov.f32 	%f47, %f193;         	// 
	add.s32 	%r239, %r226, 1;     	// 
	mov.s32 	%r240, 0;            	// 
	mov.s32 	%r241, 0;            	// 
	mov.s32 	%r242, 0;            	// 
	tex.1d.v4.f32.s32 {%f194,%f195,%f196,%f197},[tri_data_tex,{%r239,%r240,%r241,%r242}];
	.loc	3	196	0
	mov.f32 	%f52, %f194;         	// 
	mov.f32 	%f53, %f195;         	// 
	mov.f32 	%f198, %f47;         	// 
	mov.f32 	%f199, %f52;         	// 
	mov.f32 	%f200, %f53;         	// 
	bra.uni 	$Lt_1_17;            	// 
$Lt_1_19:
	.loc	3	198	0
	add.s32 	%r243, %r226, 1;     	// 
	mov.s32 	%r244, 0;            	// 
	mov.s32 	%r245, 0;            	// 
	mov.s32 	%r246, 0;            	// 
	tex.1d.v4.f32.s32 {%f201,%f202,%f203,%f204},[tri_data_tex,{%r243,%r244,%r245,%r246}];
	.loc	3	199	0
	mov.f32 	%f61, %f203;         	// 
	mov.f32 	%f62, %f204;         	// 
	add.s32 	%r247, %r226, 2;     	// 
	mov.s32 	%r248, 0;            	// 
	mov.s32 	%r249, 0;            	// 
	mov.s32 	%r250, 0;            	// 
	tex.1d.v4.f32.s32 {%f205,%f206,%f207,%f208},[tri_data_tex,{%r247,%r248,%r249,%r250}];
	.loc	3	200	0
	mov.f32 	%f67, %f205;         	// 
	mov.f32 	%f198, %f61;         	// 
	mov.f32 	%f199, %f62;         	// 
	mov.f32 	%f200, %f67;         	// 
	bra.uni 	$Lt_1_17;            	// 
$Lt_1_20:
	.loc	3	202	0
	add.s32 	%r251, %r226, 2;     	// 
	mov.s32 	%r252, 0;            	// 
	mov.s32 	%r253, 0;            	// 
	mov.s32 	%r254, 0;            	// 
	tex.1d.v4.f32.s32 {%f209,%f210,%f211,%f212},[tri_data_tex,{%r251,%r252,%r253,%r254}];
	.loc	3	203	0
	mov.f32 	%f72, %f210;         	// 
	mov.f32 	%f73, %f211;         	// 
	mov.f32 	%f74, %f212;         	// 
	mov.f32 	%f198, %f72;         	// 
	mov.f32 	%f199, %f73;         	// 
	mov.f32 	%f200, %f74;         	// 
	bra.uni 	$Lt_1_17;            	// 
$Lt_1_21:
	.loc	3	205	0
	mov.f32 	%f200, 0f00000000;   	// 0
	mov.f32 	%f199, 0f00000000;   	// 0
	mov.f32 	%f198, 0f00000000;   	// 0
$Lt_1_17:
	.loc	3	67	0
	add.f32 	%f75, %f198, %f75;   	// 
	add.f32 	%f76, %f199, %f76;   	// 
	add.f32 	%f77, %f200, %f77;   	// 
$Lt_1_225:
	.loc	3	348	0
	mov.u32 	%r255, -1;           	// 
	setp.eq.s32 	%p31, %r67, %r255;	// 
	@%p31 bra 	$Lt_1_227;         	// 
	shr.s32 	%r256, %r67, 2;      	// 
	mul.lo.s32 	%r257, %r256, 3;  	// 
	mov.s32 	%r258, %r257;        	// 
	mov.s32 	%r259, 0;            	// 
	mov.s32 	%r260, 0;            	// 
	mov.s32 	%r261, 0;            	// 
	tex.1d.v4.f32.s32 {%f213,%f214,%f215,%f216},[tri_data_tex,{%r258,%r259,%r260,%r261}];
	.loc	3	185	0
	mov.f32 	%f37, %f213;         	// 
	mov.f32 	%f38, %f214;         	// 
	mov.f32 	%f39, %f215;         	// 
	.loc	3	67	0
	add.f32 	%f40, %f37, %f40;    	// 
	add.f32 	%f41, %f38, %f41;    	// 
	add.f32 	%f42, %f39, %f42;    	// 
	.loc	3	193	0
	and.b32 	%r262, %r67, 3;      	// 
	mov.u32 	%r263, 0;            	// 
	setp.eq.s32 	%p32, %r262, %r263;	// 
	@%p32 bra 	$Lt_1_10;          	// 
	mov.u32 	%r264, 1;            	// 
	setp.eq.s32 	%p33, %r262, %r264;	// 
	@%p33 bra 	$Lt_1_11;          	// 
	mov.u32 	%r265, 2;            	// 
	setp.eq.s32 	%p34, %r262, %r265;	// 
	@%p34 bra 	$Lt_1_12;          	// 
	bra.uni 	$Lt_1_13;            	// 
$Lt_1_10:
	mov.s32 	%r266, %r257;        	// 
	mov.s32 	%r267, 0;            	// 
	mov.s32 	%r268, 0;            	// 
	mov.s32 	%r269, 0;            	// 
	tex.1d.v4.f32.s32 {%f217,%f218,%f219,%f220},[tri_data_tex,{%r266,%r267,%r268,%r269}];
	.loc	3	195	0
	mov.f32 	%f47, %f220;         	// 
	add.s32 	%r270, %r257, 1;     	// 
	mov.s32 	%r271, 0;            	// 
	mov.s32 	%r272, 0;            	// 
	mov.s32 	%r273, 0;            	// 
	tex.1d.v4.f32.s32 {%f221,%f222,%f223,%f224},[tri_data_tex,{%r270,%r271,%r272,%r273}];
	.loc	3	196	0
	mov.f32 	%f52, %f221;         	// 
	mov.f32 	%f53, %f222;         	// 
	mov.f32 	%f225, %f47;         	// 
	mov.f32 	%f226, %f52;         	// 
	mov.f32 	%f227, %f53;         	// 
	bra.uni 	$Lt_1_9;             	// 
$Lt_1_11:
	.loc	3	198	0
	add.s32 	%r274, %r257, 1;     	// 
	mov.s32 	%r275, 0;            	// 
	mov.s32 	%r276, 0;            	// 
	mov.s32 	%r277, 0;            	// 
	tex.1d.v4.f32.s32 {%f228,%f229,%f230,%f231},[tri_data_tex,{%r274,%r275,%r276,%r277}];
	.loc	3	199	0
	mov.f32 	%f61, %f230;         	// 
	mov.f32 	%f62, %f231;         	// 
	add.s32 	%r278, %r257, 2;     	// 
	mov.s32 	%r279, 0;            	// 
	mov.s32 	%r280, 0;            	// 
	mov.s32 	%r281, 0;            	// 
	tex.1d.v4.f32.s32 {%f232,%f233,%f234,%f235},[tri_data_tex,{%r278,%r279,%r280,%r281}];
	.loc	3	200	0
	mov.f32 	%f67, %f232;         	// 
	mov.f32 	%f225, %f61;         	// 
	mov.f32 	%f226, %f62;         	// 
	mov.f32 	%f227, %f67;         	// 
	bra.uni 	$Lt_1_9;             	// 
$Lt_1_12:
	.loc	3	202	0
	add.s32 	%r282, %r257, 2;     	// 
	mov.s32 	%r283, 0;            	// 
	mov.s32 	%r284, 0;            	// 
	mov.s32 	%r285, 0;            	// 
	tex.1d.v4.f32.s32 {%f236,%f237,%f238,%f239},[tri_data_tex,{%r282,%r283,%r284,%r285}];
	.loc	3	203	0
	mov.f32 	%f72, %f237;         	// 
	mov.f32 	%f73, %f238;         	// 
	mov.f32 	%f74, %f239;         	// 
	mov.f32 	%f225, %f72;         	// 
	mov.f32 	%f226, %f73;         	// 
	mov.f32 	%f227, %f74;         	// 
	bra.uni 	$Lt_1_9;             	// 
$Lt_1_13:
	.loc	3	205	0
	mov.f32 	%f227, 0f00000000;   	// 0
	mov.f32 	%f226, 0f00000000;   	// 0
	mov.f32 	%f225, 0f00000000;   	// 0
$Lt_1_9:
	.loc	3	67	0
	add.f32 	%f75, %f225, %f75;   	// 
	add.f32 	%f76, %f226, %f76;   	// 
	add.f32 	%f77, %f227, %f77;   	// 
$Lt_1_227:
	.loc	3	348	0
	mov.u32 	%r286, -1;           	// 
	setp.eq.s32 	%p35, %r68, %r286;	// 
	@%p35 bra 	$Lt_1_229;         	// 
	shr.s32 	%r287, %r68, 2;      	// 
	mul.lo.s32 	%r288, %r287, 3;  	// 
	mov.s32 	%r289, %r288;        	// 
	mov.s32 	%r290, 0;            	// 
	mov.s32 	%r291, 0;            	// 
	mov.s32 	%r292, 0;            	// 
	tex.1d.v4.f32.s32 {%f240,%f241,%f242,%f243},[tri_data_tex,{%r289,%r290,%r291,%r292}];
	.loc	3	185	0
	mov.f32 	%f37, %f240;         	// 
	mov.f32 	%f38, %f241;         	// 
	mov.f32 	%f39, %f242;         	// 
	.loc	3	67	0
	add.f32 	%f40, %f37, %f40;    	// 
	add.f32 	%f41, %f38, %f41;    	// 
	add.f32 	%f42, %f39, %f42;    	// 
	.loc	3	193	0
	and.b32 	%r293, %r68, 3;      	// 
	mov.u32 	%r294, 0;            	// 
	setp.eq.s32 	%p36, %r293, %r294;	// 
	@%p36 bra 	$Lt_1_2;           	// 
	mov.u32 	%r295, 1;            	// 
	setp.eq.s32 	%p37, %r293, %r295;	// 
	@%p37 bra 	$Lt_1_3;           	// 
	mov.u32 	%r296, 2;            	// 
	setp.eq.s32 	%p38, %r293, %r296;	// 
	@%p38 bra 	$Lt_1_4;           	// 
	bra.uni 	$Lt_1_5;             	// 
$Lt_1_2:
	mov.s32 	%r297, %r288;        	// 
	mov.s32 	%r298, 0;            	// 
	mov.s32 	%r299, 0;            	// 
	mov.s32 	%r300, 0;            	// 
	tex.1d.v4.f32.s32 {%f244,%f245,%f246,%f247},[tri_data_tex,{%r297,%r298,%r299,%r300}];
	.loc	3	195	0
	mov.f32 	%f47, %f247;         	// 
	add.s32 	%r301, %r288, 1;     	// 
	mov.s32 	%r302, 0;            	// 
	mov.s32 	%r303, 0;            	// 
	mov.s32 	%r304, 0;            	// 
	tex.1d.v4.f32.s32 {%f248,%f249,%f250,%f251},[tri_data_tex,{%r301,%r302,%r303,%r304}];
	.loc	3	196	0
	mov.f32 	%f52, %f248;         	// 
	mov.f32 	%f53, %f249;         	// 
	mov.f32 	%f252, %f47;         	// 
	mov.f32 	%f253, %f52;         	// 
	mov.f32 	%f254, %f53;         	// 
	bra.uni 	$Lt_1_1;             	// 
$Lt_1_3:
	.loc	3	198	0
	add.s32 	%r305, %r288, 1;     	// 
	mov.s32 	%r306, 0;            	// 
	mov.s32 	%r307, 0;            	// 
	mov.s32 	%r308, 0;            	// 
	tex.1d.v4.f32.s32 {%f255,%f256,%f257,%f258},[tri_data_tex,{%r305,%r306,%r307,%r308}];
	.loc	3	199	0
	mov.f32 	%f61, %f257;         	// 
	mov.f32 	%f62, %f258;         	// 
	add.s32 	%r309, %r288, 2;     	// 
	mov.s32 	%r310, 0;            	// 
	mov.s32 	%r311, 0;            	// 
	mov.s32 	%r312, 0;            	// 
	tex.1d.v4.f32.s32 {%f259,%f260,%f261,%f262},[tri_data_tex,{%r309,%r310,%r311,%r312}];
	.loc	3	200	0
	mov.f32 	%f67, %f259;         	// 
	mov.f32 	%f252, %f61;         	// 
	mov.f32 	%f253, %f62;         	// 
	mov.f32 	%f254, %f67;         	// 
	bra.uni 	$Lt_1_1;             	// 
$Lt_1_4:
	.loc	3	202	0
	add.s32 	%r313, %r288, 2;     	// 
	mov.s32 	%r314, 0;            	// 
	mov.s32 	%r315, 0;            	// 
	mov.s32 	%r316, 0;            	// 
	tex.1d.v4.f32.s32 {%f263,%f264,%f265,%f266},[tri_data_tex,{%r313,%r314,%r315,%r316}];
	.loc	3	203	0
	mov.f32 	%f72, %f264;         	// 
	mov.f32 	%f73, %f265;         	// 
	mov.f32 	%f74, %f266;         	// 
	mov.f32 	%f252, %f72;         	// 
	mov.f32 	%f253, %f73;         	// 
	mov.f32 	%f254, %f74;         	// 
	bra.uni 	$Lt_1_1;             	// 
$Lt_1_5:
	.loc	3	205	0
	mov.f32 	%f254, 0f00000000;   	// 0
	mov.f32 	%f253, 0f00000000;   	// 0
	mov.f32 	%f252, 0f00000000;   	// 0
$Lt_1_1:
	.loc	3	67	0
	add.f32 	%f75, %f252, %f75;   	// 
	add.f32 	%f76, %f253, %f76;   	// 
	add.f32 	%f77, %f254, %f77;   	// 
$Lt_1_229:
	.loc	20	1328	0
	ld.const.f32 	%f267, [pressure_factor_coeff];	// id:5087 pressure_factor_coeff+0x0
	ld.shared.f32 	%f268, [__cuda_volumes260+0];	// id:5075 __cuda_volumes260+0x0
	abs.f32 	%f269, %f268;        	// 
	mov.f32 	%f270, %f269;        	// 
	.loc	20	1250	0
	mov.f32 	%f271, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p39, %f269, %f271;	// 
	@!%p39 bra 	$Lt_1_231;        	// 
	.loc	20	1251	0
	mov.f32 	%f272, 0f3e800000;   	// 0.25
	mul.f32 	%f267, %f267, %f272; 	// 
	.loc	20	1252	0
	mov.f32 	%f273, 0f3e800000;   	// 0.25
	mul.f32 	%f270, %f269, %f273; 	// 
$Lt_1_231:
	.loc	3	356	0
	ld.const.s8 	%rh3, [opt_gravity];	// id:5088 opt_gravity+0x0
	mov.s16 	%rh4, 0;             	// 
	setp.ne.s16 	%p40, %rh3, %rh4;	// 
	div.f32 	%f274, %f267, %f270; 	// 
	@!%p40 bra 	$Lt_1_234;        	// 
	.loc	3	358	0
	ld.const.f32 	%f275, [gas_m_over_temp];	// id:5089 gas_m_over_temp+0x0
	mul.f32 	%f276, %f275, %f24;  	// 
	neg.f32 	%f277, %f276;        	// 
	mov.f32 	%f278, 0f3fb8aa3b;   	// 1.4427
	mul.f32 	%f279, %f277, %f278; 	// 
	cvt.rzi.f32.f32 	%f280, %f279;	// 
	mov.f32 	%f281, 0f7f800000;   	// ((1.0F)/(0.0F))
	mov.f32 	%f282, 0f00000000;   	// 0
	ex2.f32 	%f283, %f280;        	// 
	mov.f32 	%f284, 0f3f317200;   	// 0.693146
	mad.f32 	%f285, %f280, %f284, %f276;	// 
	mov.f32 	%f286, 0f35bfbe8e;   	// 1.42861e-06
	mad.f32 	%f287, %f280, %f286, %f285;	// 
	neg.f32 	%f288, %f287;        	// 
	mov.f32 	%f289, 0f3fb8aa3b;   	// 1.4427
	mul.f32 	%f290, %f288, %f289; 	// 
	ex2.f32 	%f291, %f290;        	// 
	mul.f32 	%f292, %f283, %f291; 	// 
	mov.f32 	%f293, 0fc2d20000;   	// -105
	setp.lt.f32 	%p41, %f277, %f293;	// 
	selp.f32 	%f294, %f282, %f292, %p41;	// 
	mov.f32 	%f295, 0f42d20000;   	// 105
	setp.gt.f32 	%p42, %f277, %f295;	// 
	selp.f32 	%f296, %f281, %f294, %p42;	// 
	mul.f32 	%f297, %f296, %f274; 	// 
	bra.uni 	$Lt_1_233;           	// 
$Lt_1_234:
	mov.f32 	%f297, %f274;        	// 
$Lt_1_233:
	@!%p40 bra 	$Lt_1_236;        	// 
	.loc	3	362	0
	ld.const.f32 	%f298, [air_particle_mass];	// id:5090 air_particle_mass+0x0
	mov.f32 	%f299, 0fbe4ccccd;   	// -0.2
	mul.f32 	%f300, %f298, %f299; 	// 
	mul.f32 	%f301, %f300, %f24;  	// 
	mov.f32 	%f302, 0f3fb8aa3b;   	// 1.4427
	mul.f32 	%f303, %f301, %f302; 	// 
	cvt.rzi.f32.f32 	%f304, %f303;	// 
	mov.f32 	%f305, 0f7f800000;   	// ((1.0F)/(0.0F))
	mov.f32 	%f306, 0f00000000;   	// 0
	ex2.f32 	%f307, %f304;        	// 
	mov.f32 	%f308, 0f3f317200;   	// 0.693146
	mul.f32 	%f309, %f304, %f308; 	// 
	sub.f32 	%f310, %f301, %f309; 	// 
	mov.f32 	%f311, 0f35bfbe8e;   	// 1.42861e-06
	mul.f32 	%f312, %f304, %f311; 	// 
	sub.f32 	%f313, %f310, %f312; 	// 
	mov.f32 	%f314, 0f3fb8aa3b;   	// 1.4427
	mul.f32 	%f315, %f313, %f314; 	// 
	ex2.f32 	%f316, %f315;        	// 
	mul.f32 	%f317, %f307, %f316; 	// 
	mov.f32 	%f318, 0fc2d20000;   	// -105
	setp.lt.f32 	%p43, %f301, %f318;	// 
	selp.f32 	%f319, %f306, %f317, %p43;	// 
	mov.f32 	%f320, 0f42d20000;   	// 105
	setp.gt.f32 	%p44, %f301, %f320;	// 
	selp.f32 	%f321, %f305, %f319, %p44;	// 
	bra.uni 	$Lt_1_235;           	// 
$Lt_1_236:
	mov.f32 	%f321, 0f3f800000;   	// 1
$Lt_1_235:
	.loc	3	371	0
	mul.f32 	%f322, %f30, %f30;   	// 
	mul.f32 	%f323, %f31, %f31;   	// 
	mul.f32 	%f324, %f32, %f32;   	// 
	mov.f32 	%f325, 0f3e2aaaab;   	// 0.166667
	mul.f32 	%f326, %f40, %f325;  	// 
	mov.f32 	%f327, 0f3e2aaaab;   	// 0.166667
	mul.f32 	%f328, %f41, %f327;  	// 
	mov.f32 	%f329, 0f3e2aaaab;   	// 0.166667
	mul.f32 	%f330, %f42, %f329;  	// 
	add.f32 	%f331, %f322, %f323; 	// 
	add.f32 	%f332, %f324, %f331; 	// 
	rsqrt.f32 	%f333, %f332;      	// 
	ld.const.f32 	%f334, [air_resistance];	// id:5091 air_resistance+0x0
	mul.f32 	%f335, %f333, %f31;  	// 
	mul.f32 	%f336, %f328, %f335; 	// 
	mul.f32 	%f337, %f333, %f30;  	// 
	mad.f32 	%f338, %f337, %f326, %f336;	// 
	mul.f32 	%f339, %f333, %f32;  	// 
	mad.f32 	%f340, %f339, %f330, %f338;	// 
	neg.f32 	%f341, %f340;        	// 
	mov.f32 	%f342, 0f00000000;   	// 0
	max.f32 	%f343, %f341, %f342; 	// 
	mul.f32 	%f344, %f334, %f343; 	// 
	neg.f32 	%f345, %f344;        	// 
	.loc	20	1328	0
	ld.const.f32 	%f346, [delta_t];	// id:5092 delta_t+0x0
	mov.f32 	%f267, %f346;        	// 
	ld.const.f32 	%f347, [point_mass];	// id:5093 point_mass+0x0
	ld.const.f32 	%f348, [gas_mass_per_vertex];	// id:5094 gas_mass_per_vertex+0x0
	add.f32 	%f349, %f348, %f347; 	// 
	mov.f32 	%f270, %f349;        	// 
	.loc	20	1250	0
	abs.f32 	%f350, %f349;        	// 
	mov.f32 	%f351, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p45, %f350, %f351;	// 
	@!%p45 bra 	$Lt_1_237;        	// 
	.loc	20	1251	0
	mov.f32 	%f352, 0f3e800000;   	// 0.25
	mul.f32 	%f267, %f346, %f352; 	// 
	.loc	20	1252	0
	mov.f32 	%f353, 0f3e800000;   	// 0.25
	mul.f32 	%f270, %f349, %f353; 	// 
$Lt_1_237:
	.loc	3	389	0
	div.f32 	%f354, %f267, %f270; 	// 
	mul.f32 	%f355, %f30, %f345;  	// 
	sub.f32 	%f356, %f321, %f297; 	// 
	mul.f32 	%f357, %f354, %f75;  	// 
	mul.f32 	%f358, %f326, %f356; 	// 
	add.f32 	%f359, %f355, %f358; 	// 
	mul.f32 	%f360, %f354, %f359; 	// 
	add.f32 	%f361, %f357, %f360; 	// 
	mov.f32 	%f362, 0f3f000000;   	// 0.5
	mul.f32 	%f363, %f361, %f362; 	// 
	add.f32 	%f364, %f363, %f30;  	// 
	mul.f32 	%f365, %f364, %f346; 	// 
	add.f32 	%f366, %f365, %f23;  	// 
	mul.f32 	%f367, %f345, %f31;  	// 
	ld.const.f32 	%f368, [gravity_mag];	// id:5095 gravity_mag+0x0
	mul.f32 	%f369, %f368, %f347; 	// 
	mul.f32 	%f370, %f354, %f76;  	// 
	mul.f32 	%f371, %f328, %f356; 	// 
	sub.f32 	%f372, %f371, %f369; 	// 
	add.f32 	%f373, %f367, %f372; 	// 
	mul.f32 	%f374, %f354, %f373; 	// 
	add.f32 	%f375, %f370, %f374; 	// 
	mov.f32 	%f376, 0f3f000000;   	// 0.5
	mul.f32 	%f377, %f375, %f376; 	// 
	add.f32 	%f378, %f377, %f31;  	// 
	mul.f32 	%f379, %f378, %f346; 	// 
	add.f32 	%f380, %f379, %f24;  	// 
	mov.f32 	%f381, %f380;        	// 
	mul.f32 	%f382, %f345, %f32;  	// 
	mul.f32 	%f383, %f354, %f77;  	// 
	mul.f32 	%f384, %f330, %f356; 	// 
	add.f32 	%f385, %f382, %f384; 	// 
	mul.f32 	%f386, %f354, %f385; 	// 
	add.f32 	%f387, %f383, %f386; 	// 
	mov.f32 	%f388, 0f3f000000;   	// 0.5
	mul.f32 	%f389, %f387, %f388; 	// 
	add.f32 	%f390, %f389, %f32;  	// 
	mul.f32 	%f391, %f390, %f346; 	// 
	add.f32 	%f392, %f391, %f25;  	// 
	.loc	3	392	0
	ld.const.f32 	%f393, [damping_v];	// id:5046 damping_v+0x0
	mad.f32 	%f394, %f357, %f393, %f360;	// 
	add.f32 	%f395, %f30, %f394;  	// 
	mad.f32 	%f396, %f370, %f393, %f374;	// 
	add.f32 	%f397, %f31, %f396;  	// 
	mad.f32 	%f398, %f383, %f393, %f386;	// 
	add.f32 	%f399, %f32, %f398;  	// 
	.loc	3	394	0
	ld.const.f32 	%f400, [platform_xmax];	// id:5098 platform_xmax+0x0
	set.ge.u32.f32 	%r317, %f400, %f366;	// 
	neg.s32 	%r318, %r317;        	// 
	ld.const.f32 	%f401, [platform_xmin];	// id:5099 platform_xmin+0x0
	set.le.u32.f32 	%r319, %f401, %f366;	// 
	neg.s32 	%r320, %r319;        	// 
	and.b32 	%r321, %r318, %r320; 	// 
	ld.const.f32 	%f402, [platform_zmax];	// id:5096 platform_zmax+0x0
	set.ge.u32.f32 	%r322, %f402, %f392;	// 
	neg.s32 	%r323, %r322;        	// 
	ld.const.f32 	%f403, [platform_zmin];	// id:5097 platform_zmin+0x0
	set.le.u32.f32 	%r324, %f403, %f392;	// 
	neg.s32 	%r325, %r324;        	// 
	and.b32 	%r326, %r323, %r325; 	// 
	and.b32 	%r327, %r321, %r326; 	// 
	mov.f32 	%f404, 0f00000000;   	// 0
	set.ge.u32.f32 	%r328, %f24, %f404;	// 
	neg.s32 	%r329, %r328;        	// 
	mov.f32 	%f405, 0f00000000;   	// 0
	set.le.u32.f32 	%r330, %f380, %f405;	// 
	neg.s32 	%r331, %r330;        	// 
	and.b32 	%r332, %r329, %r331; 	// 
	mov.s32 	%r333, 0;            	// 
	setp.ne.u32 	%p46, %r332, %r333;	// 
	selp.s32 	%r334, 1, 0, %p46;  	// 
	mov.s32 	%r335, 0;            	// 
	set.ne.u32.s32 	%r336, %r327, %r335;	// 
	neg.s32 	%r337, %r336;        	// 
	and.b32 	%r338, %r334, %r337; 	// 
	mov.u32 	%r339, 0;            	// 
	setp.eq.s32 	%p47, %r338, %r339;	// 
	@%p47 bra 	$Lt_1_239;         	// 
	.loc	20	1328	0
	sub.f32 	%f406, %f76, %f369;  	// 
	mul.f32 	%f407, %f297, %f328; 	// 
	sub.f32 	%f408, %f406, %f407; 	// 
	mov.f32 	%f409, 0f00000000;   	// 0
	min.f32 	%f410, %f408, %f409; 	// 
	neg.f32 	%f411, %f410;        	// 
	mov.f32 	%f412, 0f3d23d70a;   	// 0.04
	mul.f32 	%f413, %f411, %f412; 	// 
	mul.f32 	%f267, %f346, %f413; 	// 
	mov.f32 	%f270, %f347;        	// 
	.loc	20	1250	0
	abs.f32 	%f414, %f347;        	// 
	mov.f32 	%f415, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p48, %f414, %f415;	// 
	@!%p48 bra 	$Lt_1_241;        	// 
	.loc	20	1251	0
	mov.f32 	%f416, 0f3e800000;   	// 0.25
	mul.f32 	%f267, %f267, %f416; 	// 
	.loc	20	1252	0
	mov.f32 	%f417, 0f3e800000;   	// 0.25
	mul.f32 	%f270, %f347, %f417; 	// 
$Lt_1_241:
	.loc	3	408	0
	div.f32 	%f418, %f267, %f270; 	// 
	mul.f32 	%f419, %f395, %f395; 	// 
	mul.f32 	%f420, %f399, %f399; 	// 
	add.f32 	%f421, %f419, %f420; 	// 
	sqrt.f32 	%f422, %f421;       	// 
	setp.ge.f32 	%p49, %f418, %f422;	// 
	@!%p49 bra 	$Lt_1_244;        	// 
	mov.f32 	%f399, 0f00000000;   	// 0
	mov.f32 	%f397, 0f00000000;   	// 0
	mov.f32 	%f395, 0f00000000;   	// 0
	bra.uni 	$Lt_1_243;           	// 
$Lt_1_244:
	.loc	3	66	0
	rsqrt.f32 	%f423, %f421;      	// 
	mul.f32 	%f424, %f423, %f399; 	// 
	mul.f32 	%f425, %f418, %f424; 	// 
	sub.f32 	%f426, %f399, %f425; 	// 
	.loc	3	67	0
	mul.f32 	%f427, %f423, %f395; 	// 
	mul.f32 	%f428, %f418, %f427; 	// 
	sub.f32 	%f395, %f395, %f428; 	// 
	mov.f32 	%f399, %f426;        	// 
	.loc	3	412	0
	mov.f32 	%f397, 0f00000000;   	// 0
$Lt_1_243:
	mov.f32 	%f381, 0f00000000;   	// 0
$Lt_1_239:
	.loc	3	415	0
	mul.lo.u32 	%r340, %r48, 48;  	// 
	ld.param.u32 	%r341, [__cudaparm__Z13pass_verticesP13CUDA_Vtx_Data_vtx_data_out];	// id:5100 __cudaparm__Z13pass_verticesP13CUDA_Vtx_Data_vtx_data_out+0x0
	add.u32 	%r342, %r341, %r340; 	// 
	st.global.v2.f32 	[%r342+0], {%f326,%f328};	// 
	st.global.f32 	[%r342+8], %f330;	// id:5103
	st.global.v2.f32 	[%r342+16], {%f395,%f397};	// 
	.loc	3	416	0
	st.global.f32 	[%r342+24], %f399;	// id:5106
	st.global.v2.f32 	[%r342+32], {%f366,%f381};	// 
	.loc	3	417	0
	st.global.f32 	[%r342+40], %f392;	// id:5109
	.loc	3	418	0
	exit;                         	// 
$LDWend__Z13pass_verticesP13CUDA_Vtx_Data:
	} // _Z13pass_verticesP13CUDA_Vtx_Data

	.entry _Z12pass_unifiedP13CUDA_Vtx_DataPfS1_
	{
	.reg .u16 %rh<4>;
	.reg .u32 %r<191>;
	.reg .f32 %f<409>;
	.reg .pred %p<32>;
	.param .u32 __cudaparm__Z12pass_unifiedP13CUDA_Vtx_DataPfS1__vtx_data_out;
	.param .u32 __cudaparm__Z12pass_unifiedP13CUDA_Vtx_DataPfS1__tower_volumes_in;
	.param .u32 __cudaparm__Z12pass_unifiedP13CUDA_Vtx_DataPfS1__tower_volumes_out;
	.shared .align 4 .b8 __cuda_tri_shared528[1792];
	.shared .align 4 .b8 __cuda_volumes2320[256];
	.shared .align 4 .b8 __cuda_volumes_read2576[256];
	.loc	3	451	0
$LBB1__Z12pass_unifiedP13CUDA_Vtx_DataPfS1_:
	cvt.s32.u16 	%r1, %ctaid.x;   	// 
	cvt.s32.u16 	%r2, %ntid.x;    	// 
	mul24.lo.s32 	%r3, %r1, %r2;  	// 
	cvt.s32.u16 	%r4, %tid.x;     	// 
	add.s32 	%r5, %r3, %r4;       	// 
	ld.const.s32 	%r6, [point_count];	// id:5224 point_count+0x0
	setp.ge.s32 	%p1, %r5, %r6;   	// 
	@%p1 bra 	$Lt_2_181;          	// 
	.loc	3	479	0
	mul.lo.s32 	%r7, %r5, 3;      	// 
	add.s32 	%r8, %r7, 2;         	// 
	mov.s32 	%r9, 0;              	// 
	mov.s32 	%r10, 0;             	// 
	mov.s32 	%r11, 0;             	// 
	tex.1d.v4.f32.s32 {%f1,%f2,%f3,%f4},[vtx_data_tex,{%r8,%r9,%r10,%r11}];
	.loc	3	171	0
	mov.f32 	%f5, %f1;            	// 
	mov.f32 	%f6, %f2;            	// 
	mov.f32 	%f7, %f3;            	// 
	.loc	3	479	0
	mov.f32 	%f8, %f5;            	// 
	mov.f32 	%f9, %f6;            	// 
	mov.f32 	%f10, %f7;           	// 
	bra.uni 	$Lt_2_180;           	// 
$Lt_2_181:
	mov.f32 	%f10, 0f00000000;    	// 0
	mov.f32 	%f9, 0f00000000;     	// 0
	mov.f32 	%f8, 0f00000000;     	// 0
$Lt_2_180:
	ld.const.s32 	%r12, [tri_work_per_vtx];	// id:5220 tri_work_per_vtx+0x0
	mov.u32 	%r13, 0;             	// 
	setp.le.s32 	%p2, %r12, %r13; 	// 
	mov.f32 	%f11, 0f00000000;    	// 0
	mov.f32 	%f12, 0f00000000;    	// 0
	mov.f32 	%f13, 0f00000000;    	// 0
	mov.f32 	%f14, 0f00000000;    	// 0
	mov.f32 	%f15, 0f00000000;    	// 0
	mov.f32 	%f16, 0f00000000;    	// 0
	mov.f32 	%f17, 0f00000000;    	// 0
	@%p2 bra 	$Lt_2_290;          	// 
	mov.s32 	%r14, %r12;          	// 
	mov.u32 	%r15, 0;             	// 
	mul.lo.u32 	%r16, %r12, 32;   	// 
	mul24.lo.s32 	%r17, %r5, %r12;	// 
	mul.lo.u32 	%r18, %r17, 32;   	// 
	ld.const.u32 	%r19, [tri_work_strc];	// id:5210 tri_work_strc+0x0
	add.s32 	%r20, %r18, %r19;    	// 
	mov.u32 	%r21, __cuda_tri_shared528;	// 
	mov.s32 	%r22, %r14;          	// 
$Lt_2_184:
 //<loop> Loop body line 479, nesting depth: 1, estimated iterations: unknown
	.loc	3	484	0
	add.s32 	%r23, %r20, %r15;    	// 
	ld.global.v4.s16 	{%r24,%r25,%r26,%r27}, [%r23+0];	// 
	ld.global.v4.s16 	{%r28,%r29,%r30,%r31}, [%r23+8];	// 
	ld.global.v4.s8 	{%r32,%r33,%r34,%r35}, [%r23+20];	// 
	mov.u32 	%r36, -1;            	// 
	setp.eq.s32 	%p3, %r24, %r36; 	// 
	@%p3 bra 	$Lt_2_185;          	// 
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	491	0
	mul.lo.s32 	%r37, %r24, 3;    	// 
	add.s32 	%r38, %r37, 2;       	// 
	mov.s32 	%r39, 0;             	// 
	mov.s32 	%r40, 0;             	// 
	mov.s32 	%r41, 0;             	// 
	tex.1d.v4.f32.s32 {%f18,%f19,%f20,%f21},[vtx_data_tex,{%r38,%r39,%r40,%r41}];
	.loc	3	171	0
	mov.f32 	%f22, %f18;          	// 
	mov.f32 	%f23, %f19;          	// 
	mov.f32 	%f24, %f20;          	// 
	.loc	3	492	0
	mul.lo.s32 	%r42, %r25, 3;    	// 
	add.s32 	%r43, %r42, 2;       	// 
	mov.s32 	%r44, 0;             	// 
	mov.s32 	%r45, 0;             	// 
	mov.s32 	%r46, 0;             	// 
	tex.1d.v4.f32.s32 {%f25,%f26,%f27,%f28},[vtx_data_tex,{%r43,%r44,%r45,%r46}];
	.loc	3	171	0
	mov.f32 	%f29, %f25;          	// 
	mov.f32 	%f30, %f26;          	// 
	mov.f32 	%f31, %f27;          	// 
	.loc	3	493	0
	mul.lo.s32 	%r47, %r26, 3;    	// 
	add.s32 	%r48, %r47, 2;       	// 
	mov.s32 	%r49, 0;             	// 
	mov.s32 	%r50, 0;             	// 
	mov.s32 	%r51, 0;             	// 
	tex.1d.v4.f32.s32 {%f32,%f33,%f34,%f35},[vtx_data_tex,{%r48,%r49,%r50,%r51}];
	.loc	3	171	0
	mov.f32 	%f36, %f32;          	// 
	mov.f32 	%f37, %f33;          	// 
	mov.f32 	%f38, %f34;          	// 
	.loc	3	491	0
	add.f32 	%f39, %f30, %f37;    	// 
	sub.f32 	%f40, %f36, %f22;    	// 
	sub.f32 	%f41, %f31, %f24;    	// 
	sub.f32 	%f42, %f29, %f22;    	// 
	sub.f32 	%f43, %f38, %f24;    	// 
	add.f32 	%f44, %f39, %f23;    	// 
	mul.f32 	%f45, %f40, %f41;    	// 
	mul.f32 	%f46, %f42, %f43;    	// 
	mov.f32 	%f47, 0f3eaaaaab;    	// 0.333333
	mul.f32 	%f48, %f44, %f47;    	// 
	sub.f32 	%f49, %f45, %f46;    	// 
	mul.f32 	%f50, %f48, %f49;    	// 
	sub.f32 	%f51, %f17, %f50;    	// 
	and.b32 	%r52, %r27, 1;       	// 
	cvt.s8.s32 	%r53, %r52;       	// 
	mov.s32 	%r54, 0;             	// 
	setp.ne.s32 	%p4, %r53, %r54; 	// 
	selp.f32 	%f17, %f51, %f17, %p4;	// 
	.loc	3	507	0
	add.f32 	%f52, %f29, %f36;    	// 
	add.f32 	%f53, %f31, %f38;    	// 
	add.f32 	%f54, %f52, %f22;    	// 
	add.f32 	%f55, %f53, %f24;    	// 
	mov.f32 	%f56, 0f3eaaaaab;    	// 0.333333
	mul.f32 	%f57, %f54, %f56;    	// 
	mov.f32 	%f58, 0f3eaaaaab;    	// 0.333333
	mul.f32 	%f59, %f55, %f58;    	// 
	sub.f32 	%f60, %f48, %f23;    	// 
	sub.f32 	%f61, %f48, %f30;    	// 
	sub.f32 	%f62, %f48, %f37;    	// 
	sub.f32 	%f63, %f57, %f22;    	// 
	sub.f32 	%f64, %f57, %f29;    	// 
	sub.f32 	%f65, %f57, %f36;    	// 
	sub.f32 	%f66, %f59, %f24;    	// 
	sub.f32 	%f67, %f59, %f31;    	// 
	sub.f32 	%f68, %f59, %f38;    	// 
	mul.f32 	%f69, %f60, %f60;    	// 
	mad.f32 	%f70, %f63, %f63, %f69;	// 
	mad.f32 	%f71, %f66, %f66, %f70;	// 
	sqrt.f32 	%f72, %f71;         	// 
	mul.f32 	%f73, %f61, %f61;    	// 
	mad.f32 	%f74, %f64, %f64, %f73;	// 
	mad.f32 	%f75, %f67, %f67, %f74;	// 
	sqrt.f32 	%f76, %f75;         	// 
	add.f32 	%f77, %f72, %f76;    	// 
	mul.f32 	%f78, %f62, %f62;    	// 
	mad.f32 	%f79, %f65, %f65, %f78;	// 
	mad.f32 	%f80, %f68, %f68, %f79;	// 
	sqrt.f32 	%f81, %f80;         	// 
	add.f32 	%f82, %f77, %f81;    	// 
	add.s32 	%r55, %r15, %r18;    	// 
	add.u32 	%r56, %r19, %r55;    	// 
	ld.global.f32 	%f83, [%r56+16];	// id:5247
	sub.f32 	%f84, %f82, %f83;    	// 
	mov.f32 	%f85, 0f00000000;    	// 0
	max.f32 	%f86, %f84, %f85;    	// 
	.loc	3	510	0
	mul24.lo.u32 	%r57, %r4, 28;  	// 
	add.u32 	%r58, %r57, %r21;    	// 
	st.shared.f32 	[%r58+0], %f57;	// id:5248 __cuda_tri_shared528+0x0
	st.shared.f32 	[%r58+4], %f48;	// id:5249 __cuda_tri_shared528+0x0
	st.shared.f32 	[%r58+8], %f59;	// id:5250 __cuda_tri_shared528+0x0
	.loc	3	511	0
	sub.f32 	%f87, %f30, %f23;    	// 
	sub.f32 	%f88, %f37, %f23;    	// 
	mul.f32 	%f89, %f41, %f88;    	// 
	mul.f32 	%f90, %f43, %f87;    	// 
	sub.f32 	%f91, %f90, %f89;    	// 
	st.shared.f32 	[%r58+16], %f91;	// id:5251 __cuda_tri_shared528+0x0
	st.shared.f32 	[%r58+20], %f49;	// id:5252 __cuda_tri_shared528+0x0
	mul.f32 	%f92, %f40, %f87;    	// 
	mul.f32 	%f93, %f42, %f88;    	// 
	sub.f32 	%f94, %f93, %f92;    	// 
	st.shared.f32 	[%r58+24], %f94;	// id:5253 __cuda_tri_shared528+0x0
	.loc	3	512	0
	ld.const.f32 	%f95, [spring_constant];	// id:5254 spring_constant+0x0
	mul.f32 	%f96, %f95, %f86;    	// 
	st.shared.f32 	[%r58+12], %f96;	// id:5255 __cuda_tri_shared528+0x0
$Lt_2_185:
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	515	0
	bar.sync 	0;                  	// 
	shr.s32 	%r59, %r27, 1;       	// 
	mov.u32 	%r60, 0;             	// 
	setp.le.s32 	%p5, %r59, %r60; 	// 
	@%p5 bra 	$Lt_2_187;          	// 
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	67	0
	mul.lo.u32 	%r61, %r32, 28;   	// 
	add.u32 	%r62, %r61, %r21;    	// 
	ld.shared.f32 	%f97, [%r62+16];	// id:5256 __cuda_tri_shared528+0x0
	add.f32 	%f13, %f97, %f13;    	// 
	ld.shared.f32 	%f98, [%r62+20];	// id:5257 __cuda_tri_shared528+0x0
	add.f32 	%f12, %f98, %f12;    	// 
	ld.shared.f32 	%f99, [%r62+24];	// id:5258 __cuda_tri_shared528+0x0
	add.f32 	%f11, %f99, %f11;    	// 
	.loc	3	220	0
	mul.lo.s32 	%r63, %r28, 3;    	// 
	add.s32 	%r64, %r63, 2;       	// 
	mov.s32 	%r65, 0;             	// 
	mov.s32 	%r66, 0;             	// 
	mov.s32 	%r67, 0;             	// 
	tex.1d.v4.f32.s32 {%f100,%f101,%f102,%f103},[vtx_data_tex,{%r64,%r65,%r66,%r67}];
	.loc	3	171	0
	mov.f32 	%f104, %f100;        	// 
	mov.f32 	%f105, %f101;        	// 
	mov.f32 	%f106, %f102;        	// 
	.loc	20	1328	0
	ld.const.f32 	%f107, [rep_constant];	// id:5259 rep_constant+0x0
	mov.f32 	%f108, %f107;        	// 
	sub.f32 	%f109, %f8, %f104;   	// 
	sub.f32 	%f110, %f9, %f105;   	// 
	sub.f32 	%f111, %f10, %f106;  	// 
	mul.f32 	%f112, %f109, %f109; 	// 
	mul.f32 	%f113, %f110, %f110; 	// 
	mul.f32 	%f114, %f111, %f111; 	// 
	add.f32 	%f115, %f112, %f113; 	// 
	add.f32 	%f116, %f114, %f115; 	// 
	mov.f32 	%f117, 0f3a83126f;   	// 0.001
	max.f32 	%f118, %f116, %f117; 	// 
	mov.f32 	%f119, %f118;        	// 
	.loc	20	1250	0
	abs.f32 	%f120, %f118;        	// 
	mov.f32 	%f121, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p6, %f120, %f121;	// 
	@!%p6 bra 	$Lt_2_189;         	// 
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	20	1251	0
	mov.f32 	%f122, 0f3e800000;   	// 0.25
	mul.f32 	%f108, %f107, %f122; 	// 
	.loc	20	1252	0
	mov.f32 	%f123, 0f3e800000;   	// 0.25
	mul.f32 	%f119, %f118, %f123; 	// 
$Lt_2_189:
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	67	0
	div.f32 	%f124, %f108, %f119; 	// 
	rsqrt.f32 	%f125, %f116;      	// 
	mul.f32 	%f126, %f109, %f125; 	// 
	mad.f32 	%f16, %f124, %f126, %f16;	// 
	mul.f32 	%f127, %f110, %f125; 	// 
	mad.f32 	%f15, %f124, %f127, %f15;	// 
	mul.f32 	%f128, %f111, %f125; 	// 
	mad.f32 	%f14, %f124, %f128, %f14;	// 
	.loc	3	537	0
	ld.shared.f32 	%f129, [%r62+12];	// id:5260 __cuda_tri_shared528+0x0
	.loc	3	67	0
	ld.shared.f32 	%f130, [%r62+0];	// id:5261 __cuda_tri_shared528+0x0
	sub.f32 	%f131, %f130, %f8;   	// 
	mad.f32 	%f16, %f129, %f131, %f16;	// 
	ld.shared.f32 	%f132, [%r62+4];	// id:5262 __cuda_tri_shared528+0x0
	sub.f32 	%f133, %f132, %f9;   	// 
	mad.f32 	%f15, %f129, %f133, %f15;	// 
	ld.shared.f32 	%f134, [%r62+8];	// id:5263 __cuda_tri_shared528+0x0
	sub.f32 	%f135, %f134, %f10;  	// 
	mad.f32 	%f14, %f129, %f135, %f14;	// 
$Lt_2_187:
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	537	0
	mov.u32 	%r68, 1;             	// 
	setp.le.s32 	%p7, %r59, %r68; 	// 
	@%p7 bra 	$Lt_2_191;          	// 
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	67	0
	mul.lo.u32 	%r69, %r33, 28;   	// 
	add.u32 	%r70, %r69, %r21;    	// 
	ld.shared.f32 	%f136, [%r70+16];	// id:5264 __cuda_tri_shared528+0x0
	add.f32 	%f13, %f136, %f13;   	// 
	ld.shared.f32 	%f137, [%r70+20];	// id:5265 __cuda_tri_shared528+0x0
	add.f32 	%f12, %f137, %f12;   	// 
	ld.shared.f32 	%f138, [%r70+24];	// id:5266 __cuda_tri_shared528+0x0
	add.f32 	%f11, %f138, %f11;   	// 
	.loc	3	220	0
	mul.lo.s32 	%r71, %r29, 3;    	// 
	add.s32 	%r72, %r71, 2;       	// 
	mov.s32 	%r73, 0;             	// 
	mov.s32 	%r74, 0;             	// 
	mov.s32 	%r75, 0;             	// 
	tex.1d.v4.f32.s32 {%f139,%f140,%f141,%f142},[vtx_data_tex,{%r72,%r73,%r74,%r75}];
	.loc	3	171	0
	mov.f32 	%f104, %f139;        	// 
	mov.f32 	%f105, %f140;        	// 
	mov.f32 	%f106, %f141;        	// 
	.loc	20	1328	0
	ld.const.f32 	%f107, [rep_constant];	// id:5259 rep_constant+0x0
	mov.f32 	%f108, %f107;        	// 
	sub.f32 	%f109, %f8, %f104;   	// 
	sub.f32 	%f110, %f9, %f105;   	// 
	sub.f32 	%f111, %f10, %f106;  	// 
	mul.f32 	%f112, %f109, %f109; 	// 
	mul.f32 	%f113, %f110, %f110; 	// 
	mul.f32 	%f114, %f111, %f111; 	// 
	add.f32 	%f115, %f112, %f113; 	// 
	add.f32 	%f116, %f114, %f115; 	// 
	mov.f32 	%f143, 0f3a83126f;   	// 0.001
	max.f32 	%f144, %f116, %f143; 	// 
	mov.f32 	%f119, %f144;        	// 
	.loc	20	1250	0
	abs.f32 	%f145, %f144;        	// 
	mov.f32 	%f146, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p8, %f145, %f146;	// 
	@!%p8 bra 	$Lt_2_193;         	// 
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	20	1251	0
	mov.f32 	%f147, 0f3e800000;   	// 0.25
	mul.f32 	%f108, %f107, %f147; 	// 
	.loc	20	1252	0
	mov.f32 	%f148, 0f3e800000;   	// 0.25
	mul.f32 	%f119, %f144, %f148; 	// 
$Lt_2_193:
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	67	0
	div.f32 	%f124, %f108, %f119; 	// 
	rsqrt.f32 	%f125, %f116;      	// 
	mul.f32 	%f149, %f109, %f125; 	// 
	mad.f32 	%f16, %f124, %f149, %f16;	// 
	mul.f32 	%f150, %f110, %f125; 	// 
	mad.f32 	%f15, %f124, %f150, %f15;	// 
	mul.f32 	%f151, %f111, %f125; 	// 
	mad.f32 	%f14, %f124, %f151, %f14;	// 
	.loc	3	537	0
	ld.shared.f32 	%f152, [%r70+12];	// id:5267 __cuda_tri_shared528+0x0
	.loc	3	67	0
	ld.shared.f32 	%f153, [%r70+0];	// id:5268 __cuda_tri_shared528+0x0
	sub.f32 	%f154, %f153, %f8;   	// 
	mad.f32 	%f16, %f152, %f154, %f16;	// 
	ld.shared.f32 	%f155, [%r70+4];	// id:5269 __cuda_tri_shared528+0x0
	sub.f32 	%f156, %f155, %f9;   	// 
	mad.f32 	%f15, %f152, %f156, %f15;	// 
	ld.shared.f32 	%f157, [%r70+8];	// id:5270 __cuda_tri_shared528+0x0
	sub.f32 	%f158, %f157, %f10;  	// 
	mad.f32 	%f14, %f152, %f158, %f14;	// 
$Lt_2_191:
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	537	0
	mov.u32 	%r76, 2;             	// 
	setp.le.s32 	%p9, %r59, %r76; 	// 
	@%p9 bra 	$Lt_2_195;          	// 
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	67	0
	mul.lo.u32 	%r77, %r34, 28;   	// 
	add.u32 	%r78, %r77, %r21;    	// 
	ld.shared.f32 	%f159, [%r78+16];	// id:5271 __cuda_tri_shared528+0x0
	add.f32 	%f13, %f159, %f13;   	// 
	ld.shared.f32 	%f160, [%r78+20];	// id:5272 __cuda_tri_shared528+0x0
	add.f32 	%f12, %f160, %f12;   	// 
	ld.shared.f32 	%f161, [%r78+24];	// id:5273 __cuda_tri_shared528+0x0
	add.f32 	%f11, %f161, %f11;   	// 
	.loc	3	220	0
	mul.lo.s32 	%r79, %r30, 3;    	// 
	add.s32 	%r80, %r79, 2;       	// 
	mov.s32 	%r81, 0;             	// 
	mov.s32 	%r82, 0;             	// 
	mov.s32 	%r83, 0;             	// 
	tex.1d.v4.f32.s32 {%f162,%f163,%f164,%f165},[vtx_data_tex,{%r80,%r81,%r82,%r83}];
	.loc	3	171	0
	mov.f32 	%f104, %f162;        	// 
	mov.f32 	%f105, %f163;        	// 
	mov.f32 	%f106, %f164;        	// 
	.loc	20	1328	0
	ld.const.f32 	%f107, [rep_constant];	// id:5259 rep_constant+0x0
	mov.f32 	%f108, %f107;        	// 
	sub.f32 	%f109, %f8, %f104;   	// 
	sub.f32 	%f110, %f9, %f105;   	// 
	sub.f32 	%f111, %f10, %f106;  	// 
	mul.f32 	%f112, %f109, %f109; 	// 
	mul.f32 	%f113, %f110, %f110; 	// 
	mul.f32 	%f114, %f111, %f111; 	// 
	add.f32 	%f115, %f112, %f113; 	// 
	add.f32 	%f116, %f114, %f115; 	// 
	mov.f32 	%f166, 0f3a83126f;   	// 0.001
	max.f32 	%f167, %f116, %f166; 	// 
	mov.f32 	%f119, %f167;        	// 
	.loc	20	1250	0
	abs.f32 	%f168, %f167;        	// 
	mov.f32 	%f169, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p10, %f168, %f169;	// 
	@!%p10 bra 	$Lt_2_197;        	// 
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	20	1251	0
	mov.f32 	%f170, 0f3e800000;   	// 0.25
	mul.f32 	%f108, %f107, %f170; 	// 
	.loc	20	1252	0
	mov.f32 	%f171, 0f3e800000;   	// 0.25
	mul.f32 	%f119, %f167, %f171; 	// 
$Lt_2_197:
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	67	0
	div.f32 	%f124, %f108, %f119; 	// 
	rsqrt.f32 	%f125, %f116;      	// 
	mul.f32 	%f172, %f109, %f125; 	// 
	mad.f32 	%f16, %f124, %f172, %f16;	// 
	mul.f32 	%f173, %f110, %f125; 	// 
	mad.f32 	%f15, %f124, %f173, %f15;	// 
	mul.f32 	%f174, %f111, %f125; 	// 
	mad.f32 	%f14, %f124, %f174, %f14;	// 
	.loc	3	537	0
	ld.shared.f32 	%f175, [%r78+12];	// id:5274 __cuda_tri_shared528+0x0
	.loc	3	67	0
	ld.shared.f32 	%f176, [%r78+0];	// id:5275 __cuda_tri_shared528+0x0
	sub.f32 	%f177, %f176, %f8;   	// 
	mad.f32 	%f16, %f175, %f177, %f16;	// 
	ld.shared.f32 	%f178, [%r78+4];	// id:5276 __cuda_tri_shared528+0x0
	sub.f32 	%f179, %f178, %f9;   	// 
	mad.f32 	%f15, %f175, %f179, %f15;	// 
	ld.shared.f32 	%f180, [%r78+8];	// id:5277 __cuda_tri_shared528+0x0
	sub.f32 	%f181, %f180, %f10;  	// 
	mad.f32 	%f14, %f175, %f181, %f14;	// 
$Lt_2_195:
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	537	0
	mov.u32 	%r84, 3;             	// 
	setp.le.s32 	%p11, %r59, %r84;	// 
	@%p11 bra 	$Lt_2_199;         	// 
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	67	0
	mul.lo.u32 	%r85, %r35, 28;   	// 
	add.u32 	%r86, %r85, %r21;    	// 
	ld.shared.f32 	%f182, [%r86+16];	// id:5278 __cuda_tri_shared528+0x0
	add.f32 	%f13, %f182, %f13;   	// 
	ld.shared.f32 	%f183, [%r86+20];	// id:5279 __cuda_tri_shared528+0x0
	add.f32 	%f12, %f183, %f12;   	// 
	ld.shared.f32 	%f184, [%r86+24];	// id:5280 __cuda_tri_shared528+0x0
	add.f32 	%f11, %f184, %f11;   	// 
	.loc	3	220	0
	mul.lo.s32 	%r87, %r31, 3;    	// 
	add.s32 	%r88, %r87, 2;       	// 
	mov.s32 	%r89, 0;             	// 
	mov.s32 	%r90, 0;             	// 
	mov.s32 	%r91, 0;             	// 
	tex.1d.v4.f32.s32 {%f185,%f186,%f187,%f188},[vtx_data_tex,{%r88,%r89,%r90,%r91}];
	.loc	3	171	0
	mov.f32 	%f104, %f185;        	// 
	mov.f32 	%f105, %f186;        	// 
	mov.f32 	%f106, %f187;        	// 
	.loc	20	1328	0
	ld.const.f32 	%f107, [rep_constant];	// id:5259 rep_constant+0x0
	mov.f32 	%f108, %f107;        	// 
	sub.f32 	%f109, %f8, %f104;   	// 
	sub.f32 	%f110, %f9, %f105;   	// 
	sub.f32 	%f111, %f10, %f106;  	// 
	mul.f32 	%f112, %f109, %f109; 	// 
	mul.f32 	%f113, %f110, %f110; 	// 
	mul.f32 	%f114, %f111, %f111; 	// 
	add.f32 	%f115, %f112, %f113; 	// 
	add.f32 	%f116, %f114, %f115; 	// 
	mov.f32 	%f189, 0f3a83126f;   	// 0.001
	max.f32 	%f190, %f116, %f189; 	// 
	mov.f32 	%f119, %f190;        	// 
	.loc	20	1250	0
	abs.f32 	%f191, %f190;        	// 
	mov.f32 	%f192, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p12, %f191, %f192;	// 
	@!%p12 bra 	$Lt_2_201;        	// 
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	20	1251	0
	mov.f32 	%f193, 0f3e800000;   	// 0.25
	mul.f32 	%f108, %f107, %f193; 	// 
	.loc	20	1252	0
	mov.f32 	%f194, 0f3e800000;   	// 0.25
	mul.f32 	%f119, %f190, %f194; 	// 
$Lt_2_201:
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	67	0
	div.f32 	%f124, %f108, %f119; 	// 
	rsqrt.f32 	%f125, %f116;      	// 
	mul.f32 	%f195, %f109, %f125; 	// 
	mad.f32 	%f16, %f124, %f195, %f16;	// 
	mul.f32 	%f196, %f110, %f125; 	// 
	mad.f32 	%f15, %f124, %f196, %f15;	// 
	mul.f32 	%f197, %f111, %f125; 	// 
	mad.f32 	%f14, %f124, %f197, %f14;	// 
	.loc	3	537	0
	ld.shared.f32 	%f198, [%r86+12];	// id:5281 __cuda_tri_shared528+0x0
	.loc	3	67	0
	ld.shared.f32 	%f199, [%r86+0];	// id:5282 __cuda_tri_shared528+0x0
	sub.f32 	%f200, %f199, %f8;   	// 
	mad.f32 	%f16, %f198, %f200, %f16;	// 
	ld.shared.f32 	%f201, [%r86+4];	// id:5283 __cuda_tri_shared528+0x0
	sub.f32 	%f202, %f201, %f9;   	// 
	mad.f32 	%f15, %f198, %f202, %f15;	// 
	ld.shared.f32 	%f203, [%r86+8];	// id:5284 __cuda_tri_shared528+0x0
	sub.f32 	%f204, %f203, %f10;  	// 
	mad.f32 	%f14, %f198, %f204, %f14;	// 
$Lt_2_199:
 //<loop> Part of loop body line 479, head labeled $Lt_2_184
	.loc	3	537	0
	add.u32 	%r15, %r15, 32;      	// 
	setp.ne.s32 	%p13, %r15, %r16;	// 
	@%p13 bra 	$Lt_2_184;         	// 
	bra.uni 	$Lt_2_182;           	// 
$Lt_2_290:
$Lt_2_182:
	.loc	3	545	0
	mov.u32 	%r92, __cuda_volumes2320;	// 
	.loc	3	112	0
	mov.f32 	%f205, 0f3f000000;   	// 0.5
	mul.f32 	%f206, %f17, %f205;  	// 
	mul24.lo.u32 	%r93, %r4, 4;   	// 
	add.u32 	%r94, %r93, %r92;    	// 
	st.shared.f32 	[%r94+0], %f206;	// id:5285 __cuda_volumes2320+0x0
	mov.f32 	%f207, %f206;        	// 
	.loc	3	113	0
	bar.sync 	0;                  	// 
	mov.s32 	%r95, 8;             	// 
	setp.lt.s32 	%p14, %r4, %r95; 	// 
	@!%p14 bra 	$Lt_2_204;        	// 
	.loc	3	128	0
	add.s32 	%r96, %r4, 8;        	// 
	mul.lo.u32 	%r97, %r96, 4;    	// 
	add.u32 	%r98, %r92, %r97;    	// 
	ld.shared.f32 	%f208, [%r98+0];	// id:5286 __cuda_volumes2320+0x0
	add.f32 	%f207, %f208, %f206; 	// 
	.loc	3	129	0
	add.s32 	%r99, %r4, 16;       	// 
	mul.lo.u32 	%r100, %r99, 4;   	// 
	add.u32 	%r101, %r92, %r100;  	// 
	ld.shared.f32 	%f209, [%r101+0];	// id:5287 __cuda_volumes2320+0x0
	add.f32 	%f207, %f209, %f207; 	// 
	add.s32 	%r102, %r4, 24;      	// 
	mul.lo.u32 	%r103, %r102, 4;  	// 
	add.u32 	%r104, %r92, %r103;  	// 
	ld.shared.f32 	%f210, [%r104+0];	// id:5288 __cuda_volumes2320+0x0
	add.f32 	%f207, %f210, %f207; 	// 
	.loc	3	130	0
	add.s32 	%r105, %r4, 32;      	// 
	mul.lo.u32 	%r106, %r105, 4;  	// 
	add.u32 	%r107, %r92, %r106;  	// 
	ld.shared.f32 	%f211, [%r107+0];	// id:5289 __cuda_volumes2320+0x0
	add.f32 	%f207, %f211, %f207; 	// 
	add.s32 	%r108, %r4, 40;      	// 
	mul.lo.u32 	%r109, %r108, 4;  	// 
	add.u32 	%r110, %r92, %r109;  	// 
	ld.shared.f32 	%f212, [%r110+0];	// id:5290 __cuda_volumes2320+0x0
	add.f32 	%f207, %f212, %f207; 	// 
	add.s32 	%r111, %r4, 48;      	// 
	mul.lo.u32 	%r112, %r111, 4;  	// 
	add.u32 	%r113, %r92, %r112;  	// 
	ld.shared.f32 	%f213, [%r113+0];	// id:5291 __cuda_volumes2320+0x0
	add.f32 	%f207, %f213, %f207; 	// 
	add.s32 	%r114, %r4, 56;      	// 
	mul.lo.u32 	%r115, %r114, 4;  	// 
	add.u32 	%r116, %r92, %r115;  	// 
	ld.shared.f32 	%f214, [%r116+0];	// id:5292 __cuda_volumes2320+0x0
	add.f32 	%f207, %f214, %f207; 	// 
	.loc	3	135	0
	st.shared.f32 	[%r94+0], %f207;	// id:5293 __cuda_volumes2320+0x0
$Lt_2_204:
	.loc	3	140	0
	bar.sync 	0;                  	// 
	mov.s32 	%r117, 0;            	// 
	setp.eq.s32 	%p15, %r4, %r117;	// 
	@!%p15 bra 	$Lt_2_216;        	// 
	.loc	3	145	0
	ld.shared.f32 	%f215, [__cuda_volumes2320+4];	// id:5294 __cuda_volumes2320+0x4
	add.f32 	%f207, %f215, %f207; 	// 
	.loc	3	146	0
	ld.shared.f32 	%f216, [__cuda_volumes2320+8];	// id:5295 __cuda_volumes2320+0x8
	add.f32 	%f207, %f216, %f207; 	// 
	ld.shared.f32 	%f217, [__cuda_volumes2320+12];	// id:5296 __cuda_volumes2320+0xc
	add.f32 	%f207, %f217, %f207; 	// 
	.loc	3	147	0
	ld.shared.f32 	%f218, [__cuda_volumes2320+16];	// id:5297 __cuda_volumes2320+0x10
	add.f32 	%f207, %f218, %f207; 	// 
	ld.shared.f32 	%f219, [__cuda_volumes2320+20];	// id:5298 __cuda_volumes2320+0x14
	add.f32 	%f207, %f219, %f207; 	// 
	ld.shared.f32 	%f220, [__cuda_volumes2320+24];	// id:5299 __cuda_volumes2320+0x18
	add.f32 	%f207, %f220, %f207; 	// 
	ld.shared.f32 	%f221, [__cuda_volumes2320+28];	// id:5300 __cuda_volumes2320+0x1c
	add.f32 	%f207, %f221, %f207; 	// 
$Lt_2_216:
	@!%p15 bra 	$Lt_2_232;        	// 
	.loc	3	546	0
	ld.param.u32 	%r118, [__cudaparm__Z12pass_unifiedP13CUDA_Vtx_DataPfS1__tower_volumes_out];	// id:5301 __cudaparm__Z12pass_unifiedP13CUDA_Vtx_DataPfS1__tower_volumes_out+0x0
	mul24.lo.u32 	%r119, %r1, 4;  	// 
	add.u32 	%r120, %r118, %r119; 	// 
	st.global.f32 	[%r120+0], %f207;	// id:5302
$Lt_2_232:
	.loc	3	559	0
	cvt.s32.u16 	%r121, %nctaid.x;	// 
	and.b32 	%r122, %r121, -64;   	// 
	shr.s32 	%r123, %r121, 6;     	// 
	setp.ne.s32 	%p16, %r122, %r121;	// 
	selp.s32 	%r124, 1, 0, %p16;  	// 
	add.s32 	%r125, %r123, %r124; 	// 
	mul.lo.s32 	%r126, %r125, %r4;	// 
	add.s32 	%r127, %r125, %r126; 	// 
	min.s32 	%r128, %r127, %r121; 	// 
	setp.le.s32 	%p17, %r128, %r126;	// 
	mov.f32 	%f222, 0f00000000;   	// 0
	@%p17 bra 	$Lt_2_291;         	// 
	sub.s32 	%r129, %r128, %r126; 	// 
	mul.lo.u32 	%r130, %r126, 4;  	// 
	ld.param.u32 	%r131, [__cudaparm__Z12pass_unifiedP13CUDA_Vtx_DataPfS1__tower_volumes_in];	// id:5211 __cudaparm__Z12pass_unifiedP13CUDA_Vtx_DataPfS1__tower_volumes_in+0x0
	add.u32 	%r132, %r130, %r131; 	// 
	mul.lo.u32 	%r133, %r128, 4;  	// 
	add.u32 	%r134, %r133, %r131; 	// 
	mov.s32 	%r135, %r129;        	// 
$Lt_2_236:
 //<loop> Loop body line 559, nesting depth: 1, estimated iterations: unknown
	ld.global.f32 	%f223, [%r132+0];	// id:5303
	add.f32 	%f222, %f223, %f222; 	// 
	add.u32 	%r132, %r132, 4;     	// 
	setp.ne.u32 	%p18, %r132, %r134;	// 
	@%p18 bra 	$Lt_2_236;         	// 
	bra.uni 	$Lt_2_234;           	// 
$Lt_2_291:
$Lt_2_234:
	.loc	3	560	0
	mov.u32 	%r136, __cuda_volumes_read2576;	// 
	.loc	3	112	0
	add.u32 	%r137, %r93, %r136;  	// 
	mov.f32 	%f224, %f222;        	// 
	st.shared.f32 	[%r137+0], %f224;	// id:5304 __cuda_volumes_read2576+0x0
	mov.f32 	%f207, %f224;        	// 
	.loc	3	113	0
	bar.sync 	0;                  	// 
	@!%p14 bra 	$Lt_2_238;        	// 
	.loc	3	128	0
	add.s32 	%r138, %r4, 8;       	// 
	mul.lo.u32 	%r139, %r138, 4;  	// 
	add.u32 	%r140, %r136, %r139; 	// 
	ld.shared.f32 	%f225, [%r140+0];	// id:5305 __cuda_volumes_read2576+0x0
	add.f32 	%f207, %f225, %f224; 	// 
	.loc	3	129	0
	add.s32 	%r141, %r4, 16;      	// 
	mul.lo.u32 	%r142, %r141, 4;  	// 
	add.u32 	%r143, %r136, %r142; 	// 
	ld.shared.f32 	%f226, [%r143+0];	// id:5306 __cuda_volumes_read2576+0x0
	add.f32 	%f207, %f226, %f207; 	// 
	add.s32 	%r144, %r4, 24;      	// 
	mul.lo.u32 	%r145, %r144, 4;  	// 
	add.u32 	%r146, %r136, %r145; 	// 
	ld.shared.f32 	%f227, [%r146+0];	// id:5307 __cuda_volumes_read2576+0x0
	add.f32 	%f207, %f227, %f207; 	// 
	.loc	3	130	0
	add.s32 	%r147, %r4, 32;      	// 
	mul.lo.u32 	%r148, %r147, 4;  	// 
	add.u32 	%r149, %r136, %r148; 	// 
	ld.shared.f32 	%f228, [%r149+0];	// id:5308 __cuda_volumes_read2576+0x0
	add.f32 	%f207, %f228, %f207; 	// 
	add.s32 	%r150, %r4, 40;      	// 
	mul.lo.u32 	%r151, %r150, 4;  	// 
	add.u32 	%r152, %r136, %r151; 	// 
	ld.shared.f32 	%f229, [%r152+0];	// id:5309 __cuda_volumes_read2576+0x0
	add.f32 	%f207, %f229, %f207; 	// 
	add.s32 	%r153, %r4, 48;      	// 
	mul.lo.u32 	%r154, %r153, 4;  	// 
	add.u32 	%r155, %r136, %r154; 	// 
	ld.shared.f32 	%f230, [%r155+0];	// id:5310 __cuda_volumes_read2576+0x0
	add.f32 	%f207, %f230, %f207; 	// 
	add.s32 	%r156, %r4, 56;      	// 
	mul.lo.u32 	%r157, %r156, 4;  	// 
	add.u32 	%r158, %r136, %r157; 	// 
	ld.shared.f32 	%f231, [%r158+0];	// id:5311 __cuda_volumes_read2576+0x0
	add.f32 	%f207, %f231, %f207; 	// 
	.loc	3	135	0
	st.shared.f32 	[%r137+0], %f207;	// id:5312 __cuda_volumes_read2576+0x0
$Lt_2_238:
	.loc	3	140	0
	bar.sync 	0;                  	// 
	@!%p15 bra 	$Lt_2_264;        	// 
	.loc	3	145	0
	ld.shared.f32 	%f232, [__cuda_volumes_read2576+4];	// id:5313 __cuda_volumes_read2576+0x4
	add.f32 	%f207, %f232, %f207; 	// 
	.loc	3	146	0
	ld.shared.f32 	%f233, [__cuda_volumes_read2576+8];	// id:5314 __cuda_volumes_read2576+0x8
	add.f32 	%f207, %f233, %f207; 	// 
	ld.shared.f32 	%f234, [__cuda_volumes_read2576+12];	// id:5315 __cuda_volumes_read2576+0xc
	add.f32 	%f207, %f234, %f207; 	// 
	.loc	3	147	0
	ld.shared.f32 	%f235, [__cuda_volumes_read2576+16];	// id:5316 __cuda_volumes_read2576+0x10
	add.f32 	%f207, %f235, %f207; 	// 
	ld.shared.f32 	%f236, [__cuda_volumes_read2576+20];	// id:5317 __cuda_volumes_read2576+0x14
	add.f32 	%f207, %f236, %f207; 	// 
	ld.shared.f32 	%f237, [__cuda_volumes_read2576+24];	// id:5318 __cuda_volumes_read2576+0x18
	add.f32 	%f207, %f237, %f207; 	// 
	ld.shared.f32 	%f238, [__cuda_volumes_read2576+28];	// id:5319 __cuda_volumes_read2576+0x1c
	add.f32 	%f207, %f238, %f207; 	// 
$Lt_2_250:
	@!%p15 bra 	$Lt_2_264;        	// 
	.loc	3	154	0
	st.shared.f32 	[__cuda_volumes_read2576+0], %f207;	// id:5320 __cuda_volumes_read2576+0x0
$Lt_2_264:
	.loc	3	155	0
	bar.sync 	0;                  	// 
	setp.lt.s32 	%p19, %r5, %r6;  	// 
	@%p19 bra 	$Lt_2_266;         	// 
	bra.uni 	$LBB66__Z12pass_unifiedP13CUDA_Vtx_DataPfS1_;	// 
$Lt_2_266:
	.loc	3	575	0
	mul.lo.s32 	%r159, %r5, 3;    	// 
	add.s32 	%r160, %r159, 1;     	// 
	mov.s32 	%r161, 0;            	// 
	mov.s32 	%r162, 0;            	// 
	mov.s32 	%r163, 0;            	// 
	tex.1d.v4.f32.s32 {%f239,%f240,%f241,%f242},[vtx_data_tex,{%r160,%r161,%r162,%r163}];
	.loc	3	178	0
	mov.f32 	%f243, %f239;        	// 
	mov.f32 	%f244, %f240;        	// 
	mov.f32 	%f245, %f241;        	// 
	.loc	20	1328	0
	ld.const.f32 	%f246, [pressure_factor_coeff];	// id:5321 pressure_factor_coeff+0x0
	ld.shared.f32 	%f247, [__cuda_volumes_read2576+0];	// id:5320 __cuda_volumes_read2576+0x0
	abs.f32 	%f248, %f247;        	// 
	mov.f32 	%f249, %f248;        	// 
	.loc	20	1250	0
	mov.f32 	%f250, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p20, %f248, %f250;	// 
	@!%p20 bra 	$Lt_2_268;        	// 
	.loc	20	1251	0
	mov.f32 	%f251, 0f3e800000;   	// 0.25
	mul.f32 	%f246, %f246, %f251; 	// 
	.loc	20	1252	0
	mov.f32 	%f252, 0f3e800000;   	// 0.25
	mul.f32 	%f249, %f248, %f252; 	// 
$Lt_2_268:
	.loc	3	579	0
	ld.const.s8 	%rh1, [opt_gravity];	// id:5322 opt_gravity+0x0
	mov.s16 	%rh2, 0;             	// 
	setp.ne.s16 	%p21, %rh1, %rh2;	// 
	div.f32 	%f253, %f246, %f249; 	// 
	@!%p21 bra 	$Lt_2_271;        	// 
	.loc	3	581	0
	ld.const.f32 	%f254, [gas_m_over_temp];	// id:5323 gas_m_over_temp+0x0
	mul.f32 	%f255, %f254, %f9;   	// 
	neg.f32 	%f256, %f255;        	// 
	mov.f32 	%f257, 0f3fb8aa3b;   	// 1.4427
	mul.f32 	%f258, %f256, %f257; 	// 
	cvt.rzi.f32.f32 	%f259, %f258;	// 
	mov.f32 	%f260, 0f7f800000;   	// ((1.0F)/(0.0F))
	mov.f32 	%f261, 0f00000000;   	// 0
	ex2.f32 	%f262, %f259;        	// 
	mov.f32 	%f263, 0f3f317200;   	// 0.693146
	mad.f32 	%f264, %f259, %f263, %f255;	// 
	mov.f32 	%f265, 0f35bfbe8e;   	// 1.42861e-06
	mad.f32 	%f266, %f259, %f265, %f264;	// 
	neg.f32 	%f267, %f266;        	// 
	mov.f32 	%f268, 0f3fb8aa3b;   	// 1.4427
	mul.f32 	%f269, %f267, %f268; 	// 
	ex2.f32 	%f270, %f269;        	// 
	mul.f32 	%f271, %f262, %f270; 	// 
	mov.f32 	%f272, 0fc2d20000;   	// -105
	setp.lt.f32 	%p22, %f256, %f272;	// 
	selp.f32 	%f273, %f261, %f271, %p22;	// 
	mov.f32 	%f274, 0f42d20000;   	// 105
	setp.gt.f32 	%p23, %f256, %f274;	// 
	selp.f32 	%f275, %f260, %f273, %p23;	// 
	mul.f32 	%f276, %f275, %f253; 	// 
	bra.uni 	$Lt_2_270;           	// 
$Lt_2_271:
	mov.f32 	%f276, %f253;        	// 
$Lt_2_270:
	@!%p21 bra 	$Lt_2_273;        	// 
	.loc	3	585	0
	ld.const.f32 	%f277, [air_particle_mass];	// id:5324 air_particle_mass+0x0
	mov.f32 	%f278, 0fbe4ccccd;   	// -0.2
	mul.f32 	%f279, %f277, %f278; 	// 
	mul.f32 	%f280, %f279, %f9;   	// 
	mov.f32 	%f281, 0f3fb8aa3b;   	// 1.4427
	mul.f32 	%f282, %f280, %f281; 	// 
	cvt.rzi.f32.f32 	%f283, %f282;	// 
	mov.f32 	%f284, 0f7f800000;   	// ((1.0F)/(0.0F))
	mov.f32 	%f285, 0f00000000;   	// 0
	ex2.f32 	%f286, %f283;        	// 
	mov.f32 	%f287, 0f3f317200;   	// 0.693146
	mul.f32 	%f288, %f283, %f287; 	// 
	sub.f32 	%f289, %f280, %f288; 	// 
	mov.f32 	%f290, 0f35bfbe8e;   	// 1.42861e-06
	mul.f32 	%f291, %f283, %f290; 	// 
	sub.f32 	%f292, %f289, %f291; 	// 
	mov.f32 	%f293, 0f3fb8aa3b;   	// 1.4427
	mul.f32 	%f294, %f292, %f293; 	// 
	ex2.f32 	%f295, %f294;        	// 
	mul.f32 	%f296, %f286, %f295; 	// 
	mov.f32 	%f297, 0fc2d20000;   	// -105
	setp.lt.f32 	%p24, %f280, %f297;	// 
	selp.f32 	%f298, %f285, %f296, %p24;	// 
	mov.f32 	%f299, 0f42d20000;   	// 105
	setp.gt.f32 	%p25, %f280, %f299;	// 
	selp.f32 	%f300, %f284, %f298, %p25;	// 
	bra.uni 	$Lt_2_272;           	// 
$Lt_2_273:
	mov.f32 	%f300, 0f3f800000;   	// 1
$Lt_2_272:
	.loc	3	594	0
	mul.f32 	%f301, %f243, %f243; 	// 
	mul.f32 	%f302, %f244, %f244; 	// 
	mul.f32 	%f303, %f245, %f245; 	// 
	mov.f32 	%f304, 0f3e2aaaab;   	// 0.166667
	mul.f32 	%f305, %f13, %f304;  	// 
	mov.f32 	%f306, 0f3e2aaaab;   	// 0.166667
	mul.f32 	%f307, %f12, %f306;  	// 
	mov.f32 	%f308, 0f3e2aaaab;   	// 0.166667
	mul.f32 	%f309, %f11, %f308;  	// 
	add.f32 	%f310, %f301, %f302; 	// 
	add.f32 	%f311, %f303, %f310; 	// 
	rsqrt.f32 	%f312, %f311;      	// 
	ld.const.f32 	%f313, [air_resistance];	// id:5325 air_resistance+0x0
	mul.f32 	%f314, %f312, %f244; 	// 
	mul.f32 	%f315, %f307, %f314; 	// 
	mul.f32 	%f316, %f312, %f243; 	// 
	mad.f32 	%f317, %f316, %f305, %f315;	// 
	mul.f32 	%f318, %f312, %f245; 	// 
	mad.f32 	%f319, %f318, %f309, %f317;	// 
	neg.f32 	%f320, %f319;        	// 
	mov.f32 	%f321, 0f00000000;   	// 0
	max.f32 	%f322, %f320, %f321; 	// 
	mul.f32 	%f323, %f313, %f322; 	// 
	neg.f32 	%f324, %f323;        	// 
	.loc	20	1328	0
	ld.const.f32 	%f325, [delta_t];	// id:5326 delta_t+0x0
	mov.f32 	%f246, %f325;        	// 
	ld.const.f32 	%f326, [point_mass];	// id:5327 point_mass+0x0
	ld.const.f32 	%f327, [gas_mass_per_vertex];	// id:5328 gas_mass_per_vertex+0x0
	add.f32 	%f328, %f327, %f326; 	// 
	mov.f32 	%f249, %f328;        	// 
	.loc	20	1250	0
	abs.f32 	%f329, %f328;        	// 
	mov.f32 	%f330, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p26, %f329, %f330;	// 
	@!%p26 bra 	$Lt_2_274;        	// 
	.loc	20	1251	0
	mov.f32 	%f331, 0f3e800000;   	// 0.25
	mul.f32 	%f246, %f325, %f331; 	// 
	.loc	20	1252	0
	mov.f32 	%f332, 0f3e800000;   	// 0.25
	mul.f32 	%f249, %f328, %f332; 	// 
$Lt_2_274:
	.loc	3	612	0
	div.f32 	%f333, %f246, %f249; 	// 
	mul.f32 	%f334, %f243, %f324; 	// 
	sub.f32 	%f335, %f300, %f276; 	// 
	mul.f32 	%f336, %f333, %f16;  	// 
	mul.f32 	%f337, %f305, %f335; 	// 
	add.f32 	%f338, %f334, %f337; 	// 
	mul.f32 	%f339, %f333, %f338; 	// 
	add.f32 	%f340, %f336, %f339; 	// 
	mov.f32 	%f341, 0f3f000000;   	// 0.5
	mul.f32 	%f342, %f340, %f341; 	// 
	add.f32 	%f343, %f342, %f243; 	// 
	mul.f32 	%f344, %f343, %f325; 	// 
	add.f32 	%f345, %f344, %f8;   	// 
	mul.f32 	%f346, %f324, %f244; 	// 
	ld.const.f32 	%f347, [gravity_mag];	// id:5329 gravity_mag+0x0
	mul.f32 	%f348, %f347, %f326; 	// 
	mul.f32 	%f349, %f333, %f15;  	// 
	mul.f32 	%f350, %f307, %f335; 	// 
	sub.f32 	%f351, %f350, %f348; 	// 
	add.f32 	%f352, %f346, %f351; 	// 
	mul.f32 	%f353, %f333, %f352; 	// 
	add.f32 	%f354, %f349, %f353; 	// 
	mov.f32 	%f355, 0f3f000000;   	// 0.5
	mul.f32 	%f356, %f354, %f355; 	// 
	add.f32 	%f357, %f356, %f244; 	// 
	mul.f32 	%f358, %f357, %f325; 	// 
	add.f32 	%f359, %f358, %f9;   	// 
	mov.f32 	%f360, %f359;        	// 
	mul.f32 	%f361, %f324, %f245; 	// 
	mul.f32 	%f362, %f333, %f14;  	// 
	mul.f32 	%f363, %f309, %f335; 	// 
	add.f32 	%f364, %f361, %f363; 	// 
	mul.f32 	%f365, %f333, %f364; 	// 
	add.f32 	%f366, %f362, %f365; 	// 
	mov.f32 	%f367, 0f3f000000;   	// 0.5
	mul.f32 	%f368, %f366, %f367; 	// 
	add.f32 	%f369, %f368, %f245; 	// 
	mul.f32 	%f370, %f369, %f325; 	// 
	add.f32 	%f371, %f370, %f10;  	// 
	.loc	3	615	0
	ld.const.f32 	%f372, [damping_v];	// id:5215 damping_v+0x0
	mad.f32 	%f373, %f336, %f372, %f339;	// 
	add.f32 	%f374, %f243, %f373; 	// 
	mad.f32 	%f375, %f349, %f372, %f353;	// 
	add.f32 	%f376, %f244, %f375; 	// 
	mad.f32 	%f377, %f362, %f372, %f365;	// 
	add.f32 	%f378, %f245, %f377; 	// 
	.loc	3	617	0
	ld.const.f32 	%f379, [platform_xmax];	// id:5332 platform_xmax+0x0
	set.ge.u32.f32 	%r164, %f379, %f345;	// 
	neg.s32 	%r165, %r164;        	// 
	ld.const.f32 	%f380, [platform_xmin];	// id:5333 platform_xmin+0x0
	set.le.u32.f32 	%r166, %f380, %f345;	// 
	neg.s32 	%r167, %r166;        	// 
	and.b32 	%r168, %r165, %r167; 	// 
	ld.const.f32 	%f381, [platform_zmax];	// id:5330 platform_zmax+0x0
	set.ge.u32.f32 	%r169, %f381, %f371;	// 
	neg.s32 	%r170, %r169;        	// 
	ld.const.f32 	%f382, [platform_zmin];	// id:5331 platform_zmin+0x0
	set.le.u32.f32 	%r171, %f382, %f371;	// 
	neg.s32 	%r172, %r171;        	// 
	and.b32 	%r173, %r170, %r172; 	// 
	and.b32 	%r174, %r168, %r173; 	// 
	mov.f32 	%f383, 0f00000000;   	// 0
	set.ge.u32.f32 	%r175, %f9, %f383;	// 
	neg.s32 	%r176, %r175;        	// 
	mov.f32 	%f384, 0f00000000;   	// 0
	set.le.u32.f32 	%r177, %f359, %f384;	// 
	neg.s32 	%r178, %r177;        	// 
	and.b32 	%r179, %r176, %r178; 	// 
	mov.s32 	%r180, 0;            	// 
	setp.ne.u32 	%p27, %r179, %r180;	// 
	selp.s32 	%r181, 1, 0, %p27;  	// 
	mov.s32 	%r182, 0;            	// 
	set.ne.u32.s32 	%r183, %r174, %r182;	// 
	neg.s32 	%r184, %r183;        	// 
	and.b32 	%r185, %r181, %r184; 	// 
	mov.u32 	%r186, 0;            	// 
	setp.eq.s32 	%p28, %r185, %r186;	// 
	@%p28 bra 	$Lt_2_276;         	// 
	.loc	20	1328	0
	sub.f32 	%f385, %f15, %f348;  	// 
	mul.f32 	%f386, %f276, %f307; 	// 
	sub.f32 	%f387, %f385, %f386; 	// 
	mov.f32 	%f388, 0f00000000;   	// 0
	min.f32 	%f389, %f387, %f388; 	// 
	neg.f32 	%f390, %f389;        	// 
	mov.f32 	%f391, 0f3d23d70a;   	// 0.04
	mul.f32 	%f392, %f390, %f391; 	// 
	mul.f32 	%f246, %f325, %f392; 	// 
	mov.f32 	%f249, %f326;        	// 
	.loc	20	1250	0
	abs.f32 	%f393, %f326;        	// 
	mov.f32 	%f394, 0f7e800000;   	// 8.50706e+37
	setp.gt.f32 	%p29, %f393, %f394;	// 
	@!%p29 bra 	$Lt_2_278;        	// 
	.loc	20	1251	0
	mov.f32 	%f395, 0f3e800000;   	// 0.25
	mul.f32 	%f246, %f246, %f395; 	// 
	.loc	20	1252	0
	mov.f32 	%f396, 0f3e800000;   	// 0.25
	mul.f32 	%f249, %f326, %f396; 	// 
$Lt_2_278:
	.loc	3	631	0
	div.f32 	%f397, %f246, %f249; 	// 
	mul.f32 	%f398, %f374, %f374; 	// 
	mul.f32 	%f399, %f378, %f378; 	// 
	add.f32 	%f400, %f398, %f399; 	// 
	sqrt.f32 	%f401, %f400;       	// 
	setp.ge.f32 	%p30, %f397, %f401;	// 
	@!%p30 bra 	$Lt_2_281;        	// 
	mov.f32 	%f378, 0f00000000;   	// 0
	mov.f32 	%f376, 0f00000000;   	// 0
	mov.f32 	%f374, 0f00000000;   	// 0
	bra.uni 	$Lt_2_280;           	// 
$Lt_2_281:
	.loc	3	66	0
	rsqrt.f32 	%f402, %f400;      	// 
	mul.f32 	%f403, %f402, %f378; 	// 
	mul.f32 	%f404, %f397, %f403; 	// 
	sub.f32 	%f405, %f378, %f404; 	// 
	.loc	3	67	0
	mul.f32 	%f406, %f402, %f374; 	// 
	mul.f32 	%f407, %f397, %f406; 	// 
	sub.f32 	%f374, %f374, %f407; 	// 
	mov.f32 	%f378, %f405;        	// 
	.loc	3	635	0
	mov.f32 	%f376, 0f00000000;   	// 0
$Lt_2_280:
	mov.f32 	%f360, 0f00000000;   	// 0
$Lt_2_276:
	.loc	3	638	0
	mul.lo.u32 	%r187, %r5, 48;   	// 
	ld.param.u32 	%r188, [__cudaparm__Z12pass_unifiedP13CUDA_Vtx_DataPfS1__vtx_data_out];	// id:5334 __cudaparm__Z12pass_unifiedP13CUDA_Vtx_DataPfS1__vtx_data_out+0x0
	add.u32 	%r189, %r188, %r187; 	// 
	st.global.v2.f32 	[%r189+0], {%f305,%f307};	// 
	st.global.f32 	[%r189+8], %f309;	// id:5337
	st.global.v2.f32 	[%r189+16], {%f374,%f376};	// 
	.loc	3	639	0
	st.global.f32 	[%r189+24], %f378;	// id:5340
	st.global.v2.f32 	[%r189+32], {%f345,%f360};	// 
	.loc	3	640	0
	st.global.f32 	[%r189+40], %f371;	// id:5343
$LBB66__Z12pass_unifiedP13CUDA_Vtx_DataPfS1_:
	.loc	3	641	0
	exit;                         	// 
$LDWend__Z12pass_unifiedP13CUDA_Vtx_DataPfS1_:
	} // _Z12pass_unifiedP13CUDA_Vtx_DataPfS1_