ptx注入接管显卡

ptx_source = “””
.version 6.5
.target sm_70
.address_size 64

.visible .entry calc_force_step(
.param .u64 pos_addr,
.param .u64 vel_addr,
.param .u64 force_addr,
.param .u32 N,
.param .f32 dt
) {
.reg .b64 %rd<10>;
.reg .b32 %r<10>;
.reg .b32 %idx, %n_val;
.reg .f32 %pos, %vel, %force;
.reg .f32 %vel_new, %pos_new, %delta_t;
.reg .pred %p1;

mov.u32 %r1, %ctaid.x;
mov.u32 %r2, %ntid.x;
mul.lo.u32 %r3, %r1, %r2;
mov.u32 %r4, %tid.x;
add.u32 %idx, %r3, %r4;

ld.param.u32 %n_val, [N];
setp.ge.u32 %p1, %idx, %n_val;
@%p1 exit;

cvt.u64.u32 %rd1, %idx;
mul.wide.u32 %rd2, %idx, 4;

ld.param.u64 %rd3, [pos_addr];
add.u64 %rd4, %rd3, %rd2;
ld.global.f32 %pos, [%rd4];

mul.f32 %force, %pos, -10.0;

ld.param.u64 %rd5, [force_addr];
add.u64 %rd6, %rd5, %rd2;
st.global.f32 [%rd6], %force;

ld.param.u64 %rd7, [vel_addr];
add.u64 %rd8, %rd7, %rd2;
ld.global.f32 %vel, [%rd8];

ld.param.f32 %delta_t, [dt];

fma.rn.f32 %vel_new, %force, %delta_t, %vel; 
fma.rn.f32 %pos_new, %vel_new, %delta_t, %pos;

st.global.f32 [%rd8], %vel_new;
st.global.f32 [%rd4], %pos_new;

}
“””