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