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