.version 8.0
.target sm_90
.address_size 64

.visible .entry saxpy(
	.param .f32 alpha,
	.param .u64 x,
	.param .u64 y,
	.param .u32 n
)
{
	.reg .pred %p;
	.reg .b32 %r<6>;
	.reg .f32 %f<5>;
	.reg .b64 %rd<8>;

	ld.param.f32 %f1, [alpha];
	ld.param.u64 %rd1, [x];
	ld.param.u64 %rd2, [y];
	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 %rd3, %rd1;
	cvta.to.global.u64 %rd4, %rd2;
	mul.wide.s32 %rd5, %r5, 4;
	add.s64 %rd6, %rd3, %rd5;
	add.s64 %rd7, %rd4, %rd5;
	ld.global.f32 %f2, [%rd6];
	ld.global.f32 %f3, [%rd7];
	fma.rn.f32 %f4, %f1, %f2, %f3;
	st.global.f32 [%rd7], %f4;
DONE:
	ret;
}