trivial-kernel.ptx

//
// 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;
}