public
Created

  • Download Gist
trivial-kernel.ptx
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49
//
// Generated by LLVM NVPTX Back-End
//
 
.version 3.0
.target sm_13, texmode_independent
.address_size 32
 
 
// .globl _ZN9add_float16_cb9e1b436595b333_00E
.func _ZN9add_float16_cb9e1b436595b333_00E(
.reg .b32 _ZN9add_float16_cb9e1b436595b333_00E_param_0,
.reg .b32 _ZN9add_float16_cb9e1b436595b333_00E_param_1,
.reg .b32 _ZN9add_float16_cb9e1b436595b333_00E_param_2,
.reg .b32 _ZN9add_float16_cb9e1b436595b333_00E_param_3,
.reg .b32 _ZN9add_float16_cb9e1b436595b333_00E_param_4
) // @_ZN9add_float16_cb9e1b436595b333_00E
{
.local .align 4 .b8 __local_depot0[12];
.reg .b32 %SP;
.reg .b32 %SPL;
.reg .pred %p<396>;
.reg .s16 %rc<396>;
.reg .s16 %rs<396>;
.reg .s32 %r<396>;
.reg .s64 %rl<396>;
.reg .f32 %f<396>;
.reg .f64 %fl<396>;
 
// BB#0: // %static_allocas
mov.u32 %SP, __local_depot0;
mov.b32 %r0, _ZN9add_float16_cb9e1b436595b333_00E_param_2;
mov.b32 %r1, _ZN9add_float16_cb9e1b436595b333_00E_param_3;
mov.b32 %r2, _ZN9add_float16_cb9e1b436595b333_00E_param_4;
st.global.u32 [%SP+0], %r0;
st.global.u32 [%SP+4], %r1;
st.global.u32 [%SP+8], %r2;
// inline asm
# *z = *x + *y; (trivial-kernel.rs:3:4: 3:16)
// inline asm
ld.global.u32 %r0, [%SP+4];
ld.global.f64 %fl0, [%r0];
ld.global.u32 %r0, [%SP+0];
ld.global.f64 %fl1, [%r0];
add.f64 %fl0, %fl1, %fl0;
ld.global.u32 %r0, [%SP+8];
st.global.f64 [%r0], %fl0;
ret;
}

Please sign in to comment on this gist.

Something went wrong with that request. Please try again.