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