// Disassembling _Z12pass_unifiedP13CUDA_Vtx_DataPfS1_ (0) .entry _Z12pass_unifiedP13CUDA_Vtx_DataPfS1_ { .lmem 0 .smem 2332 .reg 46 .bar 1 cvt.u32.u16 $r25, s[0x000c] cvt.u32.u16 $r1, %ntid.y cvt.u32.u16 $r11, $r0.lo mad24.lo.s32 $r22, $r25, $r1, $r11 set.ge.s32 $p0|$o127, $r22, c0[0x0028] @$p0.equ shl.u32 $r0, $r22, 0x00000001 @$p0.equ add.u32 $r0, $r0, $r22 @$p0.equ add.u32 $r12, $r0, c1[0x0010] @$p0.equ tex.1d.s32 {$r12,$r13,_,_}, $tex0, {$r12} @$p0.ne mov.b32 $r14, $r124 @$p0.ne mov.b32 $r13, $r124 @$p0.ne mov.b32 $r12, $r124 set.ge.s32 $p0|$o127, $r124, c0[0x0018] mov.b32 $r17, $r124 mov.b32 $r18, $r124 mov.b32 $r19, $r124 mov.b32 $r15, $r124 mov.b32 $r7, $r124 mov.b32 $r16, $r124 mov.b32 $r20, $r124 @$p0.ne bra.label label6 mul24.lo.s32 $r1, $r22, c0[0x0018] mov.b32 $r0, c0[0x0018] shl.u32 $r27, $r1, 0x00000005 mov.b32 $r21, $r124 shl.u32 $r26, $r0, 0x00000005 add.u32 $r28, $r27, c0[0x0014] label0: add.u32 $r1, $r28, $r21 mov.b64 $r2, g[$r1] join.label label1 add.b32 $r0, $r1, 0x00000008 add.b32 $r1, $r1, 0x00000014 mov.u32 $r5, g[$r1] mov.b64 $r0, g[$r0] cvt.s32.s16 $r4, $r2.lo cvt.s32.s16 $r2, $r2.hi cvt.s32.s16 $r6, $r3.lo cvt.s32.s16 $r35, $r3.hi set.eq.s32 $p0|$o127, $r4, c1[0x0000] cvt.s32.s16 $r36, $r0.lo cvt.s32.s16 $r33, $r0.hi cvt.s32.s16 $r31, $r1.lo shr.s32 $r0, $r5, 0x00000008 cvt.s32.s16 $r29, $r1.hi cvt.s32.s8 $r37, $r5.lo cvt.s32.s8 $r32, $r5.hi cvt.s32.s8 $r34, $r0.lo cvt.s32.s8 $r30, $r0.hi @$p0.ne bra.label label1 shl.u32 $r0, $r4, 0x00000001 add.u32 $r5, $r0, $r4 shl.u32 $r1, $r2, 0x00000001 shl.u32 $r3, $r6, 0x00000001 and.b32 $r0, $r35, c1[0x0004] add.half.b32 $r4, $r21, $r27 add.half.b32 $r1, $r1, $r2 add.u32 $r2, $r3, $r6 cvt.s32.u16 $p0|$o127, $r0.lo// (Unknown subsubop 07) add.u32 $r3, $r4, c0[0x0014] add.b32 $r8, $r1, 0x00000002 add.b32 $r4, $r2, 0x00000002 add.b32 $r0, $r5, 0x00000002 add.b32 $r3, $r3, 0x00000010 tex.1d.s32 {$r8,$r9,_,_}, $tex0, {$r8} tex.1d.s32 {$r4,$r5,_,_}, $tex0, {$r4} tex.1d.s32 {$r0,$r1,_,_}, $tex0, {$r0} mov.u32 $r3, g[$r3] add.half.rn.f32 $r24, $r8, $r4 add.half.rn.f32 $r23, $r9, $r5 add.half.rn.f32 $r38, $r0, $r24 add.half.rn.f32 $r23, $r23, $r1 add.half.rn.f32 $r24, $r4, -$r0 add.half.rn.f32 $r40, $r10, $r6 mul.rn.f32 $r39, $r38, 0x3eaaaaab mul.rn.f32 $r23, $r23, 0x3eaaaaab add.half.rn.f32 $r38, $r8, -$r0 add.half.rn.f32 $r40, $r2, $r40 add.half.rn.f32 $r0, -$r0, $r39 add.half.rn.f32 $r41, $r23, -$r1 add.half.rn.f32 $r42, -$r8, $r39 add.half.rn.f32 $r43, $r23, -$r9 add.half.rn.f32 $r45, $r23, -$r5 mul.half.rn.f32 $r8, $r41, $r41 add.half.rn.f32 $r44, -$r4, $r39 mul.half.rn.f32 $r4, $r43, $r43 mul.rn.f32 $r43, $r45, $r45 mad.rn.f32 $r41, $r0, $r0, $r8 mul.rn.f32 $r0, $r40, 0x3eaaaaab mad.rn.f32 $r40, $r42, $r42, $r4 mad.rn.f32 $r43, $r44, $r44, $r43 add.half.rn.f32 $r4, $r10, -$r2 add.half.rn.f32 $r8, -$r2, $r0 add.half.rn.f32 $r10, -$r10, $r0 add.half.rn.f32 $r2, $r6, -$r2 add.rn.f32 $r42, -$r6, $r0 mad.rn.f32 $r8, $r8, $r8, $r41 mad.rn.f32 $r10, $r10, $r10, $r40 mul24.lo.u32 $r6, $r11, 0x0000001c mad.rn.f32 $r40, $r42, $r42, $r43 rsqrt.f32 $r8, $r8 rsqrt.f32 $r10, $r10 movsh.b32 $ofs1, $r6, 0x00000000 add.half.rn.f32 $r6, $r9, -$r1 rcp.half.f32 $r8, $r8 rcp.half.f32 $r9, $r10 add.half.rn.f32 $r5, $r5, -$r1 mov.b32 s[$ofs1+0x001c], $r39 rsqrt.f32 $r1, $r40 add.half.rn.f32 $r9, $r8, $r9 mul.half.rn.f32 $r8, $r4, $r5 mov.b32 s[$ofs1+0x0020], $r23 rcp.half.f32 $r10, $r1 mul.half.rn.f32 $r1, $r38, $r2 mad.rn.f32 $r2, $r2, $r6, -$r8 mov.b32 s[$ofs1+0x0024], $r0 add.rn.f32 $r8, $r9, $r10 mad.rn.f32 $r1, $r24, $r4, -$r1 mul.rn.f32 $r0, $r24, $r6 mov.b32 s[$ofs1+0x002c], $r2 add.rn.f32 $r3, $r8, -$r3 mad.rn.f32 $r2, $r38, $r5, -$r0 mov.b32 s[$ofs1+0x0030], $r1 max.f32 $r3, $r3, $r60// (unk0 00400000) mov.b32 $r0, $r20 mov.b32 s[$ofs1+0x0034], $r2 mul.rn.f32 $r2, $r3, c0[0x0030] @$p0.ne mad.rn.f32 $r0, -$r23, $r1, $r20 mov.b32 s[$ofs1+0x0028], $r2 mov.b32 $r20, $r0 label1: nop.join bar.sync.u32 0x00000000 shr.s32 $r3, $r35, 0x00000001 set.le.s32 $p0|$o127, $r3, $r60// (unk0 00400000) join.label label2 @$p0.ne bra.label label2 shl.u32 $r0, $r37, 0x00000005 shl.u32 $r1, $r37, 0x00000002 shl.u32 $r2, $r36, 0x00000001 sub.half.b32 $r0, $r0, $r1 add.half.b32 $r1, $r2, $r36 movsh.b32 $ofs1, $r0, 0x00000000 add.b32 $r0, $r1, 0x00000002 add.half.rn.f32 $r19, s[$ofs1+0x002c], $r19 add.half.rn.f32 $r18, s[$ofs1+0x0030], $r18 add.rn.f32 $r17, s[$ofs1+0x0034], $r17 tex.1d.s32 {$r0,$r1,_,_}, $tex0, {$r0} add.half.rn.f32 $r5, -$r1, $r13 add.half.rn.f32 $r4, $r12, -$r0 mul.half.rn.f32 $r0, $r5, $r5 add.half.rn.f32 $r2, -$r2, $r14 mad.rn.f32 $r0, $r4, $r4, $r0 mad.rn.f32 $r6, $r2, $r2, $r0 max.f32 $r0, $r6, c1[0x0008] set?68?.u16.f32.f32 $p0|$o127, $r0, c1[0x000c] mov.b32 $r1, c0[0x0054] @$p0.neu mov.b32 $r1, c1[0x003c] @$p0.neu mul.rn.f32 $r0, $r0, c1[0x003c] @$p0.neu mul.rn.f32 $r1, $r1, c0[0x0054] rcp.f32 $r0, $r0 rsqrt.f32 $r6, $r6 mul.half.rn.f32 $r0, $r1, $r0 mul.half.rn.f32 $r1, $r6, $r4 mul.half.rn.f32 $r4, $r6, $r5 mul.half.rn.f32 $r2, $r6, $r2 mad.rn.f32 $r1, $r0, $r1, $r16 mad.rn.f32 $r4, $r0, $r4, $r7 mad.rn.f32 $r6, $r0, $r2, $r15 add.half.rn.f32 $r0, s[$ofs1+0x001c], -$r12 add.half.rn.f32 $r2, s[$ofs1+0x0020], -$r13 add.rn.f32 $r5, s[$ofs1+0x0024], -$r14 mad.rn.f32 $r16, s[$ofs1+0x0028], $r0, $r1 mad.rn.f32 $r7, s[$ofs1+0x0028], $r2, $r4 mad.rn.f32 $r15, s[$ofs1+0x0028], $r5, $r6 label2: set.join.le.s32 $p0|$o127, $r3, c1[0x0004] join.label label3 @$p0.ne bra.label label3 shl.u32 $r0, $r34, 0x00000005 shl.u32 $r1, $r34, 0x00000002 shl.u32 $r2, $r33, 0x00000001 sub.half.b32 $r0, $r0, $r1 add.half.b32 $r1, $r2, $r33 movsh.b32 $ofs1, $r0, 0x00000000 add.b32 $r0, $r1, 0x00000002 add.half.rn.f32 $r19, s[$ofs1+0x002c], $r19 add.half.rn.f32 $r18, s[$ofs1+0x0030], $r18 add.rn.f32 $r17, s[$ofs1+0x0034], $r17 tex.1d.s32 {$r0,$r1,_,_}, $tex0, {$r0} add.half.rn.f32 $r5, -$r1, $r13 add.half.rn.f32 $r4, $r12, -$r0 mul.half.rn.f32 $r0, $r5, $r5 add.half.rn.f32 $r2, -$r2, $r14 mad.rn.f32 $r0, $r4, $r4, $r0 mad.rn.f32 $r6, $r2, $r2, $r0 max.f32 $r0, $r6, c1[0x0008] set?68?.u16.f32.f32 $p0|$o127, $r0, c1[0x000c] mov.b32 $r1, c0[0x0054] @$p0.neu mov.b32 $r1, c1[0x003c] @$p0.neu mul.rn.f32 $r0, $r0, c1[0x003c] @$p0.neu mul.rn.f32 $r1, $r1, c0[0x0054] rcp.f32 $r0, $r0 rsqrt.f32 $r6, $r6 mul.half.rn.f32 $r0, $r1, $r0 mul.half.rn.f32 $r1, $r6, $r4 mul.half.rn.f32 $r4, $r6, $r5 mul.half.rn.f32 $r2, $r6, $r2 mad.rn.f32 $r1, $r0, $r1, $r16 mad.rn.f32 $r4, $r0, $r4, $r7 mad.rn.f32 $r6, $r0, $r2, $r15 add.half.rn.f32 $r0, s[$ofs1+0x001c], -$r12 add.half.rn.f32 $r2, s[$ofs1+0x0020], -$r13 add.rn.f32 $r5, s[$ofs1+0x0024], -$r14 mad.rn.f32 $r16, s[$ofs1+0x0028], $r0, $r1 mad.rn.f32 $r7, s[$ofs1+0x0028], $r2, $r4 mad.rn.f32 $r15, s[$ofs1+0x0028], $r5, $r6 label3: set.join.le.s32 $p0|$o127, $r3, c1[0x0010] join.label label4 @$p0.ne bra.label label4 shl.u32 $r0, $r32, 0x00000005 shl.u32 $r1, $r32, 0x00000002 shl.u32 $r2, $r31, 0x00000001 sub.half.b32 $r0, $r0, $r1 add.half.b32 $r1, $r2, $r31 movsh.b32 $ofs1, $r0, 0x00000000 add.b32 $r0, $r1, 0x00000002 add.half.rn.f32 $r19, s[$ofs1+0x002c], $r19 add.half.rn.f32 $r18, s[$ofs1+0x0030], $r18 add.rn.f32 $r17, s[$ofs1+0x0034], $r17 tex.1d.s32 {$r0,$r1,_,_}, $tex0, {$r0} add.half.rn.f32 $r5, -$r1, $r13 add.half.rn.f32 $r4, $r12, -$r0 mul.half.rn.f32 $r0, $r5, $r5 add.half.rn.f32 $r2, -$r2, $r14 mad.rn.f32 $r0, $r4, $r4, $r0 mad.rn.f32 $r6, $r2, $r2, $r0 max.f32 $r0, $r6, c1[0x0008] set?68?.u16.f32.f32 $p0|$o127, $r0, c1[0x000c] mov.b32 $r1, c0[0x0054] @$p0.neu mov.b32 $r1, c1[0x003c] @$p0.neu mul.rn.f32 $r0, $r0, c1[0x003c] @$p0.neu mul.rn.f32 $r1, $r1, c0[0x0054] rcp.f32 $r0, $r0 rsqrt.f32 $r6, $r6 mul.half.rn.f32 $r0, $r1, $r0 mul.half.rn.f32 $r1, $r6, $r4 mul.half.rn.f32 $r4, $r6, $r5 mul.half.rn.f32 $r2, $r6, $r2 mad.rn.f32 $r1, $r0, $r1, $r16 mad.rn.f32 $r4, $r0, $r4, $r7 mad.rn.f32 $r6, $r0, $r2, $r15 add.half.rn.f32 $r0, s[$ofs1+0x001c], -$r12 add.half.rn.f32 $r2, s[$ofs1+0x0020], -$r13 add.rn.f32 $r5, s[$ofs1+0x0024], -$r14 mad.rn.f32 $r16, s[$ofs1+0x0028], $r0, $r1 mad.rn.f32 $r7, s[$ofs1+0x0028], $r2, $r4 mad.rn.f32 $r15, s[$ofs1+0x0028], $r5, $r6 label4: set.join.le.s32 $p0|$o127, $r3, c1[0x0014] join.label label5 @$p0.ne bra.label label5 shl.u32 $r0, $r30, 0x00000005 shl.u32 $r1, $r30, 0x00000002 shl.u32 $r2, $r29, 0x00000001 sub.half.b32 $r0, $r0, $r1 add.half.b32 $r1, $r2, $r29 movsh.b32 $ofs1, $r0, 0x00000000 add.b32 $r0, $r1, 0x00000002 add.half.rn.f32 $r19, s[$ofs1+0x002c], $r19 add.half.rn.f32 $r18, s[$ofs1+0x0030], $r18 add.rn.f32 $r17, s[$ofs1+0x0034], $r17 tex.1d.s32 {$r0,$r1,_,_}, $tex0, {$r0} add.half.rn.f32 $r4, -$r1, $r13 add.half.rn.f32 $r3, $r12, -$r0 mul.half.rn.f32 $r0, $r4, $r4 add.half.rn.f32 $r2, -$r2, $r14 mad.rn.f32 $r0, $r3, $r3, $r0 mad.rn.f32 $r5, $r2, $r2, $r0 max.f32 $r0, $r5, c1[0x0008] set?68?.u16.f32.f32 $p0|$o127, $r0, c1[0x000c] mov.b32 $r1, c0[0x0054] @$p0.neu mov.b32 $r1, c1[0x003c] @$p0.neu mul.rn.f32 $r0, $r0, c1[0x003c] @$p0.neu mul.rn.f32 $r1, $r1, c0[0x0054] rcp.f32 $r0, $r0 rsqrt.f32 $r5, $r5 mul.half.rn.f32 $r0, $r1, $r0 mul.half.rn.f32 $r1, $r5, $r3 mul.half.rn.f32 $r3, $r5, $r4 mul.half.rn.f32 $r2, $r5, $r2 mad.rn.f32 $r1, $r0, $r1, $r16 mad.rn.f32 $r3, $r0, $r3, $r7 mad.rn.f32 $r5, $r0, $r2, $r15 add.half.rn.f32 $r0, s[$ofs1+0x001c], -$r12 add.half.rn.f32 $r2, s[$ofs1+0x0020], -$r13 add.rn.f32 $r4, s[$ofs1+0x0024], -$r14 mad.rn.f32 $r16, s[$ofs1+0x0028], $r0, $r1 mad.rn.f32 $r7, s[$ofs1+0x0028], $r2, $r3 mad.rn.f32 $r15, s[$ofs1+0x0028], $r4, $r5 label5: nop.join add.b32 $r21, $r21, 0x00000020 set.ne.s32 $p0|$o127, $r21, $r26 @$p0.ne bra.label label0 label6: mul.rn.f32 $r0, $r20, 0x3f000000 movsh.b32 $ofs1, $r11, 0x00000002 mov.b32 s[$ofs1+0x071c], $r0 bar.sync.u32 0x00000000 set.lt.s32 $p1|$r1, $r11, c1[0x0018] cvt.s32 $p0|$o127, $r1 join.label label7 @$p1.eq bra.label label7 add.b32 $r1, $r11, 0x000001cf add.b32 $r2, $r11, 0x000001d7 movsh.b32 $ofs2, $r1, 0x00000002 movsh.b32 $ofs3, $r2, 0x00000002 add.b32 $r1, $r11, 0x000001df add.rn.f32 $r2, s[$ofs2+0x0000], $r0 add.b32 $r0, $r11, 0x000001e7 movsh.b32 $ofs2, $r1, 0x00000002 add.rn.f32 $r2, s[$ofs3+0x0000], $r2 movsh.b32 $ofs3, $r0, 0x00000002 add.b32 $r1, $r11, 0x000001ef add.rn.f32 $r2, s[$ofs2+0x0000], $r2 add.b32 $r0, $r11, 0x000001f7 movsh.b32 $ofs2, $r1, 0x00000002 add.rn.f32 $r1, s[$ofs3+0x0000], $r2 movsh.b32 $ofs3, $r0, 0x00000002 add.b32 $r0, $r11, 0x000001ff add.rn.f32 $r1, s[$ofs2+0x0000], $r1 movsh.b32 $ofs2, $r0, 0x00000002 add.half.rn.f32 $r0, s[$ofs3+0x0000], $r1 add.half.rn.f32 $r0, s[$ofs2+0x0000], $r0 mov.b32 s[$ofs1+0x071c], $r0 label7: nop.join bar.sync.u32 0x00000000 set.eq.s32 $p2|$r1, $r11, $r60// (unk0 00400000) cvt.s32 $p1|$o127, $r1 join.label label8 @$p2.eq bra.label label8 mov.b32 $ofs2, 0x00000720 add.half.rn.f32 $r0, s[$ofs2+0x0000], $r0 add.half.rn.f32 $r0, s[$ofs2+0x0004], $r0 add.half.rn.f32 $r0, s[$ofs2+0x0008], $r0 add.half.rn.f32 $r0, s[$ofs2+0x000c], $r0 add.half.rn.f32 $r0, s[$ofs2+0x0010], $r0 add.half.rn.f32 $r0, s[$ofs2+0x0014], $r0 add.rn.f32 $r0, s[$ofs2+0x0018], $r0 label8: nop.join @$p1.neu shl.u32 $r1, $r25, 0x00000002 @$p1.neu add.u32 $r1, s[0x0018], $r1 @$p1.neu mov.u32 g[$r1], $r0 cvt.u32.u16 $r0, s[0x0008] and.b32 $r1, $r0, c1[0x001c] set.ne.s32 $r2, $r1, $r0 shr.s32 $r1, $r0, 0x00000006 and.b32 $r2, $r2, c1[0x0004] add.u32 $r2, $r1, $r2 mul24.lo.u32.u16.u16 $r1, $r2.lo, $r11.hi mad24.lo.u32.u16.u16.u32 $r1, $r2.hi, $r11.lo, $r1 shl.u32 $r1, $r1, 0x00000010 mad24.lo.u32.u16.u16.u32 $r1, $r2.lo, $r11.lo, $r1 add.u32 $r2, $r1, $r2 min.s32 $r2, $r0, $r2 set.le.s32 $p2|$o127, $r2, $r1 join.label label10 mov.b32 $r0, $r124 @$p2.ne bra.label label10 shl.u32 $r1, $r1, 0x00000002 shl.u32 $r2, $r2, 0x00000002 add.half.b32 $r1, s[0x0014], $r1 add.half.b32 $r2, s[0x0014], $r2 label9: mov.u32 $r3, g[$r1] add.b32 $r1, $r1, 0x00000004 set.ne.u32 $p2|$o127, $r1, $r2 add.rn.f32 $r0, $r3, $r0 @$p2.ne bra.label label9 label10: mov.join.b32 s[$ofs1+0x081c], $r0 mov.b32 $r1, $r0 bar.sync.u32 0x00000000 join.label label11 @$p0.eq bra.label label11 add.b32 $r1, $r11, 0x0000020f add.b32 $r2, $r11, 0x00000217 movsh.b32 $ofs2, $r1, 0x00000002 movsh.b32 $ofs3, $r2, 0x00000002 add.b32 $r1, $r11, 0x0000021f add.rn.f32 $r2, s[$ofs2+0x0000], $r0 add.b32 $r0, $r11, 0x00000227 movsh.b32 $ofs2, $r1, 0x00000002 add.rn.f32 $r2, s[$ofs3+0x0000], $r2 movsh.b32 $ofs3, $r0, 0x00000002 add.b32 $r1, $r11, 0x0000022f add.rn.f32 $r2, s[$ofs2+0x0000], $r2 add.b32 $r0, $r11, 0x00000237 movsh.b32 $ofs2, $r1, 0x00000002 add.rn.f32 $r1, s[$ofs3+0x0000], $r2 movsh.b32 $ofs3, $r0, 0x00000002 add.b32 $r0, $r11, 0x0000023f add.rn.f32 $r1, s[$ofs2+0x0000], $r1 movsh.b32 $ofs2, $r0, 0x00000002 add.half.rn.f32 $r0, s[$ofs3+0x0000], $r1 add.half.rn.f32 $r1, s[$ofs2+0x0000], $r0 mov.b32 s[$ofs1+0x081c], $r1 label11: nop.join bar.sync.u32 0x00000000 join.label label12 @$p1.eq bra.label label12 mov.b32 $ofs1, 0x00000820 add.half.rn.f32 $r0, s[$ofs1+0x0000], $r1 add.half.rn.f32 $r0, s[$ofs1+0x0004], $r0 add.half.rn.f32 $r0, s[$ofs1+0x0008], $r0 add.half.rn.f32 $r0, s[$ofs1+0x000c], $r0 add.half.rn.f32 $r0, s[$ofs1+0x0010], $r0 add.half.rn.f32 $r0, s[$ofs1+0x0014], $r0 add.rn.f32 $r0, s[$ofs1+0x0018], $r0 @$p1.neu mov.b32 s[0x081c], $r0 label12: nop.join bar.sync.u32 0x00000000 set.lt.s32 $p0|$o127, $r22, c0[0x0028] @$p0.equ return shl.u32 $r0, $r22, 0x00000001 add.u32 $r0, $r0, $r22 mov.b32 $ofs1, 0x0000081c add.b32 $r0, $r0, 0x00000001 cvt.rn.abs.f32 $r5, s[$ofs1+0x0000] set?68?.u16.f32.f32 $p0|$o127, s[$ofs1+0x0000], c1[$ofs1+0x000c] tex.1d.s32 {$r0,$r1,_,_}, $tex0, {$r0} mov.b32 $r4, c0[0x0038] @$p0.neu mov.b32 $r3, c1[0x003c] @$p0.neu mul.rn.f32 $r5, $r5, c1[0x003c] @$p0.neu mul.rn.f32 $r4, $r3, c0[0x0038] mov.b16.b8 $r3.lo, c0[0x002c] rcp.f32 $r6, $r5 cvt.s16.s8 $p1|$r5.lo, $r3.lo mul.rn.f32 $r3, $r4, $r6 set.ne.s16 $p0|$o127, $r5.lo, c1[0x0020] @$p1.equ bra.label label13 mul.rn.f32 $r4, $r13, c0[0x003c] mul.rn.f32 $r5, $r4, 0x3fb8aa3b// (unk0 00008000) cvt.rzi.f32 $r6, $r5 mad.rm.f32 $r5, $r6, c1[0x0024], $r4 mad.rn.f32 $r5, $r6, 0x35bfbe8e// (No operand 4 in this instruction) mul.rn.f32 $r8, $r5, 0x3fb8aa3b// (unk0 00008000) pre.ex2.f32 $r5, $r6 pre.ex2.f32 $r6, $r8 ex2.f32 $r5, $r5 ex2.f32 $r6, $r6 mul.rn.f32 $r5, $r5, $r6 set.lt.u32.f32.f32 $p1|$o127, $r4, c1[0x002c] set.gt.u32.f32.f32 $p2|$o127, $r4, c1[0x0030] @$p1.ne mov.b32 $r5, $r124 @$p2.ne mov.b32 $r5, c1[0x0034] mul.rn.f32 $r3, $r5, $r3 bra.label label13 label13: @$p0.eq bra.label label14 mov.b32 $r4, 0xbe4ccccd mul.half.rn.f32 $r4, $r4, c0[0x0048] mul.half.rn.f32 $r4, $r4, $r13 mul.rn.f32 $r5, $r4, 0x3fb8aa3b cvt.rzi.f32 $r6, $r5 mad.rm.f32 $r5, -$r6, c1[0x0024], $r4 mad.rn.f32 $r5, $r6, 0x35bfbe8e// (No operand 4 in this instruction) mul.rn.f32 $r8, $r5, 0x3fb8aa3b pre.ex2.f32 $r5, $r6 pre.ex2.f32 $r6, $r8 ex2.f32 $r5, $r5 ex2.f32 $r6, $r6 mul.rn.f32 $r6, $r5, $r6 set.lt.u16.f32.f32 $p0|$o127, $r4, c1[0x002c] set.gt.u16.f32.f32 $p1|$o127, $r4, c1[0x0030] @$p0.ne mov.b32 $r6, $r124 @$p1.ne mov.b32 $r6, c1[0x0034] bra.label label15 label14: mov.b32 $r6, 0x3f800000 label15: mul.rn.f32 $r4, $r1, $r1 mad.rn.f32 $r4, $r0, $r0, $r4 mad.rn.f32 $r4, $r2, $r2, $r4 rsqrt.f32 $r5, $r4 mul.rn.f32 $r9, $r18, 0x3e2aaaab mul.rn.f32 $r8, $r5, $r1 mul.rn.f32 $r4, $r19, 0x3e2aaaab mul.half.rn.f32 $r10, $r5, $r0 mul.half.rn.f32 $r8, $r9, $r8 mad.rn.f32 $r8, $r4, $r10, $r8 mul.rn.f32 $r10, $r17, 0x3e2aaaab mul.rn.f32 $r5, $r5, $r2 mad.rn.f32 $r8, $r10, $r5, $r8 mov.b32 $r5, c0[0x0044] max.f32 $r8, $r8, $r60// (unk0 00400000)// (unk1 04000000) add.half.rn.f32 $r5, $r5, c0[0x0058] mul.half.rn.f32 $r8, $r8, c0[0x0040] set?68?.u16.f32.f32 $p0|$o127, $r5, c1[0x000c] mov.b32 $r11, c0[0x0050] delta.f32 $r8, $r8// (unk1 04000000) @$p0.neu mov.b32 $r11, c1[0x003c] @$p0.neu mul.rn.f32 $r5, $r5, c1[0x003c] @$p0.neu mul.rn.f32 $r11, $r11, c0[0x0050] add.half.rn.f32 $r6, $r6, -$r3 rcp.half.f32 $r5, $r5 mul.half.rn.f32 $r18, $r4, $r6 mul.half.rn.f32 $r19, $r6, $r10 mul.rn.f32 $r17, $r11, $r5 mov.b32 $r11, c0[0x004c] mad.rn.f32 $r5, $r0, $r8, $r18 mad.rn.f32 $r20, $r8, $r2, $r19 mul.half.rn.f32 $r16, $r17, $r16 mul.half.rn.f32 $r19, $r17, $r15 mul.half.rn.f32 $r18, $r17, $r5 mul.half.rn.f32 $r20, $r17, $r20 mul.half.rn.f32 $r15, $r11, c0[0x0058] add.half.rn.f32 $r5, $r16, $r18 add.rn.f32 $r21, $r19, $r20 mad.rn.f32 $r11, $r6, $r9, -$r15 mad.rm.f32 $r5, $r5, c1[0x0038], $r0 mad.rm.f32 $r6, $r21, c1[0x0038], $r2 mad.rn.f32 $r8, $r8, $r1, $r11 mad.rn.f32 $r5, $r5, c0[0x0050], $r12 mad.rn.f32 $r11, $r6, c0[0x0050], $r14 mul.half.rn.f32 $r14, $r17, $r8 mul.half.rn.f32 $r12, $r17, $r7 set.le.u16.f32.f32 $r3.lo, $r5, c0[0x0064] set.ge.u16.f32.f32 $r4.lo, $r5, c0[0x0060] set.le.u16.f32.f32 $r8.hi, $r11, c0[0x006c] set.ge.u16.f32.f32 $r10.hi, $r11, c0[0x0068] and.b32 $r8, $r6, $r8 add.rn.f32 $r6, $r12, $r14 and.b32 $r17, $r17, $r21 mad.rm.f32 $r6, $r6, c1[0x0038], $r1 and.b32 $r8, $r8, $r17 set.ge.u16.f32.f32 $p0|$o127, $r13, $r60// (unk0 00400000) mad.rn.f32 $r6, $r6, c0[0x0050], $r13 cvt.neg.s32 $r13, $r8 mov.b32 $r8, $r124 @$p0.ne set.le.u16.f32.f32 $p0|$o127, $r6, $r60// (unk0 00400000) cvt.neg.s32 $r17, $r13 mad.rn.f32 $r13, $r16, c0[0x0034], $r18 @$p0.ne mov.b32 $r8, c1[0x0004] cvt.neg.s32 $r16, $r17 mad.rn.f32 $r12, $r12, c0[0x0034], $r14 mad.rn.f32 $r14, $r19, c0[0x0034], $r20 add.rn.f32 $r0, $r0, $r13 and.b32 $p0|$o127, $r8, $r16 add.half.rn.f32 $r1, $r1, $r12 add.half.rn.f32 $r2, $r2, $r14 join.label label16 @$p0.eq bra.label label16 add.rn.f32 $r1, $r7, -$r15 mad.rn.f32 $r1, -$r3, $r9, $r1 min.f32 $r3, $r1, $r60// (unk0 00400000) mov.b32 $r1, 0x7e800000 mul.rn.f32 $r6, $r3, 0x3d23d70a// (unk0 00008000) set?33?.u16.f32.f32 $p0|$o127, $r1, c0[0x0058] mov.b32 $r3, c0[0x0058] mul.rn.f32 $r1, $r6, c0[0x0050] @$p0.neu mov.b32 $r3, c1[0x003c] @$p0.neu mul.rn.f32 $r1, $r1, c1[0x003c] @$p0.neu mul.rn.f32 $r3, $r3, c0[0x0058] mul.rn.f32 $r6, $r2, $r2 mad.rn.f32 $r6, $r0, $r0, $r6 rcp.f32 $r3, $r3 rsqrt.f32 $r7, $r6 mul.half.rn.f32 $r1, $r1, $r3 rcp.half.f32 $r3, $r7 set.ge.u16.f32.f32 $p0|$o127, $r1, $r3 @$p0.neu mov.b32 $r2, $r124 @$p0.neu mov.b32 $r1, $r124 @$p0.neu mov.b32 $r0, $r124 @$p0.eq rsqrt.f32 $r3, $r6 @$p0.eq mul.rn.f32 $r6, $r3, $r2 @$p0.eq mul.rn.f32 $r3, $r3, $r0 @$p0.eq mad.rn.f32 $r2, -$r1, $r6, $r2 @$p0.eq mad.rn.f32 $r0, -$r1, $r3, $r0 @$p0.eq mov.b32 $r1, $r124 mov.b32 $r6, $r124 label16: shl.join.u32 $r3, $r22, 0x00000005 shl.u32 $r7, $r22, 0x00000004 add.half.b32 $r3, $r3, $r7 add.half.b32 $r3, s[0x0010], $r3 mov.b32 $r8, $r4 mov.b64 g[$r3], $r8 add.b32 $r4, $r3, 0x00000008 mov.u32 g[$r4], $r10 add.b32 $r4, $r3, 0x00000010 mov.b64 g[$r4], $r0 add.b32 $r0, $r3, 0x00000018 mov.u32 g[$r0], $r2 mov.b32 $r1, $r6 add.b32 $r2, $r3, 0x00000020 mov.b32 $r0, $r5 mov.b64 g[$r2], $r0 add.b32 $r0, $r3, 0x00000028 mov.end.u32 g[$r0], $r11 #.constseg 1:0x0000 const #{ #d.32 0xffffffff, 0x00000001, 0x3a83126f, 0x7e800000 // 0000 #d.32 0x00000002, 0x00000003, 0x00000008, 0xffffffc0 // 0010 #d.32 0x00000000, 0x3f317200, 0x35bfbe8e, 0xc2d20000 // 0020 #d.32 0x42d20000, 0x7f800000, 0x3f000000, 0x3e800000 // 0030 #} } // Disassembling _Z14pass_trianglesv (1) .entry _Z14pass_trianglesv { .lmem 0 .smem 272 .reg 32 .bar 1 mov.b16 $r0.hi, s[0x000c] cvt.u32.u16 $r7, $r0.lo mad24.lo.u32.u16.u16.u32 $r20, %ntid.y, $r0.hi, $r7 set.ge.s32 $p1|$r0, $r20, c0[0x0024] cvt.s32 $p0|$o127, $r0 @$p1.neu movsh.b32 $ofs1, $r7, 0x00000002 @$p1.neu mov.b32 s[$ofs1+0x0010], $r124 bar.sync.u32 0x00000000 @$p0.neu return shl.u32 $r0, $r20, 0x00000004 add.u32 $r2, $r0, c0[0x0000] mov.b64 $r0, g[$r2] add.b32 $r3, $r2, 0x00000008 mov.u32 $r3, g[$r3] cvt.s32.s16 $r4, $r0.lo cvt.s32.s16 $r5, $r0.hi cvt.s32.s16 $r6, $r1.lo cvt.s32.s16 $r9, $r1.hi shl.u32 $r0, $r4, 0x00000001 shl.u32 $r1, $r5, 0x00000001 shl.u32 $r8, $r6, 0x00000001 shl.u32 $r10, $r9, 0x00000001 add.half.b32 $r0, $r0, $r4 add.half.b32 $r1, $r1, $r5 add.half.b32 $r6, $r8, $r6 add.half.b32 $r4, $r10, $r9 add.b32 $r5, $r2, 0x0000000c cvt.s32.s16 $r16, $r3.lo cvt.s32.s16 $r21, $r3.hi add.b32 $r0, $r0, 0x00000002 add.b32 $r12, $r4, 0x00000002 add.b32 $r4, $r1, 0x00000002 add.b32 $r8, $r6, 0x00000002 tex.1d.s32 {$r0,$r1,_,_}, $tex0, {$r0} tex.1d.s32 {$r12,$r13,_,_}, $tex0, {$r12} mov.u32 $r11, g[$r5] tex.1d.s32 {$r4,$r5,_,_}, $tex0, {$r4} tex.1d.s32 {$r8,$r9,_,_}, $tex0, {$r8} add.half.rn.f32 $r13, $r1, -$r13 add.half.rn.f32 $r12, $r0, -$r12 mul.half.rn.f32 $r3, $r13, $r13 add.half.rn.f32 $r19, $r2, -$r14 mad.rn.f32 $r3, $r12, $r12, $r3 mad.rn.f32 $r23, $r19, $r19, $r3 mov.b32 $r22, c0[0x0054] mov.b32 $r15, $r1 max.f32 $r14, $r23, c1[0x0000] mov.half.b32 $r1, $r2 mov.half.b32 $r17, $r5 mov.b32 $r2, $r6 set?68?.u16.f32.f32 $p0|$o127, $r14, c1[0x0004] mov.half.b32 $r18, $r9 mov.half.b32 $r3, $r10 @$p0.neu mov.b32 $r5, c1[0x0010] @$p0.neu mul.rn.f32 $r14, $r14, c1[0x0010] @$p0.neu mul.rn.f32 $r22, $r5, c0[0x0054] rcp.f32 $r5, $r14 rsqrt.f32 $r6, $r23 shl.u32 $r10, $r16, 0x00000001 mul.half.rn.f32 $r5, $r22, $r5 mul.half.rn.f32 $r9, $r6, $r12 add.half.b32 $r12, $r10, $r16 mul.half.rn.f32 $r10, $r6, $r13 mul.half.rn.f32 $r6, $r6, $r19 mul.half.rn.f32 $r19, $r5, $r9 add.b32 $r12, $r12, 0x00000002 mul.half.rn.f32 $r16, $r5, $r10 mul.half.rn.f32 $r5, $r5, $r6 tex.1d.s32 {$r12,$r13,_,_}, $tex0, {$r12} add.half.rn.f32 $r9, -$r13, $r17 add.half.rn.f32 $r12, $r4, -$r12 mul.half.rn.f32 $r6, $r9, $r9 add.half.rn.f32 $r13, -$r14, $r2 mad.rn.f32 $r6, $r12, $r12, $r6 mad.rn.f32 $r14, $r13, $r13, $r6 max.f32 $r6, $r14, c1[0x0000] set?68?.u16.f32.f32 $p0|$o127, $r6, c1[0x0004] mov.b32 $r10, c0[0x0054] @$p0.neu mov.b32 $r10, c1[0x0010] @$p0.neu mul.rn.f32 $r6, $r6, c1[0x0010] @$p0.neu mul.rn.f32 $r10, $r10, c0[0x0054] rcp.f32 $r6, $r6 rsqrt.f32 $r14, $r14 shl.u32 $r22, $r21, 0x00000001 mul.half.rn.f32 $r6, $r10, $r6 mul.half.rn.f32 $r10, $r14, $r12 add.half.b32 $r12, $r22, $r21 mul.half.rn.f32 $r9, $r14, $r9 mul.half.rn.f32 $r13, $r14, $r13 mul.half.rn.f32 $r22, $r6, $r10 add.b32 $r12, $r12, 0x00000002 mul.half.rn.f32 $r23, $r6, $r9 mul.half.rn.f32 $r6, $r6, $r13 tex.1d.s32 {$r12,$r13,_,_}, $tex0, {$r12} add.half.rn.f32 $r10, -$r13, $r18 add.half.rn.f32 $r21, $r8, -$r12 mul.half.rn.f32 $r9, $r10, $r10 add.half.rn.f32 $r12, -$r14, $r3 mad.rn.f32 $r9, $r21, $r21, $r9 mad.rn.f32 $r14, $r12, $r12, $r9 max.f32 $r9, $r14, c1[0x0000] set?68?.u16.f32.f32 $p0|$o127, $r9, c1[0x0004] mov.b32 $r13, c0[0x0054] @$p0.neu mov.b32 $r13, c1[0x0010] @$p0.neu mul.rn.f32 $r9, $r9, c1[0x0010] @$p0.neu mul.rn.f32 $r13, $r13, c0[0x0054] rsqrt.f32 $r26, $r14 rcp.f32 $r25, $r9 shl.u32 $r9, $r20, 0x00000005 shl.u32 $r24, $r20, 0x00000004 mul.half.rn.f32 $r14, $r26, $r21 mul.half.rn.f32 $r20, $r26, $r10 mul.half.rn.f32 $r21, $r26, $r12 add.half.rn.f32 $r10, $r8, $r4 mul.half.rn.f32 $r13, $r13, $r25 add.half.b32 $r24, $r9, $r24 add.half.rn.f32 $r27, $r8, -$r0 add.half.rn.f32 $r12, $r0, $r10 add.half.rn.f32 $r26, -$r0, $r4 add.half.rn.f32 $r9, $r3, $r2 add.rn.f32 $r25, $r17, $r18 mad.rn.f32 $r0, -$r12, 0x3eaaaaab// (unk0 00400000)// (No operand 4 in this instruction) mad.rm.f32 $r10, $r12, c1[0x0008], -$r4 mad.rn.f32 $r8, -$r12, 0x3eaaaaab// (unk0 00400000)// (No operand 4 in this instruction) add.half.rn.f32 $r4, $r15, $r25 add.half.rn.f32 $r12, $r1, $r9 add.half.rn.f32 $r29, $r3, -$r1 add.half.rn.f32 $r30, -$r1, $r2 mul.rn.f32 $r25, $r4, 0x3eaaaaab mad.rm.f32 $r4, $r12, c1[0x0008], -$r1 mad.rm.f32 $r9, $r12, c1[0x0008], -$r2 mad.rm.f32 $r12, $r12, c1[0x0008], -$r3 add.half.rn.f32 $r28, -$r15, $r25 add.half.rn.f32 $r3, -$r17, $r25 mul.half.rn.f32 $r1, $r28, $r28 mul.half.rn.f32 $r2, $r3, $r3 mad.rn.f32 $r1, $r0, $r0, $r1 mad.rn.f32 $r2, $r10, $r10, $r2 mad.rn.f32 $r1, $r4, $r4, $r1 mad.rn.f32 $r2, $r9, $r9, $r2 rsqrt.f32 $r1, $r1 rsqrt.f32 $r2, $r2 rcp.half.f32 $r1, $r1 rcp.half.f32 $r2, $r2 add.half.rn.f32 $r31, $r1, $r2 add.half.rn.f32 $r2, $r17, -$r15 add.half.rn.f32 $r1, -$r15, $r18 add.half.rn.f32 $r15, -$r18, $r25 mul.rn.f32 $r17, $r15, $r15 mad.rn.f32 $r17, $r8, $r8, $r17 mad.rn.f32 $r17, $r12, $r12, $r17 rsqrt.f32 $r17, $r17 rcp.half.f32 $r17, $r17 add.half.rn.f32 $r17, $r31, $r17 add.rn.f32 $r11, $r17, -$r11 max.f32 $r11, $r11, $r60// (unk0 00400000) mul.rn.f32 $r17, $r11, c0[0x0030] mad.rn.f32 $r18, $r17, $r10, $r22 mad.rn.f32 $r23, $r17, $r3, $r23 mul.rn.f32 $r3, $r29, $r26 mad.rn.f32 $r22, $r30, $r27, -$r3 mul.half.rn.f32 $r10, $r1, $r30 mul.half.rn.f32 $r11, $r2, $r27 mad.rn.f32 $r3, $r17, $r0, $r19 mad.rn.f32 $r0, $r2, $r29, -$r10 mad.rn.f32 $r2, $r1, $r26, -$r11 add.half.b32 $r19, $r24, c0[0x0008] mov.half.b32 $r1, $r22 mad.rn.f32 $r10, $r17, $r28, $r16 mad.rn.f32 $r11, $r17, $r4, $r5 mov.b128 g[$r19], $r0 add.b32 $r0, $r19, 0x00000010 mov.b64 g[$r0], $r10 add.b32 $r0, $r19, 0x00000018 mov.u32 g[$r0], $r18 add.b32 $r0, $r19, 0x0000001c mad.rn.f32 $r1, $r17, $r9, $r6 mov.u32 g[$r0], $r23 mul.rn.f32 $r2, $r17, $r8 add.b32 $r0, $r19, 0x00000020 mul.half.rn.f32 $r3, $r17, $r15 mul.half.rn.f32 $r5, $r17, $r12 mad.rn.f32 $r4, $r13, $r14, $r2 mov.u32 g[$r0], $r1 add.b32 $r0, $r19, 0x00000024 mad.rn.f32 $r2, $r13, $r20, $r3 mad.rn.f32 $r3, $r13, $r21, $r5 mul.rn.f32 $r5, $r22, $r25 mov.u32 g[$r0], $r4 movsh.b32 $ofs1, $r7, 0x00000002 add.b32 $r1, $r19, 0x00000028 mul.rn.f32 $r0, $r5, 0x3f000000// (unk0 00008000) mov.b64 g[$r1], $r2 mov.b32 s[$ofs1+0x0010], $r0 bar.sync.u32 0x00000000 set.ge.s32 $p0|$o127, $r7, c1[0x000c] join.label label0 @$p0.ne bra.label label0 add.b32 $r1, $r7, 0x00000008 add.b32 $r2, $r7, 0x00000010 movsh.b32 $ofs2, $r1, 0x00000002 movsh.b32 $ofs3, $r2, 0x00000002 add.b32 $r1, $r7, 0x00000018 add.rn.f32 $r2, s[$ofs2+0x0010], $r0 add.b32 $r0, $r7, 0x00000020 movsh.b32 $ofs2, $r1, 0x00000002 add.rn.f32 $r2, s[$ofs3+0x0010], $r2 movsh.b32 $ofs3, $r0, 0x00000002 add.b32 $r1, $r7, 0x00000028 add.rn.f32 $r2, s[$ofs2+0x0010], $r2 add.b32 $r0, $r7, 0x00000030 movsh.b32 $ofs2, $r1, 0x00000002 add.rn.f32 $r1, s[$ofs3+0x0010], $r2 movsh.b32 $ofs3, $r0, 0x00000002 add.b32 $r0, $r7, 0x00000038 add.rn.f32 $r1, s[$ofs2+0x0010], $r1 movsh.b32 $ofs2, $r0, 0x00000002 add.half.rn.f32 $r0, s[$ofs3+0x0010], $r1 add.half.rn.f32 $r0, s[$ofs2+0x0010], $r0 mov.b32 s[$ofs1+0x0010], $r0 label0: nop.join bar.sync.u32 0x00000000 set.ne.s32 $p0|$o127, $r7, $r60// (unk0 00400000) join.label label1 @$p0.ne bra.label label1 add.half.rn.f32 $r0, s[0x0014], $r0 add.half.rn.f32 $r0, s[0x0018], $r0 add.half.rn.f32 $r0, s[0x001c], $r0 add.half.rn.f32 $r0, s[0x0020], $r0 add.half.rn.f32 $r0, s[0x0024], $r0 add.half.rn.f32 $r0, s[0x0028], $r0 add.rn.f32 $r0, s[0x002c], $r0 label1: set.join.ne.u32 $p0|$o127, $r7, $r60// (unk0 00400000) @$p0.ne return mov.b16 $r1.lo, 0x0004 mad24.lo.u32.u16.u16.u32 $r1, s[0x000c], $r1.lo, $r3 mov.end.u32 g[$r1], $r0 #.constseg 1:0x0000 const #{ #d.32 0x3a83126f, 0x7e800000, 0x3eaaaaab, 0x00000008 // 0000 #d.32 0x3e800000 // 0010 #} } // Disassembling _Z13pass_verticesP13CUDA_Vtx_Data (2) .entry _Z13pass_verticesP13CUDA_Vtx_Data { .lmem 0 .smem 276 .reg 25 .bar 1 mov.b32 $r1, 0xffffffc0 and.b32 $r2, $r1, c0[0x0024] mov.b32 $r1, c0[0x0024] set.ne.s32 $r2, $r2, c0[0x0024] shr.s32 $r1, $r1, 0x00000006 and.b32 $r2, $r2, c1[0x0000] add.u32 $r3, $r1, $r2 and.b32 $r1, $r3, c1[0x0004] set.ne.s32 $r2, $r3, $r1 shr.s32 $r1, $r3, 0x00000006 and.b32 $r2, $r2, c1[0x0000] add.u32 $r4, $r1, $r2 cvt.u32.u16 $r2, $r0.lo mul24.lo.u32.u16.u16 $r0, $r4.lo, $r2.hi mad24.lo.u32.u16.u16.u32 $r0, $r4.hi, $r2.lo, $r0 shl.u32 $r0, $r0, 0x00000010 mad24.lo.u32.u16.u16.u32 $r1, $r4.lo, $r2.lo, $r0 add.u32 $r0, $r1, $r4 min.s32 $r3, $r3, $r0 set.le.s32 $p0|$o127, $r3, $r1 join.label label1 mov.b32 $r0, $r124 @$p0.ne bra.label label1 shl.u32 $r1, $r1, 0x00000002 shl.u32 $r3, $r3, 0x00000002 add.half.b32 $r1, $r1, c0[0x000c] add.half.b32 $r3, $r3, c0[0x000c] label0: mov.u32 $r4, g[$r1] add.b32 $r1, $r1, 0x00000004 set.ne.u32 $p0|$o127, $r1, $r3 add.rn.f32 $r0, $r4, $r0 @$p0.ne bra.label label0 label1: movsh.join.b32 $ofs1, $r2, 0x00000002 mov.b32 s[$ofs1+0x0014], $r0 mov.b32 $r1, $r0 bar.sync.u32 0x00000000 set.ge.s32 $p0|$o127, $r2, c1[0x0008] join.label label2 @$p0.ne bra.label label2 add.b32 $r1, $r2, 0x00000008 add.b32 $r3, $r2, 0x00000010 movsh.b32 $ofs2, $r1, 0x00000002 movsh.b32 $ofs3, $r3, 0x00000002 add.b32 $r1, $r2, 0x00000018 add.rn.f32 $r3, s[$ofs2+0x0014], $r0 add.b32 $r0, $r2, 0x00000020 movsh.b32 $ofs2, $r1, 0x00000002 add.rn.f32 $r3, s[$ofs3+0x0014], $r3 movsh.b32 $ofs3, $r0, 0x00000002 add.b32 $r1, $r2, 0x00000028 add.rn.f32 $r3, s[$ofs2+0x0014], $r3 add.b32 $r0, $r2, 0x00000030 movsh.b32 $ofs2, $r1, 0x00000002 add.rn.f32 $r1, s[$ofs3+0x0014], $r3 movsh.b32 $ofs3, $r0, 0x00000002 add.b32 $r0, $r2, 0x00000038 add.rn.f32 $r1, s[$ofs2+0x0014], $r1 movsh.b32 $ofs2, $r0, 0x00000002 add.half.rn.f32 $r0, s[$ofs3+0x0014], $r1 add.half.rn.f32 $r1, s[$ofs2+0x0014], $r0 mov.b32 s[$ofs1+0x0014], $r1 label2: nop.join bar.sync.u32 0x00000000 set.eq.s32 $p1|$r0, $r2, $r60// (unk0 00400000) cvt.s32 $p0|$o127, $r0 join.label label3 @$p1.eq bra.label label3 add.half.rn.f32 $r0, s[0x0018], $r1 add.half.rn.f32 $r0, s[0x001c], $r0 add.half.rn.f32 $r0, s[0x0020], $r0 add.half.rn.f32 $r0, s[0x0024], $r0 add.half.rn.f32 $r0, s[0x0028], $r0 add.half.rn.f32 $r0, s[0x002c], $r0 add.rn.f32 $r1, s[0x0030], $r0 label3: nop.join @$p0.neu mov.b32 s[0x0014], $r1 bar.sync.u32 0x00000000 mov.b16 $r0.lo, %ntid.y mad24.lo.u32.u16.u16.u32 $r2, s[0x000c], $r0.lo, $r2 shl.u32 $r0, $r2, 0x00000001 shl.u32 $r1, $r2, 0x00000004 add.half.b32 $r3, $r2, $r0 add.half.b32 $r5, $r1, c0[0x0004] mov.b64 $r0, g[$r5] join.label label9 add.b32 $r4, $r3, 0x00000002 add.b32 $r8, $r3, 0x00000001 add.b32 $r3, $r5, 0x00000008 tex.1d.s32 {$r4,$r5,_,_}, $tex0, {$r4} tex.1d.s32 {$r8,$r9,_,_}, $tex0, {$r8} mov.b64 $r12, g[$r3] cvt.u32.u16 $r14, $r0.lo cvt.u32.u16 $r22, $r0.hi cvt.u32.u16 $r24, $r1.lo cvt.u32.u16 $r21, $r1.hi mov.half.b32 $r3, $r5 mov.half.b32 $r5, $r6 mov.half.b32 $r1, $r9 mov.half.b32 $r6, $r10 set.eq.s32 $p0|$o127, $r14, c1[0x000c] cvt.u32.u16 $r20, $r12.lo cvt.u32.u16 $r19, $r12.hi cvt.u32.u16 $r10, $r13.lo cvt.u32.u16 $r7, $r13.hi @$p0.ne bra.label label8 shr.s32 $r0, $r14, 0x00000002 shl.u32 $r9, $r0, 0x00000001 add.half.b32 $r0, $r9, $r0 mov.half.b32 $r12, $r0 join.label label7 and.b32 $p0|$r16, $r14, c1[0x0010] tex.1d.s32 {$r12,$r13,_,_}, $tex1, {$r12} mov.half.b32 $r9, $r13 mov.half.b32 $r11, $r14 @$p0.eq bra.label label4 set.eq.s32 $p0|$o127, $r16, c1[0x0000] @$p0.ne bra.label label5 set.eq.s32 $p0|$o127, $r16, c1[0x0014] @$p0.ne bra.label label6 mov.b32 $r0, $r124 mov.b32 $r14, $r124 mov.b32 $r16, $r124 bra.label label7 label4: add.b32 $r14, $r0, 0x00000001 mov.b32 $r16, $r0 tex.1d.s32 {$r14,$r15,_,_}, $tex1, {$r14} tex.1d.s32 {_,$r16,_,_}, $tex1, {$r16} mov.b32 $r0, $r15 bra.label label7 label5: add.b32 $r16, $r0, 0x00000001 add.b32 $r0, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} tex.1d.s32 {$r0,_,_,_}, $tex1, {$r0} mov.b32 $r14, $r17 bra.label label7 label6: add.b32 $r16, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} mov.half.b32 $r14, $r17 mov.half.b32 $r0, $r18 label7: mov.join.b32 $r15, $r0 mov.b32 $r13, $r16 bra.label label9 label8: mov.b32 $r11, $r124 mov.b32 $r9, $r124 mov.b32 $r15, $r124 mov.b32 $r14, $r124 mov.b32 $r13, $r124 mov.b32 $r12, $r124 label9: set.join.eq.s32 $p0|$o127, $r22, c1[0x000c] join.label label14 @$p0.ne bra.label label14 shr.s32 $r0, $r22, 0x00000002 shl.u32 $r16, $r0, 0x00000001 add.half.b32 $r0, $r16, $r0 mov.half.b32 $r16, $r0 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} add.half.rn.f32 $r9, $r17, $r9 add.half.rn.f32 $r11, $r18, $r11 add.rn.f32 $r12, $r16, $r12 and.b32 $p0|$r16, $r22, c1[0x0010] join.label label13 @$p0.eq bra.label label10 set.eq.s32 $p0|$o127, $r16, c1[0x0000] @$p0.ne bra.label label11 set.eq.s32 $p0|$o127, $r16, c1[0x0014] @$p0.ne bra.label label12 mov.b32 $r0, $r124 mov.b32 $r22, $r124 mov.b32 $r16, $r124 bra.label label13 label10: add.b32 $r22, $r0, 0x00000001 mov.b32 $r16, $r0 tex.1d.s32 {$r22,$r23,_,_}, $tex1, {$r22} tex.1d.s32 {_,$r16,_,_}, $tex1, {$r16} mov.b32 $r0, $r23 bra.label label13 label11: add.b32 $r16, $r0, 0x00000001 add.b32 $r0, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} tex.1d.s32 {$r0,_,_,_}, $tex1, {$r0} mov.b32 $r22, $r17 bra.label label13 label12: add.b32 $r16, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} mov.half.b32 $r22, $r17 mov.half.b32 $r0, $r18 label13: add.join.rn.f32 $r15, $r0, $r15 add.half.rn.f32 $r13, $r16, $r13 add.half.rn.f32 $r14, $r22, $r14 label14: set.join.eq.s32 $p0|$o127, $r24, c1[0x000c] join.label label19 @$p0.ne bra.label label19 shr.s32 $r0, $r24, 0x00000002 shl.u32 $r16, $r0, 0x00000001 add.half.b32 $r0, $r16, $r0 mov.half.b32 $r16, $r0 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} add.half.rn.f32 $r9, $r17, $r9 add.half.rn.f32 $r11, $r18, $r11 and.b32 $p0|$r17, $r24, c1[0x0010] add.rn.f32 $r12, $r16, $r12 join.label label18 @$p0.eq bra.label label15 set.eq.s32 $p0|$o127, $r17, c1[0x0000] @$p0.ne bra.label label16 set.eq.s32 $p0|$o127, $r17, c1[0x0014] @$p0.ne bra.label label17 mov.b32 $r0, $r124 mov.b32 $r22, $r124 mov.b32 $r16, $r124 bra.label label18 label15: add.b32 $r22, $r0, 0x00000001 mov.b32 $r16, $r0 tex.1d.s32 {$r22,$r23,_,_}, $tex1, {$r22} tex.1d.s32 {_,$r16,_,_}, $tex1, {$r16} mov.b32 $r0, $r23 bra.label label18 label16: add.b32 $r16, $r0, 0x00000001 add.b32 $r0, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} tex.1d.s32 {$r0,_,_,_}, $tex1, {$r0} mov.b32 $r22, $r17 bra.label label18 label17: add.b32 $r16, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} mov.half.b32 $r22, $r17 mov.half.b32 $r0, $r18 label18: add.join.rn.f32 $r15, $r0, $r15 add.half.rn.f32 $r13, $r16, $r13 add.half.rn.f32 $r14, $r22, $r14 label19: set.join.eq.s32 $p0|$o127, $r21, c1[0x000c] join.label label24 @$p0.ne bra.label label24 shr.s32 $r0, $r21, 0x00000002 shl.u32 $r16, $r0, 0x00000001 add.half.b32 $r0, $r16, $r0 mov.half.b32 $r16, $r0 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} add.half.rn.f32 $r9, $r17, $r9 add.half.rn.f32 $r11, $r18, $r11 and.b32 $p0|$r17, $r21, c1[0x0010] add.rn.f32 $r12, $r16, $r12 join.label label23 @$p0.eq bra.label label20 set.eq.s32 $p0|$o127, $r17, c1[0x0000] @$p0.ne bra.label label21 set.eq.s32 $p0|$o127, $r17, c1[0x0014] @$p0.ne bra.label label22 mov.b32 $r0, $r124 mov.b32 $r22, $r124 mov.b32 $r16, $r124 bra.label label23 label20: add.b32 $r22, $r0, 0x00000001 mov.b32 $r16, $r0 tex.1d.s32 {$r22,$r23,_,_}, $tex1, {$r22} tex.1d.s32 {_,$r16,_,_}, $tex1, {$r16} mov.b32 $r0, $r23 bra.label label23 label21: add.b32 $r16, $r0, 0x00000001 add.b32 $r0, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} tex.1d.s32 {$r0,_,_,_}, $tex1, {$r0} mov.b32 $r22, $r17 bra.label label23 label22: add.b32 $r16, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} mov.half.b32 $r22, $r17 mov.half.b32 $r0, $r18 label23: add.join.rn.f32 $r15, $r0, $r15 add.half.rn.f32 $r13, $r16, $r13 add.half.rn.f32 $r14, $r22, $r14 label24: set.join.eq.s32 $p0|$o127, $r20, c1[0x000c] join.label label29 @$p0.ne bra.label label29 shr.s32 $r0, $r20, 0x00000002 shl.u32 $r16, $r0, 0x00000001 add.half.b32 $r0, $r16, $r0 mov.half.b32 $r16, $r0 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} add.half.rn.f32 $r9, $r17, $r9 add.half.rn.f32 $r11, $r18, $r11 and.b32 $p0|$r17, $r20, c1[0x0010] add.rn.f32 $r12, $r16, $r12 join.label label28 @$p0.eq bra.label label25 set.eq.s32 $p0|$o127, $r17, c1[0x0000] @$p0.ne bra.label label26 set.eq.s32 $p0|$o127, $r17, c1[0x0014] @$p0.ne bra.label label27 mov.b32 $r0, $r124 mov.b32 $r20, $r124 mov.b32 $r16, $r124 bra.label label28 label25: add.b32 $r20, $r0, 0x00000001 mov.b32 $r16, $r0 tex.1d.s32 {$r20,$r21,_,_}, $tex1, {$r20} tex.1d.s32 {_,$r16,_,_}, $tex1, {$r16} mov.b32 $r0, $r21 bra.label label28 label26: add.b32 $r16, $r0, 0x00000001 add.b32 $r0, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} tex.1d.s32 {$r0,_,_,_}, $tex1, {$r0} mov.b32 $r20, $r17 bra.label label28 label27: add.b32 $r16, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} mov.half.b32 $r20, $r17 mov.half.b32 $r0, $r18 label28: add.join.rn.f32 $r15, $r0, $r15 add.half.rn.f32 $r13, $r16, $r13 add.half.rn.f32 $r14, $r20, $r14 label29: set.join.eq.s32 $p0|$o127, $r19, c1[0x000c] join.label label34 @$p0.ne bra.label label34 shr.s32 $r0, $r19, 0x00000002 shl.u32 $r16, $r0, 0x00000001 add.half.b32 $r0, $r16, $r0 mov.half.b32 $r16, $r0 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} add.half.rn.f32 $r9, $r17, $r9 add.half.rn.f32 $r11, $r18, $r11 and.b32 $p0|$r17, $r19, c1[0x0010] add.rn.f32 $r12, $r16, $r12 join.label label33 @$p0.eq bra.label label30 set.eq.s32 $p0|$o127, $r17, c1[0x0000] @$p0.ne bra.label label31 set.eq.s32 $p0|$o127, $r17, c1[0x0014] @$p0.ne bra.label label32 mov.b32 $r0, $r124 mov.b32 $r20, $r124 mov.b32 $r16, $r124 bra.label label33 label30: add.b32 $r20, $r0, 0x00000001 mov.b32 $r16, $r0 tex.1d.s32 {$r20,$r21,_,_}, $tex1, {$r20} tex.1d.s32 {_,$r16,_,_}, $tex1, {$r16} mov.b32 $r0, $r21 bra.label label33 label31: add.b32 $r16, $r0, 0x00000001 add.b32 $r0, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} tex.1d.s32 {$r0,_,_,_}, $tex1, {$r0} mov.b32 $r20, $r17 bra.label label33 label32: add.b32 $r16, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} mov.half.b32 $r20, $r17 mov.half.b32 $r0, $r18 label33: add.join.rn.f32 $r15, $r0, $r15 add.half.rn.f32 $r13, $r16, $r13 add.half.rn.f32 $r14, $r20, $r14 label34: set.join.eq.s32 $p0|$o127, $r10, c1[0x000c] join.label label39 @$p0.ne bra.label label39 shr.s32 $r0, $r10, 0x00000002 shl.u32 $r16, $r0, 0x00000001 add.half.b32 $r0, $r16, $r0 mov.half.b32 $r16, $r0 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} add.half.rn.f32 $r9, $r17, $r9 add.half.rn.f32 $r11, $r18, $r11 and.b32 $p0|$r10, $r10, c1[0x0010] add.rn.f32 $r12, $r16, $r12 join.label label38 @$p0.eq bra.label label35 set.eq.s32 $p0|$o127, $r10, c1[0x0000] @$p0.ne bra.label label36 set.eq.s32 $p0|$o127, $r10, c1[0x0014] @$p0.ne bra.label label37 mov.b32 $r0, $r124 mov.b32 $r20, $r124 mov.b32 $r16, $r124 bra.label label38 label35: add.b32 $r20, $r0, 0x00000001 mov.b32 $r16, $r0 tex.1d.s32 {$r20,$r21,_,_}, $tex1, {$r20} tex.1d.s32 {_,$r16,_,_}, $tex1, {$r16} mov.b32 $r0, $r21 bra.label label38 label36: add.b32 $r16, $r0, 0x00000001 add.b32 $r0, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} tex.1d.s32 {$r0,_,_,_}, $tex1, {$r0} mov.b32 $r20, $r17 bra.label label38 label37: add.b32 $r16, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} mov.half.b32 $r20, $r17 mov.half.b32 $r0, $r18 label38: add.join.rn.f32 $r15, $r0, $r15 add.half.rn.f32 $r13, $r16, $r13 add.half.rn.f32 $r14, $r20, $r14 label39: set.join.eq.s32 $p0|$o127, $r7, c1[0x000c] join.label label44 @$p0.ne bra.label label44 shr.s32 $r0, $r7, 0x00000002 shl.u32 $r10, $r0, 0x00000001 add.half.b32 $r0, $r10, $r0 mov.half.b32 $r16, $r0 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} add.half.rn.f32 $r9, $r17, $r9 add.half.rn.f32 $r11, $r18, $r11 and.b32 $p0|$r7, $r7, c1[0x0010] add.rn.f32 $r12, $r16, $r12 join.label label43 @$p0.eq bra.label label40 set.eq.s32 $p0|$o127, $r7, c1[0x0000] @$p0.ne bra.label label41 set.eq.s32 $p0|$o127, $r7, c1[0x0014] @$p0.ne bra.label label42 mov.b32 $r0, $r124 mov.b32 $r20, $r124 mov.b32 $r16, $r124 bra.label label43 label40: add.b32 $r20, $r0, 0x00000001 mov.b32 $r16, $r0 tex.1d.s32 {$r20,$r21,_,_}, $tex1, {$r20} tex.1d.s32 {_,$r16,_,_}, $tex1, {$r16} mov.b32 $r0, $r21 bra.label label43 label41: add.b32 $r16, $r0, 0x00000001 add.b32 $r0, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} tex.1d.s32 {$r0,_,_,_}, $tex1, {$r0} mov.b32 $r20, $r17 bra.label label43 label42: add.b32 $r16, $r0, 0x00000002 tex.1d.s32 {$r16,$r17,_,_}, $tex1, {$r16} mov.half.b32 $r20, $r17 mov.half.b32 $r0, $r18 label43: add.join.rn.f32 $r15, $r0, $r15 add.half.rn.f32 $r13, $r16, $r13 add.half.rn.f32 $r14, $r20, $r14 label44: cvt.join.rn.abs.f32 $r0, s[0x0014] set?68?.u16.f32.f32 $p0|$o127, s[0x0014], c1[0x0018] mov.b32 $r7, c0[0x0038] @$p0.neu mov.b32 $r7, c1[0x0038] @$p0.neu mul.rn.f32 $r0, $r0, c1[0x0038] @$p0.neu mul.rn.f32 $r7, $r7, c0[0x0038] mov.b16.b8 $r10.lo, c0[0x002c] rcp.f32 $r0, $r0 cvt.s16.s8 $p1|$r10.lo, $r10.lo mul.rn.f32 $r0, $r7, $r0 set.ne.s16 $p0|$o127, $r10.lo, c1[0x001c] @$p1.equ bra.label label45 mul.rn.f32 $r7, $r3, c0[0x003c] mul.rn.f32 $r10, $r7, 0x3fb8aa3b// (unk0 00008000) cvt.rzi.f32 $r16, $r10 mad.rm.f32 $r10, $r16, c1[0x0020], $r7 mad.rn.f32 $r10, $r16, 0x35bfbe8e// (No operand 4 in this instruction) mul.rn.f32 $r17, $r10, 0x3fb8aa3b// (unk0 00008000) pre.ex2.f32 $r10, $r16 pre.ex2.f32 $r16, $r17 ex2.f32 $r10, $r10 ex2.f32 $r16, $r16 mul.rn.f32 $r10, $r10, $r16 set.lt.u32.f32.f32 $p1|$o127, $r7, c1[0x0028] set.gt.u32.f32.f32 $p2|$o127, $r7, c1[0x002c] @$p1.ne mov.b32 $r10, $r124 @$p2.ne mov.b32 $r10, c1[0x0030] mul.rn.f32 $r0, $r10, $r0 bra.label label45 label45: @$p0.eq bra.label label46 mov.b32 $r7, 0xbe4ccccd mul.half.rn.f32 $r7, $r7, c0[0x0048] mul.half.rn.f32 $r7, $r7, $r3 mul.rn.f32 $r10, $r7, 0x3fb8aa3b cvt.rzi.f32 $r16, $r10 mad.rm.f32 $r10, -$r16, c1[0x0020], $r7 mad.rn.f32 $r10, $r16, 0x35bfbe8e// (No operand 4 in this instruction) mul.rn.f32 $r17, $r10, 0x3fb8aa3b pre.ex2.f32 $r10, $r16 pre.ex2.f32 $r16, $r17 ex2.f32 $r10, $r10 ex2.f32 $r16, $r16 mul.rn.f32 $r16, $r10, $r16 set.lt.u16.f32.f32 $p0|$o127, $r7, c1[0x0028] set.gt.u16.f32.f32 $p1|$o127, $r7, c1[0x002c] @$p0.ne mov.b32 $r16, $r124 @$p1.ne mov.b32 $r16, c1[0x0030] bra.label label47 label46: mov.b32 $r16, 0x3f800000 label47: mul.rn.f32 $r7, $r1, $r1 mad.rn.f32 $r7, $r8, $r8, $r7 mad.rn.f32 $r7, $r6, $r6, $r7 rsqrt.f32 $r17, $r7 mul.rn.f32 $r10, $r9, 0x3e2aaaab mul.rn.f32 $r9, $r17, $r1 mul.rn.f32 $r7, $r12, 0x3e2aaaab mul.half.rn.f32 $r12, $r17, $r8 mul.half.rn.f32 $r9, $r10, $r9 mad.rn.f32 $r9, $r7, $r12, $r9 mul.rn.f32 $r11, $r11, 0x3e2aaaab mul.rn.f32 $r12, $r17, $r6 mad.rn.f32 $r12, $r11, $r12, $r9 mov.b32 $r9, c0[0x0044] max.f32 $r12, $r12, $r60// (unk0 00400000)// (unk1 04000000) add.half.rn.f32 $r9, $r9, c0[0x0058] mul.half.rn.f32 $r12, $r12, c0[0x0040] set?68?.u16.f32.f32 $p0|$o127, $r9, c1[0x0018] mov.b32 $r17, c0[0x0050] delta.f32 $r12, $r12// (unk1 04000000) @$p0.neu mov.b32 $r17, c1[0x0038] @$p0.neu mul.rn.f32 $r9, $r9, c1[0x0038] @$p0.neu mul.rn.f32 $r17, $r17, c0[0x0050] add.rn.f32 $r19, $r16, -$r0 mov.b32 $r16, c0[0x004c] rcp.half.f32 $r9, $r9 mul.half.rn.f32 $r20, $r7, $r19 mul.half.rn.f32 $r16, $r16, c0[0x0058] mul.half.rn.f32 $r18, $r17, $r9 mad.rn.f32 $r9, $r8, $r12, $r20 mul.rn.f32 $r21, $r19, $r11 mad.rn.f32 $r17, $r19, $r10, -$r16 mul.half.rn.f32 $r19, $r18, $r13 mul.half.rn.f32 $r20, $r18, $r9 mad.rn.f32 $r13, $r12, $r6, $r21 mad.rn.f32 $r12, $r12, $r1, $r17 mul.half.rn.f32 $r15, $r18, $r15 add.half.rn.f32 $r9, $r19, $r20 mul.half.rn.f32 $r17, $r18, $r13 mul.half.rn.f32 $r13, $r18, $r12 mad.rm.f32 $r9, $r9, c1[0x0034], $r8 add.half.rn.f32 $r12, $r15, $r17 mul.half.rn.f32 $r18, $r18, $r14 mad.rn.f32 $r4, $r9, c0[0x0050], $r4 mad.rm.f32 $r9, $r12, c1[0x0034], $r6 add.rn.f32 $r21, $r18, $r13 mad.rn.f32 $r19, $r19, c0[0x0034], $r20 mad.rn.f32 $r12, $r9, c0[0x0050], $r5 mad.rm.f32 $r5, $r21, c1[0x0034], $r1 set.le.u16.f32.f32 $r4.hi, $r4, c0[0x0064] set.ge.u16.f32.f32 $r10.lo, $r4, c0[0x0060] set.ge.u16.f32.f32 $p0|$o127, $r3, $r60// (unk0 00400000) mad.rn.f32 $r3, $r5, c0[0x0050], $r3 and.b32 $r20, $r9, $r20 set.le.u16.f32.f32 $r2.hi, $r12, c0[0x006c] set.ge.u16.f32.f32 $r4.hi, $r12, c0[0x0068] and.b32 $r5, $r5, $r9 and.b32 $r5, $r20, $r5 cvt.neg.s32 $r5, $r5 mov.b32 $r9, $r124 @$p0.ne set.le.u16.f32.f32 $p0|$o127, $r3, $r60// (unk0 00400000) cvt.neg.s32 $r5, $r5 @$p0.ne mov.b32 $r9, c1[0x0000] cvt.neg.s32 $r20, $r5 mad.rn.f32 $r13, $r18, c0[0x0034], $r13 mad.rn.f32 $r15, $r15, c0[0x0034], $r17 add.rn.f32 $r5, $r8, $r19 and.b32 $p0|$o127, $r9, $r20 add.half.rn.f32 $r8, $r1, $r13 add.half.rn.f32 $r6, $r6, $r15 join.label label48 @$p0.eq bra.label label48 add.rn.f32 $r1, $r14, -$r16 mad.rn.f32 $r0, -$r0, $r10, $r1 min.f32 $r1, $r0, $r60// (unk0 00400000) mov.b32 $r0, 0x7e800000 mul.rn.f32 $r3, $r1, 0x3d23d70a// (unk0 00008000) set?33?.u16.f32.f32 $p0|$o127, $r0, c0[0x0058] mov.b32 $r1, c0[0x0058] mul.rn.f32 $r0, $r3, c0[0x0050] @$p0.neu mov.b32 $r1, c1[0x0038] @$p0.neu mul.rn.f32 $r0, $r0, c1[0x0038] @$p0.neu mul.rn.f32 $r1, $r1, c0[0x0058] mul.rn.f32 $r3, $r6, $r6 mad.rn.f32 $r3, $r5, $r5, $r3 rcp.f32 $r1, $r1 rsqrt.f32 $r8, $r3 mul.half.rn.f32 $r0, $r0, $r1 rcp.half.f32 $r1, $r8 set.ge.u16.f32.f32 $p0|$o127, $r0, $r1 @$p0.neu mov.b32 $r6, $r124 @$p0.neu mov.b32 $r8, $r124 @$p0.neu mov.b32 $r5, $r124 @$p0.eq rsqrt.f32 $r1, $r3 @$p0.eq mul.rn.f32 $r3, $r1, $r6 @$p0.eq mul.rn.f32 $r1, $r1, $r5 @$p0.eq mad.rn.f32 $r6, -$r0, $r3, $r6 @$p0.eq mad.rn.f32 $r5, -$r0, $r1, $r5 @$p0.eq mov.b32 $r8, $r124 mov.b32 $r3, $r124 label48: shl.join.u32 $r0, $r2, 0x00000005 shl.u32 $r1, $r2, 0x00000004 add.half.b32 $r0, $r0, $r1 mov.half.b32 $r1, $r10 add.half.b32 $r2, s[0x0010], $r0 mov.half.b32 $r0, $r7 mov.b64 g[$r2], $r0 add.b32 $r0, $r2, 0x00000008 mov.u32 g[$r0], $r11 mov.b32 $r1, $r8 add.b32 $r7, $r2, 0x00000010 mov.b32 $r0, $r5 mov.b64 g[$r7], $r0 add.b32 $r0, $r2, 0x00000018 mov.u32 g[$r0], $r6 mov.b32 $r5, $r3 add.b32 $r0, $r2, 0x00000020 mov.b64 g[$r0], $r4 add.b32 $r0, $r2, 0x00000028 mov.end.u32 g[$r0], $r12 #.constseg 1:0x0000 const #{ #d.32 0x00000001, 0xffffffc0, 0x00000008, 0xffffffff // 0000 #d.32 0x00000003, 0x00000002, 0x7e800000, 0x00000000 // 0010 #d.32 0x3f317200, 0x35bfbe8e, 0xc2d20000, 0x42d20000 // 0020 #d.32 0x7f800000, 0x3f000000, 0x3e800000 // 0030 #} }