Spaces:
Running
Running
| .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; | |
| } | |