k_read

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-21458526
// Driver 375.26
// Based on LLVM 3.4svn
//

.version 5.0
.target sm_35, texmode_independent
.address_size 64

    // .globl   init_z
.const .align 4 .u32 pyopencl_defeat_cache_7afa3ed8fb6c41d786580221a26bc001;

.entry init_z(
    .param .u64 .ptr .global .align 8 init_z_param_0
)
{
    .reg .b32   %r<15>;
    .reg .b64   %rd<9>;


    ld.param.u64    %rd1, [init_z_param_0];
    mov.b32 %r1, %envreg3;
    mov.u32     %r2, %ntid.x;
    mov.u32     %r3, %ctaid.x;
    mad.lo.s32  %r4, %r3, %r2, %r1;
    mov.u32     %r5, %tid.x;
    add.s32     %r6, %r4, %r5;
    mov.b32 %r7, %envreg4;
    mov.u32     %r8, %ntid.y;
    mov.u32     %r9, %ctaid.y;
    mad.lo.s32  %r10, %r9, %r8, %r7;
    mov.u32     %r11, %tid.y;
    add.s32     %r12, %r10, %r11;
    cvt.u64.u32 %rd2, %r12;
    mov.b32 %r13, %envreg7;
    mul.lo.s32  %r14, %r8, %r13;
    mul.wide.s32    %rd3, %r14, %r6;
    add.s64     %rd4, %rd3, %rd2;
    cvt.s64.s32     %rd5, %rd4;
    shl.b64     %rd6, %rd5, 3;
    add.s64     %rd7, %rd1, %rd6;
    mov.u64     %rd8, 4613937818241073152;
    st.global.u64   [%rd7], %rd8;
    ret;
}

    // .globl   root
