// // Generated by NVIDIA NVVM Compiler // // Compiler Build ID: CL-31968024 // Cuda compilation tools, release 12.0, V12.0.76 // Based on NVVM 7.0.1 // .version 8.0 .target sm_52 .address_size 64 .const .align 16 .b8 params[1184]; .visible .func (.param .align 8 .b8 func_retval0[32]) __direct_callable__oxMain( .param .b32 __direct_callable__oxMain_param_0, .param .align 8 .b8 __direct_callable__oxMain_param_1[8] ) { .reg .pred %p<3>; .reg .b16 %rs<9>; .reg .f32 %f<72>; .reg .b32 %r<19>; .reg .b64 %rd<38>; ld.param.u32 %r2, [__direct_callable__oxMain_param_0]; ld.param.f32 %f1, [__direct_callable__oxMain_param_1]; ld.param.f32 %f2, [__direct_callable__oxMain_param_1+4]; ld.const.u64 %rd4, [params+24]; cvta.to.global.u64 %rd5, %rd4; mul.wide.u32 %rd6, %r2, 12; add.s64 %rd7, %rd5, %rd6; ld.global.s32 %rd1, [%rd7]; ld.global.s32 %rd2, [%rd7+4]; ld.global.s32 %rd3, [%rd7+8]; ld.const.u32 %r1, [params+32]; setp.gt.u32 %p1, %r1, %r2; @%p1 bra $L__BB0_2; bra.uni $L__BB0_1; $L__BB0_2: mov.f32 %f71, 0f3F800000; bra.uni $L__BB0_3; $L__BB0_1: sub.s32 %r3, %r2, %r1; ld.const.v2.u64 {%rd8, %rd9}, [params+48]; cvta.to.global.u64 %rd12, %rd8; mul.wide.u32 %rd13, %r3, 4; add.s64 %rd14, %rd12, %rd13; ld.const.u64 %rd15, [params+40]; cvta.to.global.u64 %rd16, %rd15; shl.b64 %rd17, %rd1, 3; add.s64 %rd18, %rd16, %rd17; ld.global.v2.f32 {%f5, %f6}, [%rd18]; shl.b64 %rd19, %rd2, 3; add.s64 %rd20, %rd16, %rd19; ld.global.v2.f32 {%f9, %f10}, [%rd20]; shl.b64 %rd21, %rd3, 3; add.s64 %rd22, %rd16, %rd21; ld.global.v2.f32 {%f13, %f14}, [%rd22]; mov.f32 %f17, 0f3F800000; sub.f32 %f18, %f17, %f1; sub.f32 %f19, %f18, %f2; mul.f32 %f20, %f1, %f9; mul.f32 %f21, %f1, %f10; fma.rn.f32 %f22, %f19, %f5, %f20; fma.rn.f32 %f23, %f19, %f6, %f21; fma.rn.f32 %f24, %f2, %f13, %f22; fma.rn.f32 %f25, %f2, %f14, %f23; ld.global.u32 %r4, [%rd14]; shr.u32 %r5, %r4, 16; abs.f32 %f26, %f24; cvt.rmi.f32.f32 %f27, %f26; sub.f32 %f28, %f26, %f27; abs.f32 %f29, %f25; cvt.rmi.f32.f32 %f30, %f29; sub.f32 %f31, %f29, %f30; cvta.to.global.u64 %rd23, %rd9; shl.b32 %r6, %r4, 4; cvt.u64.u32 %rd24, %r6; and.b64 %rd25, %rd24, 1048560; add.s64 %rd26, %rd23, %rd25; ld.global.v2.u32 {%r7, %r8}, [%rd26]; cvt.rn.f32.u32 %f32, %r7; mul.f32 %f33, %f28, %f32; cvt.rzi.u32.f32 %r11, %f33; cvt.rn.f32.u32 %f34, %r8; mul.f32 %f35, %f31, %f34; cvt.rzi.u32.f32 %r12, %f35; mad.lo.s32 %r13, %r7, %r12, %r11; cvt.u64.u32 %rd27, %r13; ld.global.u64 %rd28, [%rd26+8]; add.s64 %rd29, %rd28, %rd27; ld.u8 %r14, [%rd29]; and.b32 %r15, %r5, %r14; setp.eq.s32 %p2, %r15, 0; selp.f32 %f71, 0f00000000, 0f3F800000, %p2; $L__BB0_3: cvt.u32.u64 %r16, %rd1; cvt.u32.u64 %r17, %rd2; cvt.u32.u64 %r18, %rd3; ld.const.u64 %rd30, [params+16]; cvta.to.global.u64 %rd31, %rd30; mul.wide.s32 %rd32, %r16, 24; add.s64 %rd33, %rd31, %rd32; mul.wide.s32 %rd34, %r17, 24; add.s64 %rd35, %rd31, %rd34; mul.wide.s32 %rd36, %r18, 24; add.s64 %rd37, %rd31, %rd36; ld.global.f32 %f37, [%rd35]; ld.global.f32 %f38, [%rd33]; sub.f32 %f39, %f38, %f37; ld.global.f32 %f40, [%rd35+4]; ld.global.f32 %f41, [%rd33+4]; sub.f32 %f42, %f41, %f40; ld.global.f32 %f43, [%rd35+8]; ld.global.f32 %f44, [%rd33+8]; sub.f32 %f45, %f44, %f43; ld.global.f32 %f46, [%rd37]; sub.f32 %f47, %f37, %f46; ld.global.f32 %f48, [%rd37+4]; sub.f32 %f49, %f40, %f48; ld.global.f32 %f50, [%rd37+8]; sub.f32 %f51, %f43, %f50; mul.f32 %f52, %f42, %f51; mul.f32 %f53, %f45, %f49; sub.f32 %f54, %f52, %f53; mul.f32 %f55, %f45, %f47; mul.f32 %f56, %f39, %f51; sub.f32 %f57, %f55, %f56; mul.f32 %f58, %f39, %f49; mul.f32 %f59, %f42, %f47; sub.f32 %f60, %f58, %f59; mul.f32 %f61, %f57, %f57; fma.rn.f32 %f62, %f54, %f54, %f61; fma.rn.f32 %f63, %f60, %f60, %f62; sqrt.rn.f32 %f64, %f63; rcp.rn.f32 %f65, %f64; mul.f32 %f66, %f60, %f65; mul.f32 %f67, %f65, %f57; mul.f32 %f68, %f65, %f54; st.param.f32 [func_retval0+0], %f68; st.param.f32 [func_retval0+4], %f67; st.param.f32 [func_retval0+8], %f66; st.param.f32 [func_retval0+24], %f71; ret; } // .globl oxMain .visible .entry oxMain() { .reg .b64 %rd<2>; mov.u64 %rd1, __direct_callable__oxMain; // begin inline asm // end inline asm ret; }