k_explicit

//
// 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_6d4f024526da42f9951e1dd4b5a94d35;

.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 .b32   %r<15>;
    .reg .b64   %rd<9>;


    ld.param.u64    %rd1, [root_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;
}


\00\00