XFESystem

// XFESystem's own small reference kernel. Not copied from NVIDIA samples.
.version 8.7
.target sm_120
.address_size 64

.visible .entry XfeVectorAdd(
    .param .u64 InputA,
    .param .u64 InputB,
    .param .u64 Output,
    .param .u32 Count
)
{
    .reg .pred %outside;
    .reg .b32 %r<8>;
    .reg .b64 %rd<8>;
    ld.param.u64 %rd0, [InputA];
    ld.param.u64 %rd1, [InputB];
    ld.param.u64 %rd2, [Output];
    ld.param.u32 %r0, [Count];
    mov.u32 %r1, %ctaid.x;
    mov.u32 %r2, %ntid.x;
    mov.u32 %r3, %tid.x;
    mad.lo.u32 %r4, %r1, %r2, %r3;
    setp.ge.u32 %outside, %r4, %r0;
    @%outside bra Done;
    mul.wide.u32 %rd3, %r4, 4;
    add.u64 %rd4, %rd0, %rd3;
    add.u64 %rd5, %rd1, %rd3;
    add.u64 %rd6, %rd2, %rd3;
    ld.global.u32 %r5, [%rd4];
    ld.global.u32 %r6, [%rd5];
    add.u32 %r7, %r5, %r6;
    st.global.u32 [%rd6], %r7;
Done:
    ret;
}