.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_