PTX //
// Generated by LLVM NVPTX Back-End
//

.version 8.2
.target sm_86
.address_size 64

        // .globl       add_kernel_0d1d2d3de

.visible .entry add_kernel_0d1d2d3de(
        .param .u64 add_kernel_0d1d2d3de_param_0,
        .param .u64 add_kernel_0d1d2d3de_param_1,
        .param .u64 add_kernel_0d1d2d3de_param_2,
        .param .u32 add_kernel_0d1d2d3de_param_3
)
.maxntid 128, 1, 1
{
        .reg .pred      %p<7>;
        .reg .b32       %r<33>;
        .reg .f32       %f<25>;
        .reg .b64       %rd<11>;
        .loc    1 5 0
$L__func_begin0:
        .loc    1 5 0

        ld.param.u64    %rd7, [add_kernel_0d1d2d3de_param_0];
        ld.param.u64    %rd8, [add_kernel_0d1d2d3de_param_1];
$L__tmp0:
        .loc    1 8 41
        mov.u32         %r26, %tid.x;
        shl.b32         %r27, %r26, 2;
        ld.param.u64    %rd9, [add_kernel_0d1d2d3de_param_2];
        and.b32         %r28, %r27, 508;
        ld.param.u32    %r29, [add_kernel_0d1d2d3de_param_3];
        .loc    1 6 24
        mov.u32 %r1, %ctaid.x;
        .loc    1 7 24
        shl.b32         %r30, %r1, 10;
        .loc    1 8 28
        or.b32          %r31, %r30, %r28;
        or.b32          %r32, %r31, 512;
        .loc    1 9 21
        setp.lt.s32     %p1, %r31, %r29;
        setp.lt.s32     %p2, %r32, %r29;
        .loc    1 10 24
        mul.wide.s32    %rd10, %r31, 4;
        add.s64         %rd1, %rd7, %rd10;
        add.s64         %rd2, %rd1, 2048;
        .loc    1 10 16
        mov.u32 %r2, 0x0;
        mov.u32 %r3, 0x0;
        mov.u32 %r4, 0x0;
        mov.u32 %r5, 0x0;
        @%p1 ld.global.v4.b32 { %r2, %r3, %r4, %r5 }, [ %rd1 + 0 ];
        mov.b32         %f1, %r2;
        mov.b32         %f2, %r3;
        mov.b32         %f3, %r4;
        mov.b32         %f4, %r5;
        mov.u32 %r6, 0x0;
        mov.u32 %r7, 0x0;
        mov.u32 %r8, 0x0;
        mov.u32 %r9, 0x0;
        @%p2 ld.global.v4.b32 { %r6, %r7, %r8, %r9 }, [ %rd2 + 0 ];
        mov.b32         %f5, %r6;
        mov.b32         %f6, %r7;
        mov.b32         %f7, %r8;
        mov.b32         %f8, %r9;
        .loc    1 11 24
        add.s64         %rd3, %rd8, %rd10;
        add.s64         %rd4, %rd3, 2048;
        .loc    1 11 16
        mov.u32 %r10, 0x0;
        mov.u32 %r11, 0x0;
        mov.u32 %r12, 0x0;
        mov.u32 %r13, 0x0;
        @%p1 ld.global.v4.b32 { %r10, %r11, %r12, %r13 }, [ %rd3 + 0 ];
        mov.b32         %f9, %r10;
        mov.b32         %f10, %r11;
        mov.b32         %f11, %r12;
        mov.b32         %f12, %r13;
        mov.u32 %r14, 0x0;
        mov.u32 %r15, 0x0;
        mov.u32 %r16, 0x0;
        mov.u32 %r17, 0x0;
        @%p2 ld.global.v4.b32 { %r14, %r15, %r16, %r17 }, [ %rd4 + 0 ];
        mov.b32         %f13, %r14;
        mov.b32         %f14, %r15;
        mov.b32         %f15, %r16;
        mov.b32         %f16, %r17;
        .loc    1 12 17
        add.f32         %f17, %f1, %f9;
        add.f32         %f18, %f2, %f10;
        add.f32         %f19, %f3, %f11;
        add.f32         %f20, %f4, %f12;
        add.f32         %f21, %f5, %f13;
        add.f32         %f22, %f6, %f14;
        add.f32         %f23, %f7, %f15;
        add.f32         %f24, %f8, %f16;
        .loc    1 13 26
        add.s64         %rd5, %rd9, %rd10;
        add.s64         %rd6, %rd5, 2048;
        .loc    1 13 35
        mov.b32         %r18, %f17;
        mov.b32         %r19, %f18;
        mov.b32         %r20, %f19;
        mov.b32         %r21, %f20;
        @%p1 st.global.v4.b32 [ %rd5 + 0 ], { %r18, %r19, %r20, %r21 };
        mov.b32         %r22, %f21;
        mov.b32         %r23, %f22;
        mov.b32         %r24, %f23;
        mov.b32         %r25, %f24;
        @%p2 st.global.v4.b32 [ %rd6 + 0 ], { %r22, %r23, %r24, %r25 };
        .loc    1 13 4
        ret;
$L__tmp1:
$L__func_end0:

}
        .file   1 "/home/kdidi/projects/triton/kd_scripts/vector_addition.py"
        .section        .debug_abbrev
        {
.b8 1
.b8 17
.b8 1
.b8 37
.b8 8
.b8 19
.b8 5
.b8 3
.b8 8
.b8 16
.b8 6
.b8 27
.b8 8
.b8 180
.b8 66
.b8 12
.b8 17
.b8 1
.b8 18
.b8 1
.b8 0
.b8 0
.b8 2
.b8 46
.b8 0
.b8 17
.b8 1
.b8 18
.b8 1
.b8 64
.b8 10
.b8 135
.b8 64
.b8 8
.b8 3
.b8 8
.b8 58
.b8 11
.b8 59
.b8 11
.b8 63
.b8 12
.b8 0
.b8 0
.b8 0
        }
        .section        .debug_info
        {
.b32 161
.b8 2
.b8 0
.b32 .debug_abbrev
.b8 8
.b8 1
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 0
.b8 2
.b8 0
.b8 118
.b8 101
.b8 99
.b8 116
.b8 111
.b8 114
.b8 95
.b8 97
.b8 100
.b8 100
.b8 105
.b8 116
.b8 105
.b8 111
.b8 110
.b8 46
.b8 112
.b8 121
.b8 0
.b32 .debug_line
.b8 47
.b8 104
.b8 111
.b8 109
.b8 101
.b8 47
.b8 107
.b8 100
.b8 105
.b8 100
.b8 105
.b8 47
.b8 112
.b8 114
.b8 111
.b8 106
.b8 101
.b8 99
.b8 116
.b8 115
.b8 47
.b8 116
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 47
.b8 107
.b8 100
.b8 95
.b8 115
.b8 99
.b8 114
.b8 105
.b8 112
.b8 116
.b8 115
.b8 0
.b8 1
.b64 $L__func_begin0
.b64 $L__func_end0
.b8 2
.b64 $L__func_begin0
.b64 $L__func_end0
.b8 1
.b8 156
.b8 97
.b8 100
.b8 100
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 95
.b8 48
.b8 100
.b8 49
.b8 100
.b8 50
.b8 100
.b8 51
.b8 100
.b8 101
.b8 0
.b8 97
.b8 100
.b8 100
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 95
.b8 48
.b8 100
.b8 49
.b8 100
.b8 50
.b8 100
.b8 51
.b8 100
.b8 101
.b8 0
.b8 1
.b8 5
.b8 1
.b8 0
        }
        .section        .debug_pubnames
        {
.b32 $L_pubNames_end0-$L_pubNames_start0
$L__pubNames_start0:
.b8 2
.b8 0
.b32 .debug_info
.b32 165
.b32 100
.b8 97
.b8 100
.b8 100
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 95
.b8 48
.b8 100
.b8 49
.b8 100
.b8 50
.b8 100
.b8 51
.b8 100
.b8 101
.b8 0
.b32 0
$L__pubNames_end0:
        }
        .section        .debug_pubtypes
        {
.b32 $L_pubTypes_end0-$L_pubTypes_start0
$L__pubTypes_start0:
.b8 2
.b8 0
.b32 .debug_info
.b32 165
.b32 0
$L__pubTypes_end0:
        }
        .section        .debug_loc      {       }