.version 1.4
.target sm_10, map_f64_to_f32

.entry VecAdd (
    .param .u64 __cudaparm_VecAdd_A,
    .param .u64 __cudaparm_VecAdd_B,
    .param .u64 __cudaparm_VecAdd_C,
    .param .s32 __cudaparm_VecAdd_N)
{
.reg .u16 %rh<4>;
.reg .u32 %r<6>;
.reg .u64 %rd<10>;
.reg .f32 %f<5>;
.reg .pred %p<3>;
.loc 28 21 0
$LDWbegin_VecAdd:
    mov.u16 %rh1, %ctaid.x;
    mov.u16 %rh2, %ntid.x;
    mul.wide.u16 %r1, %rh1, %rh2;
    cvt.u32.u16 %r2, %tid.x;
    add.u32 %r3, %r2, %r1;
    ld.param.s32 %r4, [__cudaparm_VecAdd_N];
    setp.le.s32 %p1, %r4, %r3;
    @%p1 bra $Lt_0_1026;
.loc 28 25 0
    cvt.s64.s32 %rd1, %r3;
    mul.wide.s32 %rd2, %r3, 4;
    ld.param.u64 %rd3, [__cudaparm_VecAdd_A];
    add.u64 %rd4, %rd3, %rd2;
    ld.global.f32 %f1, [%rd4+0];
    ld.param.u64 %rd5, [__cudaparm_VecAdd_B];
    add.u64 %rd6, %rd5, %rd2;
    ld.global.f32 %f2, [%rd6+0];
    add.f32 %f3, %f1, %f2;
    ld.param.u64 %rd7, [__cudaparm_VecAdd_C];
    add.u64 %rd8, %rd7, %rd2;
    st.global.f32 [%rd8+0], %f3;
$Lt_0_1026:
.loc 28 26 0
    exit;
$LDWend_VecAdd:
} // VecAdd