| 123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475 |
- //
- // Generated by NVIDIA NVVM Compiler
- //
- // Compiler Build ID: UNKNOWN
- // Unknown Toolkit Version
- // Based on NVVM 7.0.1
- //
- .version 8.1
- .target sm_86, texmode_independent
- .address_size 32
- // .globl DynamicKernel_nop_fdiv
- .entry DynamicKernel_nop_fdiv(
- .param .u32 .ptr .global .align 8 DynamicKernel_nop_fdiv_param_0,
- .param .f64 DynamicKernel_nop_fdiv_param_1,
- .param .f64 DynamicKernel_nop_fdiv_param_2
- )
- {
- .reg .pred %p<8>;
- .reg .b32 %r<10>;
- .reg .f64 %fd<13>;
- ld.param.u32 %r2, [DynamicKernel_nop_fdiv_param_0];
- ld.param.f64 %fd5, [DynamicKernel_nop_fdiv_param_1];
- ld.param.f64 %fd6, [DynamicKernel_nop_fdiv_param_2];
- mov.b32 %r3, %envreg3;
- mov.u32 %r4, %ctaid.x;
- mov.u32 %r5, %ntid.x;
- mov.u32 %r6, %tid.x;
- add.s32 %r7, %r6, %r3;
- mad.lo.s32 %r1, %r5, %r4, %r7;
- abs.f64 %fd8, %fd6;
- setp.gtu.f64 %p1, %fd8, 0d7FF0000000000000;
- mov.f64 %fd7, 0d7FF8000000000214;
- mov.f64 %fd12, %fd7;
- @%p1 bra $L__BB0_7;
- abs.f64 %fd1, %fd5;
- setp.gtu.f64 %p2, %fd1, 0d7FF0000000000000;
- setp.neu.f64 %p3, %fd6, 0d0000000000000000;
- mov.f64 %fd12, 0d0000000000000000;
- and.pred %p4, %p3, %p2;
- @%p4 bra $L__BB0_7;
- setp.le.f64 %p5, %fd1, 0d7FF0000000000000;
- @%p5 bra $L__BB0_5;
- bra.uni $L__BB0_3;
- $L__BB0_5:
- setp.eq.f64 %p7, %fd6, 0d0000000000000000;
- mov.f64 %fd12, %fd7;
- @%p7 bra $L__BB0_7;
- div.rn.f64 %fd12, %fd5, %fd6;
- bra.uni $L__BB0_7;
- $L__BB0_3:
- setp.eq.f64 %p6, %fd6, 0d0000000000000000;
- mov.f64 %fd12, %fd7;
- @%p6 bra $L__BB0_7;
- rcp.rn.f64 %fd12, %fd6;
- $L__BB0_7:
- shl.b32 %r8, %r1, 3;
- add.s32 %r9, %r2, %r8;
- st.global.f64 [%r9], %fd12;
- ret;
- }
- ��
|