Homec4science
Diffusion lammps (efficient_neuronet) Browse

lammps/lib/gpu/gayberne.ptx
6b1774dc2052efficient_neuronet

Last Change
Raw File

gayberne.ptx
View Options

.version 2.3
.target sm_20
.address_size 64
// compiled with /usr/local/cuda/open64/lib//be
// nvopencc 4.0 built on 2011-05-12
//-----------------------------------------------------------
// Compiling /tmp/tmpxft_00009b6f_00000000-9_lal_gayberne.cpp3.i (/home/sjplimp/ccBI#.YH8XBK)
//-----------------------------------------------------------
//-----------------------------------------------------------
// Options:
//-----------------------------------------------------------
// Target:ptx, ISA:sm_20, Endian:little, Pointer Size:64
// -O3 (Optimization level)
// -g0 (Debug level)
// -m2 (Report advisories)
//-----------------------------------------------------------
.file 1 "<command-line>"
.file 2 "/tmp/tmpxft_00009b6f_00000000-8_lal_gayberne.cudafe2.gpu"
.file 3 "/usr/lib/gcc/x86_64-redhat-linux/4.4.5/include/stddef.h"
.file 4 "/usr/local/cuda/include/crt/device_runtime.h"
.file 5 "/usr/local/cuda/include/host_defines.h"
.file 6 "/usr/local/cuda/include/builtin_types.h"
.file 7 "/usr/local/cuda/include/device_types.h"
.file 8 "/usr/local/cuda/include/driver_types.h"
.file 9 "/usr/local/cuda/include/surface_types.h"
.file 10 "/usr/local/cuda/include/texture_types.h"
.file 11 "/usr/local/cuda/include/vector_types.h"
.file 12 "/usr/local/cuda/include/device_launch_parameters.h"
.file 13 "/usr/local/cuda/include/crt/storage_class.h"
.file 14 "/usr/include/bits/types.h"
.file 15 "/usr/include/time.h"
.file 16 "lal_ellipsoid_extra.h"
.file 17 "lal_gayberne.cu"
.file 18 "/usr/local/cuda/include/common_functions.h"
.file 19 "/usr/local/cuda/include/math_functions.h"
.file 20 "/usr/local/cuda/include/math_constants.h"
.file 21 "/usr/local/cuda/include/device_functions.h"
.file 22 "/usr/local/cuda/include/sm_11_atomic_functions.h"
.file 23 "/usr/local/cuda/include/sm_12_atomic_functions.h"
.file 24 "/usr/local/cuda/include/sm_13_double_functions.h"
.file 25 "/usr/local/cuda/include/sm_20_atomic_functions.h"
.file 26 "/usr/local/cuda/include/sm_20_intrinsics.h"
.file 27 "/usr/local/cuda/include/surface_functions.h"
.file 28 "/usr/local/cuda/include/texture_fetch_functions.h"
.file 29 "/usr/local/cuda/include/math_functions_dbl_ptx3.h"
.entry kernel_ellipsoid (
.param .u64 __cudaparm_kernel_ellipsoid_x_,
.param .u64 __cudaparm_kernel_ellipsoid_q,
.param .u64 __cudaparm_kernel_ellipsoid_shape,
.param .u64 __cudaparm_kernel_ellipsoid_well,
.param .u64 __cudaparm_kernel_ellipsoid_gum,
.param .u64 __cudaparm_kernel_ellipsoid_sig_eps,
.param .s32 __cudaparm_kernel_ellipsoid_ntypes,
.param .u64 __cudaparm_kernel_ellipsoid_lshape,
.param .u64 __cudaparm_kernel_ellipsoid_dev_nbor,
.param .s32 __cudaparm_kernel_ellipsoid_stride,
.param .u64 __cudaparm_kernel_ellipsoid_ans,
.param .s32 __cudaparm_kernel_ellipsoid_astride,
.param .u64 __cudaparm_kernel_ellipsoid_engv,
.param .u64 __cudaparm_kernel_ellipsoid_err_flag,
.param .s32 __cudaparm_kernel_ellipsoid_eflag,
.param .s32 __cudaparm_kernel_ellipsoid_vflag,
.param .s32 __cudaparm_kernel_ellipsoid_inum,
.param .s32 __cudaparm_kernel_ellipsoid_t_per_atom)
{
.reg .u32 %r<67>;
.reg .u64 %rd<83>;
.reg .f32 %f<898>;
.reg .pred %p<35>;
.shared .align 16 .b8 __cuda___cuda_local_var_32950_33_non_const_sp_lj128[16];
.shared .align 4 .b8 __cuda___cuda_local_var_33207_55_non_const_red_acc144[3584];
// __cuda_local_var_32957_10_non_const_f = 48
// __cuda_local_var_32961_10_non_const_tor = 64
// __cuda_local_var_32965_9_non_const_virial = 16
.loc 17 91 0
$LDWbegin_kernel_ellipsoid:
.loc 17 96 0
ld.param.u64 %rd1, [__cudaparm_kernel_ellipsoid_gum];
ldu.global.f32 %f1, [%rd1+12];
.loc 17 97 0
ld.global.f32 %f2, [%rd1+16];
.loc 17 98 0
ld.global.f32 %f3, [%rd1+20];
.loc 17 99 0
ld.global.f32 %f4, [%rd1+24];
st.shared.v4.f32 [__cuda___cuda_local_var_32950_33_non_const_sp_lj128+0], {%f1,%f2,%f3,%f4};
.loc 17 112 0
mov.f32 %f5, 0f00000000; // 0
mov.f32 %f6, %f5;
mov.f32 %f7, 0f00000000; // 0
mov.f32 %f8, %f7;
mov.f32 %f9, 0f00000000; // 0
mov.f32 %f10, %f9;
mov.f32 %f11, 0f00000000; // 0
mov.f32 %f12, %f11;
mov.f32 %f13, 0f00000000; // 0
mov.f32 %f14, %f13;
mov.f32 %f15, 0f00000000; // 0
mov.f32 %f16, %f15;
ld.param.s32 %r1, [__cudaparm_kernel_ellipsoid_t_per_atom];
cvt.s32.u32 %r2, %tid.x;
div.s32 %r3, %r2, %r1;
cvt.s32.u32 %r4, %ntid.x;
div.s32 %r5, %r4, %r1;
cvt.s32.u32 %r6, %ctaid.x;
mul.lo.s32 %r7, %r6, %r5;
add.s32 %r8, %r3, %r7;
ld.param.s32 %r9, [__cudaparm_kernel_ellipsoid_inum];
setp.le.s32 %p1, %r9, %r8;
@%p1 bra $Lt_0_55298;
.loc 17 117 0
cvt.s64.s32 %rd2, %r8;
mul.wide.s32 %rd3, %r8, 4;
ld.param.u64 %rd4, [__cudaparm_kernel_ellipsoid_dev_nbor];
add.u64 %rd5, %rd4, %rd3;
ld.global.s32 %r10, [%rd5+0];
ld.param.s32 %r11, [__cudaparm_kernel_ellipsoid_stride];
cvt.s64.s32 %rd6, %r11;
mul.wide.s32 %rd7, %r11, 4;
add.u64 %rd8, %rd7, %rd5;
ld.global.s32 %r12, [%rd8+0];
.loc 17 120 0
cvt.s64.s32 %rd9, %r10;
mul.wide.s32 %rd10, %r10, 16;
ld.param.u64 %rd11, [__cudaparm_kernel_ellipsoid_x_];
add.u64 %rd12, %rd10, %rd11;
ld.global.v4.f32 {%f17,%f18,%f19,%f20}, [%rd12+0];
.loc 17 123 0
cvt.rzi.ftz.s32.f32 %r13, %f20;
cvt.s64.s32 %rd13, %r13;
mul.wide.s32 %rd14, %r13, 16;
ld.param.u64 %rd15, [__cudaparm_kernel_ellipsoid_shape];
add.u64 %rd16, %rd14, %rd15;
ld.global.v4.f32 {%f21,%f22,%f23,_}, [%rd16+0];
.loc 17 126 0
ld.param.u64 %rd17, [__cudaparm_kernel_ellipsoid_q];
add.u64 %rd18, %rd10, %rd17;
ld.global.v4.f32 {%f24,%f25,%f26,%f27}, [%rd18+0];
.loc 17 129 0
ld.param.u64 %rd19, [__cudaparm_kernel_ellipsoid_well];
add.u64 %rd20, %rd14, %rd19;
ld.global.v4.f32 {%f28,%f29,%f30,_}, [%rd20+0];
.loc 17 130 0
cvt.s32.s64 %r14, %rd6;
sub.s32 %r15, %r1, 1;
and.b32 %r16, %r15, %r2;
add.u64 %rd21, %rd7, %rd8;
mul.lo.s32 %r17, %r14, %r16;
cvt.s64.s32 %rd22, %r17;
mul.wide.s32 %rd23, %r17, 4;
add.u64 %rd24, %rd21, %rd23;
mov.s64 %rd25, %rd24;
mul.lo.s32 %r18, %r14, %r12;
cvt.s64.s32 %rd26, %r18;
mul.wide.s32 %rd27, %r18, 4;
add.u64 %rd28, %rd21, %rd27;
setp.ge.u64 %p2, %rd24, %rd28;
@%p2 bra $Lt_0_56834;
ld.param.s32 %r19, [__cudaparm_kernel_ellipsoid_eflag];
mov.s32 %r20, 0;
setp.gt.s32 %p3, %r19, %r20;
ld.param.s32 %r21, [__cudaparm_kernel_ellipsoid_vflag];
mov.s32 %r22, 0;
setp.gt.s32 %p4, %r21, %r22;
add.ftz.f32 %f31, %f25, %f25;
add.ftz.f32 %f32, %f27, %f27;
mul.ftz.f32 %f33, %f24, %f24;
mul.ftz.f32 %f34, %f25, %f25;
mul.ftz.f32 %f35, %f26, %f26;
mul.ftz.f32 %f36, %f27, %f27;
add.ftz.f32 %f37, %f26, %f26;
ld.param.s32 %r23, [__cudaparm_kernel_ellipsoid_ntypes];
mul.lo.s32 %r24, %r23, %r13;
mul.ftz.f32 %f38, %f31, %f26;
mul.ftz.f32 %f39, %f31, %f27;
mul.ftz.f32 %f40, %f31, %f24;
mul.ftz.f32 %f41, %f32, %f24;
add.ftz.f32 %f42, %f33, %f34;
sub.ftz.f32 %f43, %f33, %f34;
mul.ftz.f32 %f44, %f37, %f24;
mul.ftz.f32 %f45, %f37, %f27;
sub.ftz.f32 %f46, %f38, %f41;
add.ftz.f32 %f47, %f38, %f41;
sub.ftz.f32 %f48, %f42, %f35;
add.ftz.f32 %f49, %f35, %f43;
sub.ftz.f32 %f50, %f43, %f35;
add.ftz.f32 %f51, %f39, %f44;
sub.ftz.f32 %f52, %f39, %f44;
sub.ftz.f32 %f53, %f45, %f40;
add.ftz.f32 %f54, %f40, %f45;
ld.param.u64 %rd29, [__cudaparm_kernel_ellipsoid_lshape];
mul.lo.u64 %rd30, %rd13, 4;
add.u64 %rd31, %rd29, %rd30;
mul.ftz.f32 %f55, %f46, %f22;
mul.ftz.f32 %f56, %f46, %f29;
mul.ftz.f32 %f57, %f47, %f21;
mul.ftz.f32 %f58, %f47, %f28;
sub.ftz.f32 %f59, %f48, %f36;
sub.ftz.f32 %f60, %f49, %f36;
add.ftz.f32 %f61, %f36, %f50;
mul.ftz.f32 %f62, %f51, %f23;
mul.ftz.f32 %f63, %f51, %f30;
add.ftz.f32 %f64, %f51, %f51;
mul.ftz.f32 %f65, %f52, %f21;
mul.ftz.f32 %f66, %f52, %f28;
mul.ftz.f32 %f67, %f53, %f23;
mul.ftz.f32 %f68, %f53, %f30;
add.ftz.f32 %f69, %f53, %f53;
mul.ftz.f32 %f70, %f54, %f22;
mul.ftz.f32 %f71, %f54, %f29;
mul.ftz.f32 %f72, %f46, %f55;
mul.ftz.f32 %f73, %f54, %f55;
mul.ftz.f32 %f74, %f46, %f56;
mul.ftz.f32 %f75, %f54, %f56;
mul.ftz.f32 %f76, %f59, %f21;
mul.ftz.f32 %f77, %f59, %f28;
mul.ftz.f32 %f78, %f60, %f22;
mul.ftz.f32 %f79, %f55, %f60;
mul.ftz.f32 %f80, %f60, %f29;
mul.ftz.f32 %f81, %f56, %f60;
mul.ftz.f32 %f82, %f61, %f23;
mul.ftz.f32 %f83, %f61, %f30;
add.ftz.f32 %f84, %f61, %f61;
mul.ftz.f32 %f85, %f46, %f70;
mul.ftz.f32 %f86, %f60, %f70;
mul.ftz.f32 %f87, %f54, %f70;
mul.ftz.f32 %f88, %f46, %f71;
mul.ftz.f32 %f89, %f60, %f71;
mul.ftz.f32 %f90, %f54, %f71;
fma.rn.ftz.f32 %f91, %f59, %f76, %f72;
fma.rn.ftz.f32 %f92, %f76, %f52, %f73;
fma.rn.ftz.f32 %f93, %f59, %f77, %f74;
fma.rn.ftz.f32 %f94, %f77, %f52, %f75;
mul.ftz.f32 %f95, %f46, %f78;
mul.ftz.f32 %f96, %f60, %f78;
mul.ftz.f32 %f97, %f54, %f78;
fma.rn.ftz.f32 %f98, %f76, %f47, %f79;
mul.ftz.f32 %f99, %f46, %f80;
mul.ftz.f32 %f100, %f60, %f80;
mul.ftz.f32 %f101, %f54, %f80;
fma.rn.ftz.f32 %f102, %f77, %f47, %f81;
fma.rn.ftz.f32 %f103, %f59, %f65, %f85;
fma.rn.ftz.f32 %f104, %f47, %f65, %f86;
fma.rn.ftz.f32 %f105, %f52, %f65, %f87;
fma.rn.ftz.f32 %f106, %f59, %f66, %f88;
fma.rn.ftz.f32 %f107, %f47, %f66, %f89;
fma.rn.ftz.f32 %f108, %f52, %f66, %f90;
fma.rn.ftz.f32 %f109, %f51, %f62, %f91;
fma.rn.ftz.f32 %f110, %f62, %f61, %f92;
fma.rn.ftz.f32 %f111, %f51, %f63, %f93;
fma.rn.ftz.f32 %f112, %f63, %f61, %f94;
fma.rn.ftz.f32 %f113, %f59, %f57, %f95;
fma.rn.ftz.f32 %f114, %f47, %f57, %f96;
fma.rn.ftz.f32 %f115, %f57, %f52, %f97;
fma.rn.ftz.f32 %f116, %f62, %f53, %f98;
fma.rn.ftz.f32 %f117, %f59, %f58, %f99;
fma.rn.ftz.f32 %f118, %f47, %f58, %f100;
fma.rn.ftz.f32 %f119, %f58, %f52, %f101;
fma.rn.ftz.f32 %f120, %f63, %f53, %f102;
fma.rn.ftz.f32 %f121, %f51, %f82, %f103;
fma.rn.ftz.f32 %f122, %f53, %f82, %f104;
fma.rn.ftz.f32 %f123, %f61, %f82, %f105;
fma.rn.ftz.f32 %f124, %f51, %f83, %f106;
fma.rn.ftz.f32 %f125, %f53, %f83, %f107;
fma.rn.ftz.f32 %f126, %f61, %f83, %f108;
fma.rn.ftz.f32 %f127, %f51, %f67, %f113;
fma.rn.ftz.f32 %f128, %f53, %f67, %f114;
fma.rn.ftz.f32 %f129, %f67, %f61, %f115;
fma.rn.ftz.f32 %f130, %f51, %f68, %f117;
fma.rn.ftz.f32 %f131, %f53, %f68, %f118;
fma.rn.ftz.f32 %f132, %f68, %f61, %f119;
ld.param.u64 %rd32, [__cudaparm_kernel_ellipsoid_sig_eps];
mov.f32 %f133, 0f00000000; // 0
mov.f32 %f134, 0f00000000; // 0
mov.f32 %f135, 0f00000000; // 0
mov.f32 %f136, 0f00000000; // 0
mov.f32 %f137, 0f00000000; // 0
mov.f32 %f138, 0f00000000; // 0
mov.f32 %f139, 0f00000000; // 0
mov.u64 %rd33, __cuda___cuda_local_var_32950_33_non_const_sp_lj128;
$Lt_0_40962:
//<loop> Loop body line 130, nesting depth: 1, estimated iterations: unknown
.loc 17 135 0
ld.global.s32 %r25, [%rd25+0];
.loc 17 136 0
shr.s32 %r26, %r25, 30;
and.b32 %r27, %r26, 3;
cvt.s64.s32 %rd34, %r27;
mul.wide.s32 %rd35, %r27, 4;
add.u64 %rd36, %rd33, %rd35;
ld.shared.f32 %f140, [%rd36+0];
.loc 17 139 0
and.b32 %r28, %r25, 1073741823;
cvt.s64.s32 %rd37, %r28;
mul.wide.s32 %rd38, %r28, 16;
add.u64 %rd39, %rd38, %rd11;
ld.global.v4.f32 {%f141,%f142,%f143,%f144}, [%rd39+0];
.loc 17 153 0
add.u64 %rd40, %rd38, %rd17;
ld.global.v4.f32 {%f145,%f146,%f147,%f148}, [%rd40+0];
.loc 17 162 0
cvt.rzi.ftz.s32.f32 %r29, %f144;
cvt.s64.s32 %rd41, %r29;
mul.wide.s32 %rd42, %r29, 16;
add.u64 %rd43, %rd42, %rd15;
ld.global.v4.f32 {%f149,%f150,%f151,_}, [%rd43+0];
.loc 16 299 0
sub.ftz.f32 %f152, %f141, %f17;
mov.f32 %f153, %f152;
.loc 16 300 0
add.ftz.f32 %f154, %f146, %f146;
add.ftz.f32 %f155, %f148, %f148;
mul.ftz.f32 %f156, %f145, %f145;
mul.ftz.f32 %f157, %f146, %f146;
mul.ftz.f32 %f158, %f147, %f147;
mul.ftz.f32 %f159, %f148, %f148;
add.ftz.f32 %f160, %f147, %f147;
mul.ftz.f32 %f161, %f154, %f147;
mul.ftz.f32 %f162, %f154, %f148;
mul.ftz.f32 %f163, %f155, %f145;
add.ftz.f32 %f164, %f156, %f157;
mul.ftz.f32 %f165, %f160, %f145;
sub.ftz.f32 %f166, %f161, %f163;
sub.ftz.f32 %f167, %f164, %f158;
add.ftz.f32 %f168, %f162, %f165;
mul.ftz.f32 %f169, %f166, %f150;
sub.ftz.f32 %f170, %f167, %f159;
mul.ftz.f32 %f171, %f168, %f151;
mul.ftz.f32 %f172, %f166, %f169;
mul.ftz.f32 %f173, %f170, %f149;
fma.rn.ftz.f32 %f174, %f170, %f173, %f172;
fma.rn.ftz.f32 %f175, %f168, %f171, %f174;
add.ftz.f32 %f176, %f109, %f175;
mov.f32 %f177, %f176;
.loc 16 301 0
mul.ftz.f32 %f178, %f154, %f145;
sub.ftz.f32 %f179, %f156, %f157;
mul.ftz.f32 %f180, %f160, %f148;
add.ftz.f32 %f181, %f161, %f163;
add.ftz.f32 %f182, %f158, %f179;
sub.ftz.f32 %f183, %f180, %f178;
mul.ftz.f32 %f184, %f181, %f149;
sub.ftz.f32 %f185, %f182, %f159;
mul.ftz.f32 %f186, %f183, %f151;
mul.ftz.f32 %f187, %f185, %f150;
mul.ftz.f32 %f188, %f166, %f187;
fma.rn.ftz.f32 %f189, %f170, %f184, %f188;
fma.rn.ftz.f32 %f190, %f168, %f186, %f189;
add.ftz.f32 %f191, %f127, %f190;
mov.f32 %f192, %f191;
.loc 16 302 0
sub.ftz.f32 %f193, %f179, %f158;
sub.ftz.f32 %f194, %f162, %f165;
add.ftz.f32 %f195, %f178, %f180;
add.ftz.f32 %f196, %f159, %f193;
mul.ftz.f32 %f197, %f194, %f149;
mul.ftz.f32 %f198, %f195, %f150;
mul.ftz.f32 %f199, %f196, %f151;
mul.ftz.f32 %f200, %f166, %f198;
fma.rn.ftz.f32 %f201, %f170, %f197, %f200;
fma.rn.ftz.f32 %f202, %f168, %f199, %f201;
add.ftz.f32 %f203, %f121, %f202;
mov.f32 %f204, %f203;
.loc 16 303 0
sub.ftz.f32 %f205, %f142, %f18;
mov.f32 %f206, %f205;
.loc 16 304 0
mul.ftz.f32 %f207, %f169, %f185;
fma.rn.ftz.f32 %f208, %f173, %f181, %f207;
fma.rn.ftz.f32 %f209, %f171, %f183, %f208;
add.ftz.f32 %f210, %f116, %f209;
mov.f32 %f211, %f210;
.loc 16 305 0
mul.ftz.f32 %f212, %f185, %f187;
fma.rn.ftz.f32 %f213, %f181, %f184, %f212;
fma.rn.ftz.f32 %f214, %f183, %f186, %f213;
add.ftz.f32 %f215, %f128, %f214;
mov.f32 %f216, %f215;
.loc 16 306 0
mul.ftz.f32 %f217, %f185, %f198;
fma.rn.ftz.f32 %f218, %f181, %f197, %f217;
fma.rn.ftz.f32 %f219, %f183, %f199, %f218;
add.ftz.f32 %f220, %f122, %f219;
mov.f32 %f221, %f220;
.loc 16 307 0
sub.ftz.f32 %f222, %f143, %f19;
mov.f32 %f223, %f222;
.loc 16 308 0
mul.ftz.f32 %f224, %f195, %f169;
fma.rn.ftz.f32 %f225, %f173, %f194, %f224;
fma.rn.ftz.f32 %f226, %f171, %f196, %f225;
add.ftz.f32 %f227, %f110, %f226;
mov.f32 %f228, %f227;
.loc 16 309 0
mul.ftz.f32 %f229, %f195, %f187;
fma.rn.ftz.f32 %f230, %f184, %f194, %f229;
fma.rn.ftz.f32 %f231, %f186, %f196, %f230;
add.ftz.f32 %f232, %f129, %f231;
mov.f32 %f233, %f232;
.loc 16 310 0
mul.ftz.f32 %f234, %f195, %f198;
fma.rn.ftz.f32 %f235, %f194, %f197, %f234;
fma.rn.ftz.f32 %f236, %f196, %f199, %f235;
add.ftz.f32 %f237, %f123, %f236;
mov.f32 %f238, %f237;
abs.ftz.f32 %f239, %f210;
abs.ftz.f32 %f240, %f176;
setp.gt.ftz.f32 %p5, %f239, %f240;
@!%p5 bra $Lt_0_41218;
.loc 16 314 0
mov.f32 %f177, %f210;
mov.f32 %f211, %f176;
.loc 16 315 0
mov.f32 %f192, %f215;
mov.f32 %f216, %f191;
.loc 16 316 0
mov.f32 %f204, %f220;
mov.f32 %f221, %f203;
.loc 16 317 0
mov.f32 %f153, %f205;
mov.f32 %f206, %f152;
$Lt_0_41218:
mov.f32 %f241, %f177;
abs.ftz.f32 %f242, %f241;
abs.ftz.f32 %f243, %f227;
setp.lt.ftz.f32 %p6, %f242, %f243;
@!%p6 bra $Lt_0_41730;
.loc 16 321 0
mov.f32 %f177, %f227;
mov.f32 %f228, %f241;
.loc 16 322 0
mov.f32 %f244, %f192;
mov.f32 %f192, %f232;
mov.f32 %f233, %f244;
.loc 16 323 0
mov.f32 %f245, %f204;
mov.f32 %f204, %f237;
mov.f32 %f238, %f245;
.loc 16 324 0
mov.f32 %f246, %f153;
mov.f32 %f153, %f222;
mov.f32 %f223, %f246;
$Lt_0_41730:
mov.f32 %f247, %f177;
mov.f32 %f248, 0f00000000; // 0
setp.neu.ftz.f32 %p7, %f247, %f248;
@!%p7 bra $Lt_0_42498;
bra.uni $Lt_0_43266;
$Lt_0_42498:
mov.f32 %f249, 0f00000000; // 0
setp.neu.ftz.f32 %p8, %f211, %f249;
@!%p8 bra $Lt_0_43010;
.loc 16 338 0
mov.f32 %f177, %f211;
mov.f32 %f211, %f247;
.loc 16 339 0
mov.f32 %f250, %f192;
mov.f32 %f192, %f216;
mov.f32 %f216, %f250;
.loc 16 340 0
mov.f32 %f251, %f204;
mov.f32 %f204, %f221;
mov.f32 %f221, %f251;
.loc 16 341 0
mov.f32 %f252, %f153;
mov.f32 %f153, %f206;
mov.f32 %f206, %f252;
bra.uni $Lt_0_43266;
$Lt_0_43010:
mov.f32 %f253, 0f00000000; // 0
setp.neu.ftz.f32 %p9, %f228, %f253;
@!%p9 bra $Lt_0_43522;
.loc 16 346 0
mov.f32 %f177, %f228;
mov.f32 %f228, %f247;
.loc 16 347 0
mov.f32 %f254, %f192;
mov.f32 %f192, %f233;
mov.f32 %f233, %f254;
.loc 16 348 0
mov.f32 %f255, %f204;
mov.f32 %f204, %f238;
mov.f32 %f238, %f255;
.loc 16 349 0
mov.f32 %f256, %f153;
mov.f32 %f153, %f223;
mov.f32 %f223, %f256;
bra.uni $Lt_0_43266;
$Lt_0_43522:
.loc 16 352 0
mov.s32 %r30, 2;
ld.param.u64 %rd44, [__cudaparm_kernel_ellipsoid_err_flag];
st.global.s32 [%rd44+0], %r30;
$Lt_0_43266:
$Lt_0_42754:
$Lt_0_42242:
.loc 16 355 0
div.approx.ftz.f32 %f257, %f211, %f177;
mul.ftz.f32 %f258, %f192, %f257;
sub.ftz.f32 %f259, %f216, %f258;
mov.f32 %f216, %f259;
.loc 16 356 0
mul.ftz.f32 %f260, %f204, %f257;
sub.ftz.f32 %f261, %f221, %f260;
mov.f32 %f221, %f261;
.loc 16 357 0
mul.ftz.f32 %f262, %f153, %f257;
sub.ftz.f32 %f263, %f206, %f262;
mov.f32 %f206, %f263;
.loc 16 359 0
div.approx.ftz.f32 %f264, %f228, %f177;
mul.ftz.f32 %f265, %f192, %f264;
sub.ftz.f32 %f233, %f233, %f265;
.loc 16 360 0
mul.ftz.f32 %f266, %f204, %f264;
sub.ftz.f32 %f238, %f238, %f266;
.loc 16 361 0
mul.ftz.f32 %f267, %f153, %f264;
sub.ftz.f32 %f223, %f223, %f267;
abs.ftz.f32 %f268, %f259;
abs.ftz.f32 %f269, %f233;
setp.lt.ftz.f32 %p10, %f268, %f269;
@!%p10 bra $Lt_0_43778;
.loc 16 366 0
mov.f32 %f216, %f233;
mov.f32 %f233, %f259;
.loc 16 367 0
mov.f32 %f221, %f238;
mov.f32 %f238, %f261;
.loc 16 368 0
mov.f32 %f206, %f223;
mov.f32 %f223, %f263;
$Lt_0_43778:
mov.f32 %f270, %f216;
mov.f32 %f271, 0f00000000; // 0
setp.neu.ftz.f32 %p11, %f270, %f271;
@!%p11 bra $Lt_0_44546;
bra.uni $Lt_0_44802;
$Lt_0_44546:
mov.f32 %f272, 0f00000000; // 0
setp.neu.ftz.f32 %p12, %f233, %f272;
@!%p12 bra $Lt_0_44802;
.loc 16 383 0
mov.f32 %f216, %f233;
mov.f32 %f233, %f270;
.loc 16 384 0
mov.f32 %f273, %f221;
mov.f32 %f221, %f238;
mov.f32 %f238, %f273;
.loc 16 385 0
mov.f32 %f274, %f206;
mov.f32 %f206, %f223;
mov.f32 %f223, %f274;
$Lt_0_44802:
$Lt_0_44290:
.loc 16 390 0
div.approx.ftz.f32 %f275, %f233, %f216;
mul.ftz.f32 %f276, %f221, %f275;
sub.ftz.f32 %f238, %f238, %f276;
.loc 16 391 0
mul.ftz.f32 %f277, %f206, %f275;
sub.ftz.f32 %f223, %f223, %f277;
mov.f32 %f278, 0f00000000; // 0
setp.eq.ftz.f32 %p13, %f238, %f278;
@!%p13 bra $Lt_0_45314;
.loc 16 394 0
mov.s32 %r31, 2;
ld.param.u64 %rd45, [__cudaparm_kernel_ellipsoid_err_flag];
st.global.s32 [%rd45+0], %r31;
$Lt_0_45314:
.loc 17 179 0
div.approx.ftz.f32 %f279, %f223, %f238;
mul.ftz.f32 %f280, %f205, %f205;
mul.ftz.f32 %f281, %f279, %f221;
fma.rn.ftz.f32 %f282, %f152, %f152, %f280;
sub.ftz.f32 %f283, %f206, %f281;
fma.rn.ftz.f32 %f284, %f222, %f222, %f282;
div.approx.ftz.f32 %f285, %f283, %f216;
rsqrt.approx.ftz.f32 %f286, %f284;
mul.ftz.f32 %f287, %f285, %f192;
fma.rn.ftz.f32 %f288, %f204, %f279, %f287;
sub.ftz.f32 %f289, %f153, %f288;
div.approx.ftz.f32 %f290, %f289, %f177;
mul.ftz.f32 %f291, %f286, %f290;
.loc 17 191 0
mul.ftz.f32 %f292, %f285, %f286;
mul.ftz.f32 %f293, %f286, %f205;
mul.ftz.f32 %f294, %f286, %f152;
mul.ftz.f32 %f295, %f286, %f222;
mul.ftz.f32 %f296, %f279, %f286;
mul.ftz.f32 %f297, %f292, %f293;
fma.rn.ftz.f32 %f298, %f294, %f291, %f297;
fma.rn.ftz.f32 %f299, %f295, %f296, %f298;
mov.f32 %f300, 0f3f000000; // 0.5
mul.ftz.f32 %f301, %f299, %f300;
rsqrt.approx.ftz.f32 %f302, %f301;
.loc 17 195 0
rcp.approx.ftz.f32 %f303, %f286;
mul.ftz.f32 %f304, %f303, %f291;
.loc 17 200 0
add.s32 %r32, %r29, %r24;
cvt.s64.s32 %rd46, %r32;
mul.wide.s32 %rd47, %r32, 8;
add.u64 %rd48, %rd32, %rd47;
ld.global.v2.f32 {%f305,%f306}, [%rd48+0];
.loc 17 202 0
sub.ftz.f32 %f307, %f303, %f302;
ld.global.f32 %f308, [%rd1+0];
fma.rn.ftz.f32 %f309, %f308, %f305, %f307;
.loc 17 209 0
div.approx.ftz.f32 %f310, %f305, %f309;
mul.ftz.f32 %f311, %f310, %f310;
mul.ftz.f32 %f312, %f310, %f311;
mul.ftz.f32 %f313, %f312, %f312;
mul.ftz.f32 %f314, %f313, %f313;
mul.ftz.f32 %f315, %f310, %f313;
add.ftz.f32 %f316, %f314, %f314;
mul.ftz.f32 %f317, %f310, %f316;
sub.ftz.f32 %f318, %f317, %f315;
div.approx.ftz.f32 %f319, %f318, %f305;
mov.f32 %f320, 0f41c00000; // 24
mul.ftz.f32 %f321, %f319, %f320;
mul.ftz.f32 %f322, %f306, %f321;
.loc 17 214 0
mul.ftz.f32 %f323, %f302, %f322;
mul.ftz.f32 %f324, %f323, %f302;
mul.ftz.f32 %f325, %f324, %f302;
mov.f32 %f326, 0f3f000000; // 0.5
mul.ftz.f32 %f327, %f325, %f326;
mul.ftz.f32 %f328, %f327, %f286;
mul.ftz.f32 %f329, %f292, %f303;
mul.ftz.f32 %f330, %f296, %f303;
mul.ftz.f32 %f331, %f286, %f328;
mul.ftz.f32 %f332, %f293, %f329;
fma.rn.ftz.f32 %f333, %f294, %f304, %f332;
fma.rn.ftz.f32 %f334, %f295, %f330, %f333;
mul.ftz.f32 %f335, %f294, %f334;
sub.ftz.f32 %f336, %f304, %f335;
mul.ftz.f32 %f337, %f331, %f336;
fma.rn.ftz.f32 %f338, %f294, %f322, %f337;
.loc 17 215 0
mul.ftz.f32 %f339, %f293, %f334;
sub.ftz.f32 %f340, %f329, %f339;
mul.ftz.f32 %f341, %f331, %f340;
fma.rn.ftz.f32 %f342, %f293, %f322, %f341;
.loc 17 216 0
mul.ftz.f32 %f343, %f295, %f334;
sub.ftz.f32 %f344, %f330, %f343;
mul.ftz.f32 %f345, %f331, %f344;
fma.rn.ftz.f32 %f346, %f295, %f322, %f345;
.loc 17 226 0
mul.ftz.f32 %f347, %f122, %f329;
mul.ftz.f32 %f348, %f330, %f331;
mul.ftz.f32 %f349, %f329, %f331;
mul.ftz.f32 %f350, %f329, %f128;
fma.rn.ftz.f32 %f351, %f304, %f121, %f347;
fma.rn.ftz.f32 %f352, %f304, %f127, %f350;
fma.rn.ftz.f32 %f353, %f330, %f123, %f351;
fma.rn.ftz.f32 %f354, %f330, %f129, %f352;
mul.ftz.f32 %f355, %f348, %f354;
neg.ftz.f32 %f356, %f349;
fma.rn.ftz.f32 %f357, %f356, %f353, %f355;
mul.ftz.f32 %f358, %f116, %f329;
mul.ftz.f32 %f359, %f331, %f304;
fma.rn.ftz.f32 %f360, %f109, %f304, %f358;
fma.rn.ftz.f32 %f361, %f330, %f110, %f360;
mul.ftz.f32 %f362, %f359, %f353;
neg.ftz.f32 %f363, %f348;
fma.rn.ftz.f32 %f364, %f361, %f363, %f362;
mul.ftz.f32 %f365, %f349, %f361;
neg.ftz.f32 %f366, %f359;
fma.rn.ftz.f32 %f367, %f366, %f354, %f365;
.loc 17 233 0
ld.global.f32 %f368, [%rd31+0];
mul.lo.u64 %rd49, %rd41, 4;
add.u64 %rd50, %rd29, %rd49;
ld.global.f32 %f369, [%rd50+0];
add.ftz.f32 %f370, %f368, %f368;
mul.ftz.f32 %f371, %f369, %f370;
.loc 17 234 0
mul.ftz.f32 %f372, %f210, %f203;
mul.ftz.f32 %f373, %f227, %f203;
mul.ftz.f32 %f374, %f220, %f176;
mul.ftz.f32 %f375, %f210, %f191;
mul.ftz.f32 %f376, %f227, %f191;
mul.ftz.f32 %f377, %f215, %f176;
mul.ftz.f32 %f378, %f374, %f232;
mul.ftz.f32 %f379, %f237, %f377;
sub.ftz.f32 %f380, %f379, %f378;
mul.ftz.f32 %f381, %f237, %f375;
sub.ftz.f32 %f382, %f380, %f381;
fma.rn.ftz.f32 %f383, %f232, %f372, %f382;
fma.rn.ftz.f32 %f384, %f220, %f376, %f383;
mul.ftz.f32 %f385, %f215, %f373;
sub.ftz.f32 %f386, %f384, %f385;
.loc 17 235 0
ld.global.f32 %f387, [%rd1+4];
.loc 17 240 0
mul.ftz.f32 %f388, %f232, %f372;
sub.ftz.f32 %f389, %f388, %f378;
mul.ftz.f32 %f390, %f215, %f373;
sub.ftz.f32 %f391, %f389, %f390;
fma.rn.ftz.f32 %f392, %f220, %f376, %f391;
mul.ftz.f32 %f393, %f237, %f375;
sub.ftz.f32 %f394, %f392, %f393;
fma.rn.ftz.f32 %f395, %f237, %f377, %f394;
.loc 17 241 0
div.approx.ftz.f32 %f396, %f371, %f386;
lg2.approx.ftz.f32 %f397, %f396;
mul.ftz.f32 %f398, %f397, %f387;
ex2.approx.ftz.f32 %f399, %f398;
mul.ftz.f32 %f400, %f399, %f387;
neg.ftz.f32 %f401, %f400;
.loc 17 274 0
add.u64 %rd51, %rd42, %rd19;
ld.global.v4.f32 {%f402,%f403,%f404,_}, [%rd51+0];
.loc 16 299 0
mul.ftz.f32 %f405, %f294, %f303;
mov.f32 %f153, %f405;
.loc 16 300 0
mul.ftz.f32 %f406, %f166, %f403;
mul.ftz.f32 %f407, %f168, %f404;
mul.ftz.f32 %f408, %f166, %f406;
mul.ftz.f32 %f409, %f170, %f402;
fma.rn.ftz.f32 %f410, %f170, %f409, %f408;
fma.rn.ftz.f32 %f411, %f168, %f407, %f410;
add.ftz.f32 %f412, %f111, %f411;
mov.f32 %f177, %f412;
.loc 16 301 0
mul.ftz.f32 %f413, %f181, %f402;
mul.ftz.f32 %f414, %f183, %f404;
mul.ftz.f32 %f415, %f185, %f403;
mul.ftz.f32 %f416, %f166, %f415;
fma.rn.ftz.f32 %f417, %f170, %f413, %f416;
fma.rn.ftz.f32 %f418, %f168, %f414, %f417;
add.ftz.f32 %f419, %f130, %f418;
mov.f32 %f192, %f419;
.loc 16 302 0
mul.ftz.f32 %f420, %f194, %f402;
mul.ftz.f32 %f421, %f195, %f403;
mul.ftz.f32 %f422, %f196, %f404;
mul.ftz.f32 %f423, %f166, %f421;
fma.rn.ftz.f32 %f424, %f170, %f420, %f423;
fma.rn.ftz.f32 %f425, %f168, %f422, %f424;
add.ftz.f32 %f426, %f124, %f425;
mov.f32 %f204, %f426;
.loc 16 303 0
mul.ftz.f32 %f427, %f293, %f303;
mov.f32 %f206, %f427;
.loc 16 304 0
mul.ftz.f32 %f428, %f406, %f185;
fma.rn.ftz.f32 %f429, %f409, %f181, %f428;
fma.rn.ftz.f32 %f430, %f407, %f183, %f429;
add.ftz.f32 %f431, %f120, %f430;
mov.f32 %f211, %f431;
.loc 16 305 0
mul.ftz.f32 %f432, %f185, %f415;
fma.rn.ftz.f32 %f433, %f181, %f413, %f432;
fma.rn.ftz.f32 %f434, %f183, %f414, %f433;
add.ftz.f32 %f216, %f131, %f434;
.loc 16 306 0
mul.ftz.f32 %f435, %f185, %f421;
fma.rn.ftz.f32 %f436, %f181, %f420, %f435;
fma.rn.ftz.f32 %f437, %f183, %f422, %f436;
add.ftz.f32 %f221, %f125, %f437;
.loc 16 307 0
mul.ftz.f32 %f438, %f295, %f303;
mov.f32 %f223, %f438;
.loc 16 308 0
mul.ftz.f32 %f439, %f195, %f406;
fma.rn.ftz.f32 %f440, %f409, %f194, %f439;
fma.rn.ftz.f32 %f441, %f407, %f196, %f440;
add.ftz.f32 %f442, %f112, %f441;
mov.f32 %f228, %f442;
.loc 16 309 0
mul.ftz.f32 %f443, %f195, %f415;
fma.rn.ftz.f32 %f444, %f413, %f194, %f443;
fma.rn.ftz.f32 %f445, %f414, %f196, %f444;
add.ftz.f32 %f233, %f132, %f445;
.loc 16 310 0
mul.ftz.f32 %f446, %f195, %f421;
fma.rn.ftz.f32 %f447, %f194, %f420, %f446;
fma.rn.ftz.f32 %f448, %f196, %f422, %f447;
add.ftz.f32 %f238, %f126, %f448;
abs.ftz.f32 %f449, %f431;
abs.ftz.f32 %f450, %f412;
setp.gt.ftz.f32 %p14, %f449, %f450;
@!%p14 bra $Lt_0_45826;
.loc 16 314 0
mov.f32 %f177, %f431;
mov.f32 %f211, %f412;
.loc 16 315 0
mov.f32 %f192, %f216;
mov.f32 %f216, %f419;
.loc 16 316 0
mov.f32 %f204, %f221;
mov.f32 %f221, %f426;
.loc 16 317 0
mov.f32 %f153, %f427;
mov.f32 %f206, %f405;
$Lt_0_45826:
mov.f32 %f451, %f177;
abs.ftz.f32 %f452, %f451;
abs.ftz.f32 %f453, %f442;
setp.lt.ftz.f32 %p15, %f452, %f453;
@!%p15 bra $Lt_0_46338;
.loc 16 321 0
mov.f32 %f177, %f442;
mov.f32 %f228, %f451;
.loc 16 322 0
mov.f32 %f454, %f192;
mov.f32 %f192, %f233;
mov.f32 %f233, %f454;
.loc 16 323 0
mov.f32 %f455, %f204;
mov.f32 %f204, %f238;
mov.f32 %f238, %f455;
.loc 16 324 0
mov.f32 %f456, %f153;
mov.f32 %f153, %f438;
mov.f32 %f223, %f456;
$Lt_0_46338:
mov.f32 %f457, %f177;
mov.f32 %f458, 0f00000000; // 0
setp.neu.ftz.f32 %p16, %f457, %f458;
@!%p16 bra $Lt_0_47106;
bra.uni $Lt_0_47874;
$Lt_0_47106:
mov.f32 %f459, 0f00000000; // 0
setp.neu.ftz.f32 %p17, %f211, %f459;
@!%p17 bra $Lt_0_47618;
.loc 16 338 0
mov.f32 %f177, %f211;
mov.f32 %f211, %f457;
.loc 16 339 0
mov.f32 %f460, %f192;
mov.f32 %f192, %f216;
mov.f32 %f216, %f460;
.loc 16 340 0
mov.f32 %f461, %f204;
mov.f32 %f204, %f221;
mov.f32 %f221, %f461;
.loc 16 341 0
mov.f32 %f462, %f153;
mov.f32 %f153, %f206;
mov.f32 %f206, %f462;
bra.uni $Lt_0_47874;
$Lt_0_47618:
mov.f32 %f463, 0f00000000; // 0
setp.neu.ftz.f32 %p18, %f228, %f463;
@!%p18 bra $Lt_0_48130;
.loc 16 346 0
mov.f32 %f177, %f228;
mov.f32 %f228, %f457;
.loc 16 347 0
mov.f32 %f464, %f192;
mov.f32 %f192, %f233;
mov.f32 %f233, %f464;
.loc 16 348 0
mov.f32 %f465, %f204;
mov.f32 %f204, %f238;
mov.f32 %f238, %f465;
.loc 16 349 0
mov.f32 %f466, %f153;
mov.f32 %f153, %f223;
mov.f32 %f223, %f466;
bra.uni $Lt_0_47874;
$Lt_0_48130:
.loc 16 352 0
mov.s32 %r33, 2;
ld.param.u64 %rd52, [__cudaparm_kernel_ellipsoid_err_flag];
st.global.s32 [%rd52+0], %r33;
$Lt_0_47874:
$Lt_0_47362:
$Lt_0_46850:
.loc 16 355 0
div.approx.ftz.f32 %f467, %f211, %f177;
mul.ftz.f32 %f468, %f192, %f467;
sub.ftz.f32 %f469, %f216, %f468;
mov.f32 %f216, %f469;
.loc 16 356 0
mul.ftz.f32 %f470, %f204, %f467;
sub.ftz.f32 %f471, %f221, %f470;
mov.f32 %f221, %f471;
.loc 16 357 0
mul.ftz.f32 %f472, %f153, %f467;
sub.ftz.f32 %f473, %f206, %f472;
mov.f32 %f206, %f473;
.loc 16 359 0
div.approx.ftz.f32 %f474, %f228, %f177;
mul.ftz.f32 %f475, %f192, %f474;
sub.ftz.f32 %f233, %f233, %f475;
.loc 16 360 0
mul.ftz.f32 %f476, %f204, %f474;
sub.ftz.f32 %f238, %f238, %f476;
.loc 16 361 0
mul.ftz.f32 %f477, %f153, %f474;
sub.ftz.f32 %f223, %f223, %f477;
abs.ftz.f32 %f478, %f469;
abs.ftz.f32 %f479, %f233;
setp.lt.ftz.f32 %p19, %f478, %f479;
@!%p19 bra $Lt_0_48386;
.loc 16 366 0
mov.f32 %f216, %f233;
mov.f32 %f233, %f469;
.loc 16 367 0
mov.f32 %f221, %f238;
mov.f32 %f238, %f471;
.loc 16 368 0
mov.f32 %f206, %f223;
mov.f32 %f223, %f473;
$Lt_0_48386:
mov.f32 %f480, %f216;
mov.f32 %f481, 0f00000000; // 0
setp.neu.ftz.f32 %p20, %f480, %f481;
@!%p20 bra $Lt_0_49154;
bra.uni $Lt_0_49410;
$Lt_0_49154:
mov.f32 %f482, 0f00000000; // 0
setp.neu.ftz.f32 %p21, %f233, %f482;
@!%p21 bra $Lt_0_49410;
.loc 16 383 0
mov.f32 %f216, %f233;
mov.f32 %f233, %f480;
.loc 16 384 0
mov.f32 %f483, %f221;
mov.f32 %f221, %f238;
mov.f32 %f238, %f483;
.loc 16 385 0
mov.f32 %f484, %f206;
mov.f32 %f206, %f223;
mov.f32 %f223, %f484;
$Lt_0_49410:
$Lt_0_48898:
.loc 16 390 0
div.approx.ftz.f32 %f485, %f233, %f216;
mul.ftz.f32 %f486, %f221, %f485;
sub.ftz.f32 %f238, %f238, %f486;
.loc 16 391 0
mul.ftz.f32 %f487, %f206, %f485;
sub.ftz.f32 %f223, %f223, %f487;
mov.f32 %f488, 0f00000000; // 0
setp.eq.ftz.f32 %p22, %f238, %f488;
@!%p22 bra $Lt_0_49922;
.loc 16 394 0
mov.s32 %r34, 2;
ld.param.u64 %rd53, [__cudaparm_kernel_ellipsoid_err_flag];
st.global.s32 [%rd53+0], %r34;
$Lt_0_49922:
.loc 17 286 0
div.approx.ftz.f32 %f489, %f223, %f238;
mul.ftz.f32 %f490, %f489, %f221;
sub.ftz.f32 %f491, %f206, %f490;
div.approx.ftz.f32 %f492, %f491, %f216;
mul.ftz.f32 %f493, %f492, %f192;
fma.rn.ftz.f32 %f494, %f204, %f489, %f493;
sub.ftz.f32 %f495, %f153, %f494;
div.approx.ftz.f32 %f496, %f495, %f177;
mul.ftz.f32 %f497, %f286, %f496;
.loc 17 293 0
mul.ftz.f32 %f498, %f492, %f286;
mul.ftz.f32 %f499, %f489, %f286;
mul.ftz.f32 %f500, %f286, %f405;
mul.ftz.f32 %f501, %f286, %f427;
mul.ftz.f32 %f502, %f286, %f438;
mul.ftz.f32 %f503, %f498, %f501;
fma.rn.ftz.f32 %f504, %f500, %f497, %f503;
fma.rn.ftz.f32 %f505, %f502, %f499, %f504;
add.ftz.f32 %f506, %f505, %f505;
ld.global.f32 %f507, [%rd1+8];
.loc 17 296 0
mul.ftz.f32 %f508, %f303, %f497;
.loc 17 301 0
mov.f32 %f509, 0fbf800000; // -1
add.ftz.f32 %f510, %f507, %f509;
lg2.approx.ftz.f32 %f511, %f506;
mul.ftz.f32 %f512, %f511, %f507;
ex2.approx.ftz.f32 %f513, %f512;
mov.f32 %f514, 0fc0800000; // -4
mul.ftz.f32 %f515, %f286, %f514;
mul.ftz.f32 %f516, %f286, %f515;
lg2.approx.ftz.f32 %f517, %f513;
div.approx.ftz.f32 %f518, %f510, %f507;
mul.ftz.f32 %f519, %f517, %f518;
ex2.approx.ftz.f32 %f520, %f519;
mul.ftz.f32 %f521, %f516, %f507;
mul.ftz.f32 %f522, %f520, %f521;
.loc 17 303 0
mul.ftz.f32 %f523, %f498, %f303;
mul.ftz.f32 %f524, %f499, %f303;
mul.ftz.f32 %f525, %f523, %f501;
fma.rn.ftz.f32 %f526, %f500, %f508, %f525;
fma.rn.ftz.f32 %f527, %f502, %f524, %f526;
mul.ftz.f32 %f528, %f500, %f527;
sub.ftz.f32 %f529, %f508, %f528;
mul.ftz.f32 %f530, %f522, %f529;
.loc 17 304 0
mul.ftz.f32 %f531, %f501, %f527;
sub.ftz.f32 %f532, %f523, %f531;
mul.ftz.f32 %f533, %f522, %f532;
.loc 17 305 0
mul.ftz.f32 %f534, %f502, %f527;
sub.ftz.f32 %f535, %f524, %f534;
mul.ftz.f32 %f536, %f522, %f535;
.loc 17 310 0
mul.ftz.f32 %f537, %f125, %f523;
mul.ftz.f32 %f538, %f523, %f131;
fma.rn.ftz.f32 %f539, %f508, %f124, %f537;
fma.rn.ftz.f32 %f540, %f508, %f130, %f538;
fma.rn.ftz.f32 %f541, %f524, %f126, %f539;
fma.rn.ftz.f32 %f542, %f524, %f132, %f540;
mul.ftz.f32 %f543, %f523, %f541;
mul.ftz.f32 %f544, %f542, %f524;
sub.ftz.f32 %f545, %f544, %f543;
mul.ftz.f32 %f546, %f120, %f523;
fma.rn.ftz.f32 %f547, %f111, %f508, %f546;
fma.rn.ftz.f32 %f548, %f524, %f112, %f547;
mul.ftz.f32 %f549, %f524, %f548;
mul.ftz.f32 %f550, %f508, %f541;
sub.ftz.f32 %f551, %f550, %f549;
mul.ftz.f32 %f552, %f542, %f508;
mul.ftz.f32 %f553, %f548, %f523;
sub.ftz.f32 %f554, %f553, %f552;
.loc 17 312 0
mul.ftz.f32 %f555, %f516, %f545;
.loc 17 313 0
mul.ftz.f32 %f556, %f516, %f551;
.loc 17 314 0
mul.ftz.f32 %f557, %f516, %f554;
.loc 16 396 0
mov.f32 %f558, 0f40800000; // 4
mul.ftz.f32 %f559, %f306, %f558;
mul.ftz.f32 %f560, %f399, %f140;
sub.ftz.f32 %f561, %f314, %f313;
mul.ftz.f32 %f562, %f513, %f560;
mul.ftz.f32 %f563, %f559, %f561;
fma.rn.ftz.f32 %f564, %f563, %f562, %f139;
selp.f32 %f139, %f564, %f139, %p3;
mul.ftz.f32 %f565, %f562, %f338;
mul.ftz.f32 %f566, %f562, %f342;
mul.ftz.f32 %f567, %f562, %f346;
mul.ftz.f32 %f568, %f399, %f563;
mul.ftz.f32 %f569, %f568, %f140;
neg.ftz.f32 %f570, %f569;
mul.ftz.f32 %f571, %f530, %f570;
sub.ftz.f32 %f572, %f571, %f565;
mul.ftz.f32 %f573, %f533, %f570;
sub.ftz.f32 %f574, %f573, %f566;
mul.ftz.f32 %f575, %f536, %f570;
sub.ftz.f32 %f576, %f575, %f567;
@!%p4 bra $Lt_0_50690;
.loc 17 326 0
add.ftz.f32 %f138, %f572, %f138;
.loc 17 327 0
mul.ftz.f32 %f577, %f303, %f500;
neg.ftz.f32 %f578, %f577;
mov.f32 %f579, %f6;
fma.rn.ftz.f32 %f580, %f578, %f572, %f579;
mov.f32 %f6, %f580;
.loc 17 329 0
add.ftz.f32 %f137, %f574, %f137;
.loc 17 330 0
mul.ftz.f32 %f581, %f303, %f501;
neg.ftz.f32 %f582, %f581;
mov.f32 %f583, %f8;
fma.rn.ftz.f32 %f584, %f582, %f574, %f583;
mov.f32 %f8, %f584;
.loc 17 331 0
mov.f32 %f585, %f12;
fma.rn.ftz.f32 %f586, %f578, %f574, %f585;
mov.f32 %f12, %f586;
.loc 17 333 0
add.ftz.f32 %f136, %f576, %f136;
.loc 17 334 0
mov.f32 %f587, %f10;
mul.ftz.f32 %f588, %f303, %f502;
neg.ftz.f32 %f589, %f588;
fma.rn.ftz.f32 %f590, %f589, %f576, %f587;
mov.f32 %f10, %f590;
.loc 17 335 0
mov.f32 %f591, %f14;
fma.rn.ftz.f32 %f592, %f578, %f576, %f591;
mov.f32 %f14, %f592;
.loc 17 336 0
fma.rn.ftz.f32 %f15, %f582, %f576, %f15;
mov.f32 %f16, %f15;
bra.uni $Lt_0_50434;
$Lt_0_50690:
.loc 17 338 0
add.ftz.f32 %f138, %f572, %f138;
.loc 17 339 0
add.ftz.f32 %f137, %f574, %f137;
.loc 17 340 0
add.ftz.f32 %f136, %f576, %f136;
$Lt_0_50434:
.loc 17 347 0
rcp.approx.ftz.f32 %f593, %f395;
mul.ftz.f32 %f594, %f513, %f399;
mul.ftz.f32 %f595, %f594, %f140;
neg.ftz.f32 %f596, %f595;
mul.ftz.f32 %f597, %f513, %f563;
mul.ftz.f32 %f598, %f54, %f176;
mul.ftz.f32 %f599, %f60, %f176;
add.ftz.f32 %f600, %f176, %f176;
mul.ftz.f32 %f601, %f52, %f176;
mul.ftz.f32 %f602, %f47, %f176;
mul.ftz.f32 %f603, %f69, %f176;
mul.ftz.f32 %f604, %f61, %f176;
add.ftz.f32 %f605, %f227, %f227;
mul.ftz.f32 %f606, %f46, %f227;
mul.ftz.f32 %f607, %f59, %f227;
mul.ftz.f32 %f608, %f52, %f227;
mul.ftz.f32 %f609, %f47, %f227;
mul.ftz.f32 %f610, %f54, %f210;
add.ftz.f32 %f611, %f210, %f210;
mul.ftz.f32 %f612, %f46, %f210;
mul.ftz.f32 %f613, %f52, %f210;
mul.ftz.f32 %f614, %f51, %f210;
mul.ftz.f32 %f615, %f84, %f210;
mul.ftz.f32 %f616, %f46, %f203;
mul.ftz.f32 %f617, %f59, %f203;
mul.ftz.f32 %f618, %f51, %f203;
mul.ftz.f32 %f619, %f69, %f203;
mul.ftz.f32 %f620, %f227, %f220;
mul.ftz.f32 %f621, %f61, %f220;
add.ftz.f32 %f622, %f237, %f237;
mul.ftz.f32 %f623, %f237, %f210;
mul.ftz.f32 %f624, %f59, %f237;
mul.ftz.f32 %f625, %f597, %f140;
mul.ftz.f32 %f626, %f600, %f237;
mul.ftz.f32 %f627, %f60, %f605;
mul.ftz.f32 %f628, %f605, %f203;
mul.ftz.f32 %f629, %f237, %f191;
mul.ftz.f32 %f630, %f54, %f191;
mul.ftz.f32 %f631, %f220, %f191;
mul.ftz.f32 %f632, %f61, %f191;
add.ftz.f32 %f633, %f215, %f215;
mul.ftz.f32 %f634, %f227, %f215;
mul.ftz.f32 %f635, %f232, %f210;
mul.ftz.f32 %f636, %f53, %f232;
mul.ftz.f32 %f637, %f611, %f191;
mul.ftz.f32 %f638, %f52, %f611;
mul.ftz.f32 %f639, %f616, %f215;
mul.ftz.f32 %f640, %f617, %f215;
mul.ftz.f32 %f641, %f618, %f232;
mul.ftz.f32 %f642, %f618, %f215;
mul.ftz.f32 %f643, %f622, %f176;
mul.ftz.f32 %f644, %f624, %f191;
neg.ftz.f32 %f645, %f625;
mul.ftz.f32 %f646, %f46, %f629;
mul.ftz.f32 %f647, %f633, %f176;
mul.ftz.f32 %f648, %f61, %f633;
mul.ftz.f32 %f649, %f46, %f631;
sub.ftz.f32 %f650, %f649, %f639;
mul.ftz.f32 %f651, %f59, %f631;
sub.ftz.f32 %f652, %f651, %f640;
mul.ftz.f32 %f653, %f51, %f629;
sub.ftz.f32 %f654, %f653, %f641;
mul.ftz.f32 %f655, %f51, %f631;
sub.ftz.f32 %f656, %f655, %f642;
mul.ftz.f32 %f657, %f232, %f617;
sub.ftz.f32 %f658, %f657, %f644;
mul.ftz.f32 %f659, %f232, %f616;
sub.ftz.f32 %f660, %f659, %f646;
mul.ftz.f32 %f661, %f60, %f374;
sub.ftz.f32 %f662, %f650, %f661;
mul.ftz.f32 %f663, %f47, %f374;
sub.ftz.f32 %f664, %f652, %f663;
mul.ftz.f32 %f665, %f237, %f603;
sub.ftz.f32 %f666, %f654, %f665;
mul.ftz.f32 %f667, %f53, %f374;
sub.ftz.f32 %f668, %f656, %f667;
fma.rn.ftz.f32 %f669, %f47, %f626, %f658;
fma.rn.ftz.f32 %f670, %f60, %f643, %f660;
fma.rn.ftz.f32 %f671, %f60, %f372, %f662;
fma.rn.ftz.f32 %f672, %f47, %f372, %f664;
fma.rn.ftz.f32 %f673, %f176, %f621, %f666;
fma.rn.ftz.f32 %f674, %f53, %f372, %f668;
mul.ftz.f32 %f675, %f220, %f601;
sub.ftz.f32 %f676, %f669, %f675;
mul.ftz.f32 %f677, %f220, %f598;
sub.ftz.f32 %f678, %f670, %f677;
fma.rn.ftz.f32 %f679, %f54, %f647, %f671;
mul.ftz.f32 %f680, %f232, %f602;
sub.ftz.f32 %f681, %f672, %f680;
fma.rn.ftz.f32 %f682, %f227, %f619, %f673;
mul.ftz.f32 %f683, %f51, %f634;
sub.ftz.f32 %f684, %f674, %f683;
mul.ftz.f32 %f685, %f47, %f628;
sub.ftz.f32 %f686, %f676, %f685;
mul.ftz.f32 %f687, %f203, %f627;
sub.ftz.f32 %f688, %f678, %f687;
mul.ftz.f32 %f689, %f232, %f599;
sub.ftz.f32 %f690, %f679, %f689;
mul.ftz.f32 %f691, %f59, %f634;
sub.ftz.f32 %f692, %f681, %f691;
fma.rn.ftz.f32 %f693, %f237, %f614, %f682;
mul.ftz.f32 %f694, %f176, %f636;
sub.ftz.f32 %f695, %f684, %f694;
fma.rn.ftz.f32 %f696, %f203, %f613, %f686;
mul.ftz.f32 %f697, %f46, %f623;
sub.ftz.f32 %f698, %f688, %f697;
fma.rn.ftz.f32 %f699, %f60, %f376, %f690;
fma.rn.ftz.f32 %f700, %f52, %f647, %f692;
mul.ftz.f32 %f701, %f61, %f372;
sub.ftz.f32 %f702, %f693, %f701;
fma.rn.ftz.f32 %f703, %f176, %f648, %f695;
mul.ftz.f32 %f704, %f59, %f623;
sub.ftz.f32 %f705, %f696, %f704;
fma.rn.ftz.f32 %f706, %f46, %f620, %f698;
mul.ftz.f32 %f707, %f215, %f606;
sub.ftz.f32 %f708, %f699, %f707;
mul.ftz.f32 %f709, %f191, %f638;
sub.ftz.f32 %f710, %f700, %f709;
mul.ftz.f32 %f711, %f51, %f620;
sub.ftz.f32 %f712, %f702, %f711;
fma.rn.ftz.f32 %f713, %f51, %f635, %f703;
fma.rn.ftz.f32 %f714, %f220, %f607, %f705;
fma.rn.ftz.f32 %f715, %f203, %f610, %f706;
mul.ftz.f32 %f716, %f54, %f637;
sub.ftz.f32 %f717, %f708, %f716;
fma.rn.ftz.f32 %f718, %f59, %f635, %f710;
fma.rn.ftz.f32 %f719, %f232, %f604, %f712;
fma.rn.ftz.f32 %f720, %f53, %f376, %f713;
fma.rn.ftz.f32 %f721, %f191, %f608, %f714;
mul.ftz.f32 %f722, %f232, %f598;
sub.ftz.f32 %f723, %f715, %f722;
fma.rn.ftz.f32 %f724, %f232, %f612, %f717;
fma.rn.ftz.f32 %f725, %f191, %f609, %f718;
mul.ftz.f32 %f726, %f227, %f632;
sub.ftz.f32 %f727, %f726, %f719;
mul.ftz.f32 %f728, %f191, %f615;
sub.ftz.f32 %f729, %f720, %f728;
mul.ftz.f32 %f730, %f232, %f601;
sub.ftz.f32 %f731, %f721, %f730;
fma.rn.ftz.f32 %f732, %f227, %f630, %f723;
mul.ftz.f32 %f733, %f724, %f22;
mul.ftz.f32 %f734, %f725, %f21;
mul.ftz.f32 %f735, %f727, %f23;
mul.ftz.f32 %f736, %f729, %f23;
mul.ftz.f32 %f737, %f731, %f21;
mul.ftz.f32 %f738, %f732, %f22;
mul.ftz.f32 %f739, %f593, %f733;
mul.ftz.f32 %f740, %f593, %f734;
mul.ftz.f32 %f741, %f593, %f735;
mul.ftz.f32 %f742, %f593, %f736;
mul.ftz.f32 %f743, %f593, %f737;
mul.ftz.f32 %f744, %f593, %f738;
mul.ftz.f32 %f745, %f739, %f401;
mul.ftz.f32 %f746, %f740, %f401;
mul.ftz.f32 %f747, %f741, %f401;
mul.ftz.f32 %f748, %f742, %f401;
mul.ftz.f32 %f749, %f743, %f401;
mul.ftz.f32 %f750, %f744, %f401;
mul.ftz.f32 %f751, %f569, %f555;
mul.ftz.f32 %f752, %f52, %f749;
mul.ftz.f32 %f753, %f47, %f746;
sub.ftz.f32 %f754, %f753, %f752;
mul.ftz.f32 %f755, %f54, %f750;
mul.ftz.f32 %f756, %f745, %f60;
sub.ftz.f32 %f757, %f756, %f755;
add.ftz.f32 %f758, %f754, %f757;
mul.ftz.f32 %f759, %f61, %f747;
mul.ftz.f32 %f760, %f748, %f53;
sub.ftz.f32 %f761, %f760, %f759;
add.ftz.f32 %f762, %f758, %f761;
mul.ftz.f32 %f763, %f762, %f645;
sub.ftz.f32 %f764, %f763, %f751;
fma.rn.ftz.f32 %f765, %f357, %f596, %f764;
add.ftz.f32 %f135, %f135, %f765;
.loc 17 348 0
mul.ftz.f32 %f766, %f54, %f227;
mul.ftz.f32 %f767, %f53, %f210;
mul.ftz.f32 %f768, %f47, %f203;
mul.ftz.f32 %f769, %f60, %f203;
add.ftz.f32 %f770, %f220, %f220;
mul.ftz.f32 %f771, %f47, %f191;
mul.ftz.f32 %f772, %f60, %f191;
mul.ftz.f32 %f773, %f53, %f191;
mul.ftz.f32 %f774, %f52, %f215;
mul.ftz.f32 %f775, %f54, %f215;
mul.ftz.f32 %f776, %f215, %f203;
mul.ftz.f32 %f777, %f232, %f203;
mul.ftz.f32 %f778, %f64, %f232;
mul.ftz.f32 %f779, %f59, %f770;
mul.ftz.f32 %f780, %f46, %f770;
mul.ftz.f32 %f781, %f52, %f631;
mul.ftz.f32 %f782, %f633, %f237;
mul.ftz.f32 %f783, %f51, %f633;
mul.ftz.f32 %f784, %f775, %f203;
mul.ftz.f32 %f785, %f61, %f776;
fma.rn.ftz.f32 %f786, %f59, %f782, %f781;
mul.ftz.f32 %f787, %f46, %f782;
sub.ftz.f32 %f788, %f787, %f784;
mul.ftz.f32 %f789, %f61, %f631;
sub.ftz.f32 %f790, %f789, %f785;
mul.ftz.f32 %f791, %f203, %f774;
sub.ftz.f32 %f792, %f786, %f791;
fma.rn.ftz.f32 %f793, %f54, %f631, %f788;
fma.rn.ftz.f32 %f794, %f237, %f783, %f790;
mul.ftz.f32 %f795, %f232, %f779;
sub.ftz.f32 %f796, %f792, %f795;
mul.ftz.f32 %f797, %f232, %f780;
sub.ftz.f32 %f798, %f793, %f797;
mul.ftz.f32 %f799, %f237, %f773;
sub.ftz.f32 %f800, %f794, %f799;
fma.rn.ftz.f32 %f801, %f232, %f768, %f796;
fma.rn.ftz.f32 %f802, %f232, %f769, %f798;
fma.rn.ftz.f32 %f803, %f53, %f777, %f800;
mul.ftz.f32 %f804, %f237, %f771;
sub.ftz.f32 %f805, %f801, %f804;
mul.ftz.f32 %f806, %f237, %f772;
sub.ftz.f32 %f807, %f802, %f806;
mul.ftz.f32 %f808, %f220, %f778;
sub.ftz.f32 %f809, %f803, %f808;
mul.ftz.f32 %f810, %f47, %f623;
sub.ftz.f32 %f811, %f805, %f810;
mul.ftz.f32 %f812, %f60, %f623;
sub.ftz.f32 %f813, %f807, %f812;
mul.ftz.f32 %f814, %f237, %f767;
sub.ftz.f32 %f815, %f809, %f814;
fma.rn.ftz.f32 %f816, %f47, %f620, %f811;
fma.rn.ftz.f32 %f817, %f60, %f620, %f813;
fma.rn.ftz.f32 %f818, %f53, %f620, %f815;
fma.rn.ftz.f32 %f819, %f232, %f613, %f816;
mul.ftz.f32 %f820, %f215, %f766;
sub.ftz.f32 %f821, %f817, %f820;
mul.ftz.f32 %f822, %f61, %f634;
sub.ftz.f32 %f823, %f818, %f822;
mul.ftz.f32 %f824, %f215, %f608;
sub.ftz.f32 %f825, %f819, %f824;
fma.rn.ftz.f32 %f826, %f232, %f610, %f821;
fma.rn.ftz.f32 %f827, %f61, %f635, %f823;
mul.ftz.f32 %f828, %f825, %f21;
mul.ftz.f32 %f829, %f826, %f22;
mul.ftz.f32 %f830, %f827, %f23;
mul.ftz.f32 %f831, %f593, %f828;
mul.ftz.f32 %f832, %f593, %f829;
mul.ftz.f32 %f833, %f593, %f830;
mul.ftz.f32 %f834, %f831, %f401;
mul.ftz.f32 %f835, %f832, %f401;
mul.ftz.f32 %f836, %f833, %f401;
mul.ftz.f32 %f837, %f569, %f556;
mul.ftz.f32 %f838, %f46, %f745;
mul.ftz.f32 %f839, %f835, %f54;
sub.ftz.f32 %f840, %f839, %f838;
mul.ftz.f32 %f841, %f59, %f746;
mul.ftz.f32 %f842, %f834, %f52;
sub.ftz.f32 %f843, %f842, %f841;
add.ftz.f32 %f844, %f840, %f843;
mul.ftz.f32 %f845, %f51, %f748;
mul.ftz.f32 %f846, %f836, %f61;
sub.ftz.f32 %f847, %f846, %f845;
add.ftz.f32 %f848, %f844, %f847;
mul.ftz.f32 %f849, %f848, %f645;
sub.ftz.f32 %f850, %f849, %f837;
fma.rn.ftz.f32 %f851, %f364, %f596, %f850;
add.ftz.f32 %f134, %f134, %f851;
.loc 17 349 0
mul.ftz.f32 %f852, %f569, %f557;
mul.ftz.f32 %f853, %f47, %f834;
mul.ftz.f32 %f854, %f59, %f749;
sub.ftz.f32 %f855, %f854, %f853;
mul.ftz.f32 %f856, %f60, %f835;
mul.ftz.f32 %f857, %f750, %f46;
sub.ftz.f32 %f858, %f857, %f856;
add.ftz.f32 %f859, %f855, %f858;
mul.ftz.f32 %f860, %f53, %f836;
mul.ftz.f32 %f861, %f747, %f51;
sub.ftz.f32 %f862, %f861, %f860;
add.ftz.f32 %f863, %f859, %f862;
mul.ftz.f32 %f864, %f863, %f645;
sub.ftz.f32 %f865, %f864, %f852;
fma.rn.ftz.f32 %f866, %f367, %f596, %f865;
add.ftz.f32 %f133, %f133, %f866;
mul.lo.s32 %r35, %r14, %r1;
cvt.s64.s32 %rd54, %r35;
mul.wide.s32 %rd55, %r35, 4;
add.u64 %rd25, %rd25, %rd55;
setp.gt.u64 %p23, %rd28, %rd25;
@%p23 bra $Lt_0_40962;
bra.uni $Lt_0_40450;
$Lt_0_56834:
mov.f32 %f133, 0f00000000; // 0
mov.f32 %f134, 0f00000000; // 0
mov.f32 %f135, 0f00000000; // 0
mov.f32 %f136, 0f00000000; // 0
mov.f32 %f137, 0f00000000; // 0
mov.f32 %f138, 0f00000000; // 0
mov.f32 %f139, 0f00000000; // 0
$Lt_0_40450:
mov.u32 %r36, 1;
setp.le.s32 %p24, %r1, %r36;
@%p24 bra $Lt_0_53250;
.loc 17 352 0
mov.u64 %rd56, __cuda___cuda_local_var_33207_55_non_const_red_acc144;
cvt.s64.s32 %rd57, %r2;
mul.wide.s32 %rd58, %r2, 4;
add.u64 %rd59, %rd56, %rd58;
mov.f32 %f867, %f138;
st.shared.f32 [%rd59+0], %f867;
mov.f32 %f868, %f137;
st.shared.f32 [%rd59+512], %f868;
mov.f32 %f869, %f136;
st.shared.f32 [%rd59+1024], %f869;
mov.f32 %f870, %f135;
st.shared.f32 [%rd59+1536], %f870;
mov.f32 %f871, %f134;
st.shared.f32 [%rd59+2048], %f871;
mov.f32 %f872, %f133;
st.shared.f32 [%rd59+2560], %f872;
shr.s32 %r37, %r1, 31;
mov.s32 %r38, 1;
and.b32 %r39, %r37, %r38;
add.s32 %r40, %r39, %r1;
shr.s32 %r41, %r40, 1;
mov.s32 %r42, %r41;
mov.u32 %r43, 0;
setp.ne.u32 %p25, %r41, %r43;
@!%p25 bra $Lt_0_51714;
$Lt_0_52226:
setp.ge.u32 %p26, %r16, %r42;
@%p26 bra $Lt_0_52482;
add.u32 %r44, %r2, %r42;
cvt.u64.u32 %rd60, %r44;
mul.wide.u32 %rd61, %r44, 4;
add.u64 %rd62, %rd56, %rd61;
ld.shared.f32 %f873, [%rd62+0];
add.ftz.f32 %f867, %f873, %f867;
st.shared.f32 [%rd59+0], %f867;
ld.shared.f32 %f874, [%rd62+512];
add.ftz.f32 %f868, %f874, %f868;
st.shared.f32 [%rd59+512], %f868;
ld.shared.f32 %f875, [%rd62+1024];
add.ftz.f32 %f869, %f875, %f869;
st.shared.f32 [%rd59+1024], %f869;
ld.shared.f32 %f876, [%rd62+1536];
add.ftz.f32 %f870, %f876, %f870;
st.shared.f32 [%rd59+1536], %f870;
ld.shared.f32 %f877, [%rd62+2048];
add.ftz.f32 %f871, %f877, %f871;
st.shared.f32 [%rd59+2048], %f871;
ld.shared.f32 %f878, [%rd62+2560];
add.ftz.f32 %f872, %f878, %f872;
st.shared.f32 [%rd59+2560], %f872;
$Lt_0_52482:
shr.u32 %r42, %r42, 1;
mov.u32 %r45, 0;
setp.ne.u32 %p27, %r42, %r45;
@%p27 bra $Lt_0_52226;
$Lt_0_51714:
mov.f32 %f138, %f867;
mov.f32 %f137, %f868;
mov.f32 %f136, %f869;
mov.f32 %f135, %f870;
mov.f32 %f134, %f871;
mov.f32 %f133, %f872;
ld.param.s32 %r46, [__cudaparm_kernel_ellipsoid_eflag];
mov.s32 %r47, 0;
set.gt.u32.s32 %r48, %r46, %r47;
neg.s32 %r49, %r48;
ld.param.s32 %r50, [__cudaparm_kernel_ellipsoid_vflag];
mov.s32 %r51, 0;
set.gt.u32.s32 %r52, %r50, %r51;
neg.s32 %r53, %r52;
or.b32 %r54, %r49, %r53;
mov.u32 %r55, 0;
setp.eq.s32 %p28, %r54, %r55;
@%p28 bra $Lt_0_53250;
mov.f32 %f867, %f6;
st.shared.f32 [%rd59+0], %f867;
mov.f32 %f868, %f8;
st.shared.f32 [%rd59+512], %f868;
mov.f32 %f869, %f10;
st.shared.f32 [%rd59+1024], %f869;
mov.f32 %f870, %f12;
st.shared.f32 [%rd59+1536], %f870;
mov.f32 %f871, %f14;
st.shared.f32 [%rd59+2048], %f871;
mov.f32 %f872, %f15;
st.shared.f32 [%rd59+2560], %f872;
mov.f32 %f879, %f139;
st.shared.f32 [%rd59+3072], %f879;
mov.s32 %r56, %r41;
@!%p25 bra $Lt_0_53762;
$Lt_0_54274:
setp.ge.u32 %p29, %r16, %r56;
@%p29 bra $Lt_0_54530;
add.u32 %r57, %r2, %r56;
cvt.u64.u32 %rd63, %r57;
mul.wide.u32 %rd64, %r57, 4;
add.u64 %rd65, %rd56, %rd64;
ld.shared.f32 %f880, [%rd65+0];
add.ftz.f32 %f867, %f880, %f867;
st.shared.f32 [%rd59+0], %f867;
ld.shared.f32 %f881, [%rd65+512];
add.ftz.f32 %f868, %f881, %f868;
st.shared.f32 [%rd59+512], %f868;
ld.shared.f32 %f882, [%rd65+1024];
add.ftz.f32 %f869, %f882, %f869;
st.shared.f32 [%rd59+1024], %f869;
ld.shared.f32 %f883, [%rd65+1536];
add.ftz.f32 %f870, %f883, %f870;
st.shared.f32 [%rd59+1536], %f870;
ld.shared.f32 %f884, [%rd65+2048];
add.ftz.f32 %f871, %f884, %f871;
st.shared.f32 [%rd59+2048], %f871;
ld.shared.f32 %f885, [%rd65+2560];
add.ftz.f32 %f872, %f885, %f872;
st.shared.f32 [%rd59+2560], %f872;
ld.shared.f32 %f886, [%rd65+3072];
add.ftz.f32 %f879, %f886, %f879;
st.shared.f32 [%rd59+3072], %f879;
$Lt_0_54530:
shr.u32 %r56, %r56, 1;
mov.u32 %r58, 0;
setp.ne.u32 %p30, %r56, %r58;
@%p30 bra $Lt_0_54274;
$Lt_0_53762:
mov.f32 %f6, %f867;
mov.f32 %f8, %f868;
mov.f32 %f10, %f869;
mov.f32 %f12, %f870;
mov.f32 %f14, %f871;
mov.f32 %f16, %f872;
mov.f32 %f139, %f879;
$Lt_0_53250:
$Lt_0_51202:
mov.u32 %r59, 0;
setp.ne.s32 %p31, %r16, %r59;
@%p31 bra $Lt_0_55298;
ld.param.u64 %rd66, [__cudaparm_kernel_ellipsoid_engv];
add.u64 %rd67, %rd66, %rd3;
ld.param.s32 %r60, [__cudaparm_kernel_ellipsoid_astride];
ld.param.s32 %r61, [__cudaparm_kernel_ellipsoid_eflag];
mov.u32 %r62, 0;
setp.le.s32 %p32, %r61, %r62;
@%p32 bra $Lt_0_55810;
st.global.f32 [%rd67+0], %f139;
cvt.s64.s32 %rd68, %r60;
mul.wide.s32 %rd69, %r60, 4;
add.u64 %rd67, %rd67, %rd69;
$Lt_0_55810:
ld.param.s32 %r63, [__cudaparm_kernel_ellipsoid_vflag];
mov.u32 %r64, 0;
setp.le.s32 %p33, %r63, %r64;
@%p33 bra $Lt_0_56322;
mov.f32 %f887, %f6;
st.global.f32 [%rd67+0], %f887;
cvt.s64.s32 %rd70, %r60;
mul.wide.s32 %rd71, %r60, 4;
add.u64 %rd72, %rd71, %rd67;
mov.f32 %f888, %f8;
st.global.f32 [%rd72+0], %f888;
add.u64 %rd73, %rd71, %rd72;
mov.f32 %f889, %f10;
st.global.f32 [%rd73+0], %f889;
add.u64 %rd74, %rd71, %rd73;
mov.f32 %f890, %f12;
st.global.f32 [%rd74+0], %f890;
add.u64 %rd67, %rd71, %rd74;
mov.f32 %f891, %f14;
st.global.f32 [%rd67+0], %f891;
mov.f32 %f892, %f16;
add.u64 %rd75, %rd71, %rd67;
st.global.f32 [%rd75+0], %f892;
$Lt_0_56322:
ld.param.u64 %rd76, [__cudaparm_kernel_ellipsoid_ans];
mul.lo.u64 %rd77, %rd2, 16;
add.u64 %rd78, %rd76, %rd77;
mov.f32 %f893, %f894;
st.global.v4.f32 [%rd78+0], {%f138,%f137,%f136,%f893};
add.s32 %r65, %r8, %r60;
cvt.s64.s32 %rd79, %r65;
mul.wide.s32 %rd80, %r65, 16;
add.u64 %rd81, %rd76, %rd80;
mov.f32 %f895, %f896;
st.global.v4.f32 [%rd81+0], {%f135,%f134,%f133,%f895};
$Lt_0_55298:
$Lt_0_39938:
.loc 17 355 0
exit;
$LDWend_kernel_ellipsoid:
} // kernel_ellipsoid
c4science · Help