//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-33281558
// Cuda compilation tools, release 12.3, V12.3.52
// Based on NVVM 7.0.1
//

.version 8.3
.target sm_89
.address_size 64


.func  (.param .b32 func_retval0) _Z6add_opff(
	.param .b32 _Z6add_opff_param_0,
	.param .b32 _Z6add_opff_param_1
)
{
	.reg .f32 	%f<4>;


	ld.param.f32 	%f1, [_Z6add_opff_param_0];
	ld.param.f32 	%f2, [_Z6add_opff_param_1];
	add.f32 	%f3, %f1, %f2;
	st.param.f32 	[func_retval0+0], %f3;
	ret;

}
	// .globl	_Z3addPfS_S_m
.visible .entry _Z3addPfS_S_m(
	.param .u64 _Z3addPfS_S_m_param_0,
	.param .u64 _Z3addPfS_S_m_param_1,
	.param .u64 _Z3addPfS_S_m_param_2,
	.param .u64 _Z3addPfS_S_m_param_3
)
{
	.reg .pred 	%p<2>;
	.reg .f32 	%f<4>;
	.reg .b32 	%r<5>;
	.reg .b64 	%rd<13>;


	ld.param.u64 	%rd2, [_Z3addPfS_S_m_param_0];
	ld.param.u64 	%rd3, [_Z3addPfS_S_m_param_1];
	ld.param.u64 	%rd4, [_Z3addPfS_S_m_param_2];
	ld.param.u64 	%rd5, [_Z3addPfS_S_m_param_3];
	mov.u32 	%r1, %ctaid.x;
	mov.u32 	%r2, %ntid.x;
	mov.u32 	%r3, %tid.x;
	mad.lo.s32 	%r4, %r1, %r2, %r3;
	cvt.u64.u32 	%rd1, %r4;
	setp.ge.u64 	%p1, %rd1, %rd5;
	@%p1 bra 	$L__BB1_2;

	cvta.to.global.u64 	%rd6, %rd2;
	shl.b64 	%rd7, %rd1, 2;
	add.s64 	%rd8, %rd6, %rd7;
	cvta.to.global.u64 	%rd9, %rd3;
	add.s64 	%rd10, %rd9, %rd7;
	ld.global.f32 	%f1, [%rd10];
	ld.global.f32 	%f2, [%rd8];
	{ // callseq 0, 0
	.reg .b32 temp_param_reg;
	.param .b32 param0;
	st.param.f32 	[param0+0], %f2;
	.param .b32 param1;
	st.param.f32 	[param1+0], %f1;
	.param .b32 retval0;
	call.uni (retval0), 
	_Z6add_opff, 
	(
	param0, 
	param1
	);
	ld.param.f32 	%f3, [retval0+0];
	} // callseq 0
	cvta.to.global.u64 	%rd11, %rd4;
	add.s64 	%rd12, %rd11, %rd7;
	st.global.f32 	[%rd12], %f3;

$L__BB1_2:
	ret;

}