.version 8.0
.target sm_80
.address_size 64

.visible .entry vecadd(
	.param .u64 a,
	.param .u64 b,
	.param .u64 c,
	.param .u32 n
)
{
	.reg .pred %p;
	.reg .b32 %r<6>;
	.reg .f32 %f<4>;
	.reg .b64 %rd<10>;

	ld.param.u64 %rd1, [a];
	ld.param.u64 %rd2, [b];
	ld.param.u64 %rd3, [c];
	ld.param.u32 %r1, [n];
	mov.u32 %r2, %ctaid.x;
	mov.u32 %r3, %ntid.x;
	mov.u32 %r4, %tid.x;
	mad.lo.s32 %r5, %r2, %r3, %r4;
	setp.ge.s32 %p, %r5, %r1;
	@%p bra DONE;
	cvta.to.global.u64 %rd4, %rd1;
	cvta.to.global.u64 %rd5, %rd2;
	cvta.to.global.u64 %rd6, %rd3;
	mul.wide.s32 %rd7, %r5, 4;
	add.s64 %rd8, %rd4, %rd7;
	ld.global.f32 %f1, [%rd8];
	add.s64 %rd8, %rd5, %rd7;
	ld.global.f32 %f2, [%rd8];
	add.f32 %f3, %f1, %f2;
	add.s64 %rd9, %rd6, %rd7;
	st.global.f32 [%rd9], %f3;
DONE:
	ret;
}