// // 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<70>; .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 %f69, 0f3F800000; sub.f32 %f39, %f69, %f1; sub.f32 %f68, %f39, %f2; 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 {%f8, %f9}, [%rd18]; shl.b64 %rd19, %rd2, 3; add.s64 %rd20, %rd16, %rd19; ld.global.v2.f32 {%f12, %f13}, [%rd20]; shl.b64 %rd21, %rd3, 3; add.s64 %rd22, %rd16, %rd21; ld.global.v2.f32 {%f16, %f17}, [%rd22]; mov.f32 %f20, 0f3F800000; sub.f32 %f21, %f20, %f1; sub.f32 %f68, %f21, %f2; mul.f32 %f22, %f1, %f12; mul.f32 %f23, %f1, %f13; fma.rn.f32 %f24, %f68, %f8, %f22; fma.rn.f32 %f25, %f68, %f9, %f23; fma.rn.f32 %f26, %f2, %f16, %f24; fma.rn.f32 %f27, %f2, %f17, %f25; ld.global.u32 %r4, [%rd14]; shr.u32 %r5, %r4, 16; abs.f32 %f28, %f26; cvt.rmi.f32.f32 %f29, %f28; sub.f32 %f30, %f28, %f29; abs.f32 %f31, %f27; cvt.rmi.f32.f32 %f32, %f31; sub.f32 %f33, %f31, %f32; 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 %f34, %r7; mul.f32 %f35, %f30, %f34; cvt.rzi.u32.f32 %r11, %f35; cvt.rn.f32.u32 %f36, %r8; mul.f32 %f37, %f33, %f36; cvt.rzi.u32.f32 %r12, %f37; 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 %f69, 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 %f40, [%rd33+12]; ld.global.f32 %f41, [%rd33+16]; ld.global.f32 %f42, [%rd33+20]; ld.global.f32 %f43, [%rd35+12]; mul.f32 %f44, %f1, %f43; ld.global.f32 %f45, [%rd35+16]; mul.f32 %f46, %f1, %f45; ld.global.f32 %f47, [%rd35+20]; mul.f32 %f48, %f1, %f47; fma.rn.f32 %f49, %f68, %f40, %f44; fma.rn.f32 %f50, %f68, %f41, %f46; fma.rn.f32 %f51, %f68, %f42, %f48; ld.global.f32 %f52, [%rd37+12]; ld.global.f32 %f53, [%rd37+16]; ld.global.f32 %f54, [%rd37+20]; fma.rn.f32 %f55, %f2, %f52, %f49; fma.rn.f32 %f56, %f2, %f53, %f50; fma.rn.f32 %f57, %f2, %f54, %f51; mul.f32 %f58, %f56, %f56; fma.rn.f32 %f59, %f55, %f55, %f58; fma.rn.f32 %f60, %f57, %f57, %f59; sqrt.rn.f32 %f61, %f60; rcp.rn.f32 %f62, %f61; mul.f32 %f63, %f57, %f62; mul.f32 %f64, %f56, %f62; mul.f32 %f65, %f55, %f62; st.param.f32 [func_retval0+0], %f65; st.param.f32 [func_retval0+4], %f64; st.param.f32 [func_retval0+8], %f63; st.param.f32 [func_retval0+24], %f69; ret; } // .globl oxMain .visible .entry oxMain() { .reg .b64 %rd<2>; mov.u64 %rd1, __direct_callable__oxMain; // begin inline asm // end inline asm ret; }