.entry root(
    .param .u64 .ptr .global .align 8 root_param_0
)
{
    .reg .pred  %p<21>;
    .reg .f32   %f<3>;
    .reg .b32   %r<61>;
    .reg .f64   %fd<217>;
    .reg .b64   %rd<8>;


    ld.param.u64    %rd2, [root_param_0];
    mov.b32 %r18, %envreg3;
    mov.u32     %r19, %ctaid.x;
    mov.u32     %r20, %ntid.x;
    mad.lo.s32  %r21, %r19, %r20, %r18;
    mov.u32     %r22, %tid.x;
    add.s32     %r23, %r21, %r22;
    mov.u32     %r24, %ctaid.y;
    mov.u32     %r25, %ntid.y;
    mov.b32 %r26, %envreg4;
    mad.lo.s32  %r27, %r24, %r25, %r26;
    mov.u32     %r28, %tid.y;
    add.s32     %r29, %r27, %r28;
    mov.b32 %r30, %envreg7;
    mul.lo.s32  %r31, %r25, %r30;
    mad.lo.s32  %r32, %r31, %r23, %r29;
    mul.wide.s32    %rd3, %r32, 8;
    add.s64     %rd1, %rd2, %rd3;
    ld.global.f64   %fd25, [%rd1];
    cvt.rzi.s32.f64 %r33, %fd25;
    add.s32     %r34, %r33, -1;
    cvt.rn.f64.s32  %fd26, %r34;
    mov.f64     %fd27, 0d4032000000000000;
    div.rn.f64  %fd1, %fd27, %fd26;
    setp.eq.f64 %p1, %fd1, 0d3FF0000000000000;
    mov.f64     %fd216, 0d3FF0000000000000;
    @%p1 bra    BB1_26;

    abs.f64     %fd2, %fd1;
    setp.gtu.f64    %p2, %fd2, 0d7FF0000000000000;
    @%p2 bra    BB1_25;
    bra.uni     BB1_2;

BB1_25:
    add.f64     %fd216, %fd1, 0d3FE0000000000000;
    bra.uni     BB1_26;

BB1_2:
    setp.eq.f64 %p3, %fd1, 0d7FF0000000000000;
    @%p3 bra    BB1_24;
    bra.uni     BB1_3;

BB1_24:
    mov.f64     %fd213, 0d3FE0000000000000;
    {
    .reg .b32 %temp;
    mov.b64     {%temp, %r55}, %fd213;
    }
    setp.gt.s32 %p20, %r55, -1;
    selp.f64    %fd216, 0d7FF0000000000000, 0d0000000000000000, %p20;
    bra.uni     BB1_26;

BB1_3:
    mov.f64     %fd28, 0d3FE0000000000000;
    mul.rn.f64  %fd29, %fd28, %fd28;
    cvt.rzi.f64.f64 %fd30, %fd29;
    mov.f64     %fd31, 0d4000000000000000;
    mul.rn.f64  %fd32, %fd31, %fd30;
    sub.f64     %fd33, %fd28, %fd32;
    abs.f64     %fd3, %fd33;
    setp.eq.f64 %p4, %fd1, 0d0000000000000000;
    @%p4 bra    BB1_23;
    bra.uni     BB1_4;

BB1_23:
    setp.eq.f64 %p19, %fd3, 0d3FF0000000000000;
    selp.f64    %fd216, %fd1, 0d0000000000000000, %p19;
    bra.uni     BB1_26;

BB1_4:
    setp.eq.f64 %p5, %fd1, 0dFFF0000000000000;
    @%p5 bra    BB1_21;
    bra.uni     BB1_5;

BB1_21:
    neg.f64     %fd216, %fd1;
    setp.neu.f64    %p18, %fd3, 0d3FF0000000000000;
    @%p18 bra   BB1_26;

    mov.b64      %rd6, %fd216;
    xor.b64     %rd7, %rd6, -9223372036854775808;
    mov.b64      %fd216, %rd7;
    bra.uni     BB1_26;

BB1_5:
    setp.geu.f64    %p6, %fd1, 0d0000000000000000;
    @%p6 bra    BB1_7;

    cvt.rzi.f64.f64 %fd36, %fd28;
    setp.neu.f64    %p7, %fd36, 0d3FE0000000000000;
    mov.f64     %fd216, 0dFFF8000000000000;
    @%p7 bra    BB1_26;

BB1_7:
    // inline asm
    {
    .reg    .b32 lo;
    mov.b64     {lo, %r57}, %fd2;
    }
    // inline asm
    // inline asm
    {
    .reg    .b32 hi;
    mov.b64     {%r56, hi}, %fd2;
    }
    // inline asm
    bfe.u32     %r58, %r57, 20, 11;
    setp.ne.s32 %p8, %r58, 0;
    @%p8 bra    BB1_9;

    mov.f64     %fd41, 0d4350000000000000;
    mul.rn.f64  %fd40, %fd2, %fd41;
    // inline asm
    {
    .reg    .b32 lo;
    mov.b64     {lo, %r57}, %fd40;
    }
    // inline asm
    // inline asm
    {
    .reg    .b32 hi;
    mov.b64     {%r56, hi}, %fd40;
    }
    // inline asm
    bfe.u32     %r39, %r57, 20, 11;
    add.s32     %r58, %r39, -54;

BB1_9:
    add.s32     %r59, %r58, -1023;
    and.b32     %r42, %r57, -2146435073;
    or.b32      %r41, %r42, 1072693248;
    // inline asm
    mov.b64     %fd214, {%r56, %r41};
    // inline asm
    setp.lt.u32 %p9, %r41, 1073127583;
    @%p9 bra    BB1_11;

    // inline asm
    {
    .reg    .b32 hi;
    mov.b64     {%r43, hi}, %fd214;
    }
    // inline asm
    // inline asm
    {
    .reg    .b32 lo;
    mov.b64     {lo, %r44}, %fd214;
    }
    // inline asm
    add.s32     %r46, %r44, -1048576;
    // inline asm
    mov.b64     %fd214, {%r43, %r46};
    // inline asm
    add.s32     %r59, %r58, -1022;

BB1_11:
    add.f64     %fd130, %fd214, 0d3FF0000000000000;
    rcp.rn.f64  %fd131, %fd130;
    add.f64     %fd72, %fd214, 0dBFF0000000000000;
    mul.rn.f64  %fd132, %fd72, %fd131;
    add.f64     %fd120, %fd132, %fd132;
    mul.rn.f64  %fd68, %fd120, %fd120;
    mov.f64     %fd47, 0d3EB0F5FF7D2CAFE2;
    mov.f64     %fd49, 0d3ED0F5D241AD3B5A;
    // inline asm
    fma.rn.f64  %fd46, %fd47, %fd68, %fd49;
    // inline asm
    mov.f64     %fd53, 0d3EF3B20A75488A3F;
    // inline asm
    fma.rn.f64  %fd50, %fd46, %fd68, %fd53;
    // inline asm
    mov.f64     %fd57, 0d3F1745CDE4FAECD5;
    // inline asm
    fma.rn.f64  %fd54, %fd50, %fd68, %fd57;
    // inline asm
    mov.f64     %fd61, 0d3F3C71C7258A578B;
    // inline asm
    fma.rn.f64  %fd58, %fd54, %fd68, %fd61;
    // inline asm
    mov.f64     %fd65, 0d3F6249249242B910;
    // inline asm
    fma.rn.f64  %fd62, %fd58, %fd68, %fd65;
    // inline asm
    mov.f64     %fd69, 0d3F89999999999DFB;
    // inline asm
    fma.rn.f64  %fd66, %fd62, %fd68, %fd69;
    // inline asm
    mul.rn.f64  %fd133, %fd66, %fd68;
    sub.f64     %fd134, %fd72, %fd120;
    mul.rn.f64  %fd73, %fd31, %fd134;
    neg.f64     %fd71, %fd120;
    // inline asm
    fma.rn.f64  %fd70, %fd71, %fd72, %fd73;
    // inline asm
    mul.rn.f64  %fd116, %fd131, %fd70;
    add.f64     %fd136, %fd133, 0d3FB5555555555555;
    mov.f64     %fd137, 0d3FB5555555555555;
    sub.f64     %fd138, %fd137, %fd136;
    add.f64     %fd139, %fd133, %fd138;
    add.f64     %fd140, %fd139, 0d0000000000000000;
    add.f64     %fd141, %fd140, 0dBC46A4CB00B9E7B0;
    add.f64     %fd83, %fd136, %fd141;
    sub.f64     %fd142, %fd136, %fd83;
    add.f64     %fd87, %fd141, %fd142;
    mul.rn.f64  %fd143, %fd83, %fd120;
    neg.f64     %fd77, %fd143;
    // inline asm
    fma.rn.f64  %fd74, %fd83, %fd120, %fd77;
    // inline asm
    // inline asm
    fma.rn.f64  %fd78, %fd87, %fd116, %fd74;
    // inline asm
    // inline asm
    fma.rn.f64  %fd82, %fd83, %fd116, %fd78;
    // inline asm
    // inline asm
    fma.rn.f64  %fd86, %fd87, %fd120, %fd82;
    // inline asm
    add.f64     %fd99, %fd143, %fd86;
    sub.f64     %fd144, %fd143, %fd99;
    add.f64     %fd103, %fd86, %fd144;
    mul.rn.f64  %fd145, %fd99, %fd120;
    neg.f64     %fd93, %fd145;
    // inline asm
    fma.rn.f64  %fd90, %fd99, %fd120, %fd93;
    // inline asm
    // inline asm
    fma.rn.f64  %fd94, %fd103, %fd116, %fd90;
    // inline asm
    // inline asm
    fma.rn.f64  %fd98, %fd99, %fd116, %fd94;
    // inline asm
    // inline asm
    fma.rn.f64  %fd102, %fd103, %fd120, %fd98;
    // inline asm
    add.f64     %fd115, %fd145, %fd102;
    sub.f64     %fd146, %fd145, %fd115;
    add.f64     %fd119, %fd102, %fd146;
    mul.rn.f64  %fd147, %fd115, %fd120;
    neg.f64     %fd109, %fd147;
    // inline asm
    fma.rn.f64  %fd106, %fd115, %fd120, %fd109;
    // inline asm
    // inline asm
    fma.rn.f64  %fd110, %fd119, %fd116, %fd106;
    // inline asm
    // inline asm
    fma.rn.f64  %fd114, %fd115, %fd116, %fd110;
    // inline asm
    // inline asm
    fma.rn.f64  %fd118, %fd119, %fd120, %fd114;
    // inline asm
    add.f64     %fd148, %fd147, %fd118;
    sub.f64     %fd149, %fd147, %fd148;
    add.f64     %fd150, %fd118, %fd149;
    add.f64     %fd151, %fd120, %fd148;
    sub.f64     %fd152, %fd120, %fd151;
    add.f64     %fd153, %fd148, %fd152;
    add.f64     %fd154, %fd150, %fd153;
    add.f64     %fd155, %fd116, %fd154;
    add.f64     %fd156, %fd151, %fd155;
    sub.f64     %fd157, %fd151, %fd156;
    add.f64     %fd158, %fd155, %fd157;
    cvt.rn.f64.s32  %fd159, %r59;
    mov.f64     %fd160, 0d3FE62E42FEFA3000;
    mul.rn.f64  %fd161, %fd159, %fd160;
    mov.f64     %fd162, 0d3D53DE6AF278ECE6;
    mul.rn.f64  %fd163, %fd159, %fd162;
    add.f64     %fd164, %fd161, %fd156;
    sub.f64     %fd165, %fd161, %fd164;
    add.f64     %fd166, %fd156, %fd165;
    add.f64     %fd167, %fd158, %fd166;
    add.f64     %fd168, %fd163, %fd167;
    add.f64     %fd123, %fd164, %fd168;
    sub.f64     %fd169, %fd164, %fd123;
    add.f64     %fd127, %fd168, %fd169;
    mul.rn.f64  %fd170, %fd123, %fd28;
    neg.f64     %fd125, %fd170;
    // inline asm
    fma.rn.f64  %fd122, %fd123, %fd28, %fd125;
    // inline asm
    // inline asm
    fma.rn.f64  %fd126, %fd127, %fd28, %fd122;
    // inline asm
    add.f64     %fd7, %fd170, %fd126;
    sub.f64     %fd171, %fd170, %fd7;
    add.f64     %fd8, %fd126, %fd171;
    {
    .reg .b32 %temp;
    mov.b64     {%temp, %r13}, %fd7;
    }
    mov.b32      %f1, %r13;
    abs.f32     %f2, %f1;
    setp.lt.f32 %p10, %f2, 0f40874911;
    @%p10 bra   BB1_13;
    bra.uni     BB1_12;

BB1_13:
    mov.f64     %fd175, 0d3FF71547652B82FE;
    mul.rn.f64  %fd176, %fd7, %fd175;
    mov.f64     %fd177, 0d4338000000000000;
    add.rn.f64  %fd178, %fd176, %fd177;
    {
    .reg .b32 %temp;
    mov.b64     {%r14, %temp}, %fd178;
    }
    mov.f64     %fd179, 0dC338000000000000;
    add.rn.f64  %fd180, %fd178, %fd179;
    mov.f64     %fd181, 0dBFE62E42FEFA39EF;
    fma.rn.f64  %fd182, %fd180, %fd181, %fd7;
    mov.f64     %fd183, 0dBC7ABC9E3B39803F;
    fma.rn.f64  %fd184, %fd180, %fd183, %fd182;
    mov.f64     %fd185, 0d3E928AF3FCA213EA;
    mov.f64     %fd186, 0d3E5ADE1569CE2BDF;
    fma.rn.f64  %fd187, %fd186, %fd184, %fd185;
    mov.f64     %fd188, 0d3EC71DEE62401315;
    fma.rn.f64  %fd189, %fd187, %fd184, %fd188;
    mov.f64     %fd190, 0d3EFA01997C89EB71;
    fma.rn.f64  %fd191, %fd189, %fd184, %fd190;
    mov.f64     %fd192, 0d3F2A01A014761F65;
    fma.rn.f64  %fd193, %fd191, %fd184, %fd192;
    mov.f64     %fd194, 0d3F56C16C1852B7AF;
    fma.rn.f64  %fd195, %fd193, %fd184, %fd194;
    mov.f64     %fd196, 0d3F81111111122322;
    fma.rn.f64  %fd197, %fd195, %fd184, %fd196;
    mov.f64     %fd198, 0d3FA55555555502A1;
    fma.rn.f64  %fd199, %fd197, %fd184, %fd198;
    mov.f64     %fd200, 0d3FC5555555555511;
    fma.rn.f64  %fd201, %fd199, %fd184, %fd200;
    mov.f64     %fd202, 0d3FE000000000000B;
    fma.rn.f64  %fd203, %fd201, %fd184, %fd202;
    mov.f64     %fd204, 0d3FF0000000000000;
    fma.rn.f64  %fd205, %fd203, %fd184, %fd204;
    fma.rn.f64  %fd215, %fd205, %fd184, %fd204;
    abs.s32     %r47, %r14;
    setp.lt.s32 %p13, %r47, 1023;
    @%p13 bra   BB1_15;
    bra.uni     BB1_14;

BB1_15:
    shl.b32     %r53, %r14, 20;
    add.s32     %r60, %r53, 1072693248;
    bra.uni     BB1_16;

BB1_12:
    setp.lt.s32 %p11, %r13, 0;
    selp.f64    %fd172, 0d0000000000000000, 0d7FF0000000000000, %p11;
    abs.f64     %fd173, %fd7;
    setp.gtu.f64    %p12, %fd173, 0d7FF0000000000000;
    add.f64     %fd174, %fd7, %fd7;
    selp.f64    %fd216, %fd174, %fd172, %p12;
    bra.uni     BB1_17;

BB1_14:
    add.s32     %r48, %r14, 2046;
    shl.b32     %r49, %r48, 19;
    and.b32     %r50, %r49, -1048576;
    shl.b32     %r51, %r48, 20;
    sub.s32     %r60, %r51, %r50;
    mov.u32     %r52, 0;
    mov.b64     %fd206, {%r52, %r50};
    mul.f64     %fd215, %fd215, %fd206;

BB1_16:
    mov.u32     %r54, 0;
    mov.b64     %fd207, {%r54, %r60};
    mul.f64     %fd216, %fd215, %fd207;

BB1_17:
    abs.f64     %fd208, %fd216;
    setp.eq.f64 %p14, %fd208, 0d7FF0000000000000;
    @%p14 bra   BB1_19;

    // inline asm
    fma.rn.f64  %fd216, %fd216, %fd8, %fd216;
    // inline asm

BB1_19:
    setp.neu.f64    %p15, %fd3, 0d3FF0000000000000;
    or.pred     %p17, %p6, %p15;
    @%p17 bra   BB1_26;

    mov.b64      %rd4, %fd216;
    xor.b64     %rd5, %rd4, -9223372036854775808;
    mov.b64      %fd216, %rd5;

BB1_26:
    st.global.f64   [%rd1], %fd216;
    ret;
}


\00\00