matrix_operations_cuda 0.1.2

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-32267302
// Cuda compilation tools, release 12.0, V12.0.140
// Based on NVVM 7.0.1
//

.version 8.0
.target sm_52
.address_size 64

	// .globl	add

.visible .entry add(
	.param .u64 add_param_0,
	.param .u64 add_param_1,
	.param .u64 add_param_2
)
{
	.reg .f32 	%f<4>;
	.reg .b32 	%r<8>;
	.reg .b64 	%rd<11>;


	ld.param.u64 	%rd1, [add_param_0];
	ld.param.u64 	%rd2, [add_param_1];
	ld.param.u64 	%rd3, [add_param_2];
	cvta.to.global.u64 	%rd4, %rd3;
	cvta.to.global.u64 	%rd5, %rd2;
	cvta.to.global.u64 	%rd6, %rd1;
	mov.u32 	%r1, %tid.x;
	mov.u32 	%r2, %tid.y;
	mov.u32 	%r3, %ntid.x;
	mov.u32 	%r4, %ctaid.x;
	mov.u32 	%r5, %ntid.y;
	mad.lo.s32 	%r6, %r5, %r4, %r2;
	mad.lo.s32 	%r7, %r6, %r3, %r1;
	mul.wide.s32 	%rd7, %r7, 4;
	add.s64 	%rd8, %rd6, %rd7;
	ld.global.f32 	%f1, [%rd8];
	add.s64 	%rd9, %rd5, %rd7;
	ld.global.f32 	%f2, [%rd9];
	add.f32 	%f3, %f1, %f2;
	add.s64 	%rd10, %rd4, %rd7;
	st.global.f32 	[%rd10], %f3;
	ret;

}
	// .globl	sub
.visible .entry sub(
	.param .u64 sub_param_0,
	.param .u64 sub_param_1,
	.param .u64 sub_param_2
)
{
	.reg .f32 	%f<4>;
	.reg .b32 	%r<8>;
	.reg .b64 	%rd<11>;


	ld.param.u64 	%rd1, [sub_param_0];
	ld.param.u64 	%rd2, [sub_param_1];
	ld.param.u64 	%rd3, [sub_param_2];
	cvta.to.global.u64 	%rd4, %rd3;
	cvta.to.global.u64 	%rd5, %rd2;
	cvta.to.global.u64 	%rd6, %rd1;
	mov.u32 	%r1, %tid.x;
	mov.u32 	%r2, %tid.y;
	mov.u32 	%r3, %ntid.x;
	mov.u32 	%r4, %ctaid.x;
	mov.u32 	%r5, %ntid.y;
	mad.lo.s32 	%r6, %r5, %r4, %r2;
	mad.lo.s32 	%r7, %r6, %r3, %r1;
	mul.wide.s32 	%rd7, %r7, 4;
	add.s64 	%rd8, %rd6, %rd7;
	ld.global.f32 	%f1, [%rd8];
	add.s64 	%rd9, %rd5, %rd7;
	ld.global.f32 	%f2, [%rd9];
	sub.f32 	%f3, %f1, %f2;
	add.s64 	%rd10, %rd4, %rd7;
	st.global.f32 	[%rd10], %f3;
	ret;

}
	// .globl	dot
.visible .entry dot(
	.param .u64 dot_param_0,
	.param .u64 dot_param_1,
	.param .u64 dot_param_2,
	.param .u32 dot_param_3,
	.param .u32 dot_param_4,
	.param .u32 dot_param_5,
	.param .u32 dot_param_6
)
{
	.reg .pred 	%p<6>;
	.reg .f32 	%f<30>;
	.reg .b32 	%r<33>;
	.reg .b64 	%rd<35>;


	ld.param.u64 	%rd18, [dot_param_0];
	ld.param.u64 	%rd19, [dot_param_1];
	ld.param.u64 	%rd17, [dot_param_2];
	ld.param.u32 	%r13, [dot_param_4];
	ld.param.u32 	%r14, [dot_param_5];
	ld.param.u32 	%r15, [dot_param_6];
	cvta.to.global.u64 	%rd1, %rd19;
	cvta.to.global.u64 	%rd2, %rd18;
	mov.u32 	%r16, %ntid.y;
	mov.u32 	%r17, %ctaid.x;
	mov.u32 	%r18, %tid.y;
	mad.lo.s32 	%r19, %r16, %r17, %r18;
	mov.u32 	%r20, %ntid.x;
	mov.u32 	%r21, %tid.x;
	mad.lo.s32 	%r1, %r19, %r20, %r21;
	div.s32 	%r3, %r1, %r15;
	mul.lo.s32 	%r22, %r3, %r15;
	sub.s32 	%r2, %r1, %r22;
	setp.lt.s32 	%p1, %r14, 1;
	mov.f32 	%f29, 0f00000000;
	@%p1 bra 	$L__BB2_7;

	add.s32 	%r24, %r14, -1;
	and.b32  	%r32, %r14, 3;
	setp.lt.u32 	%p2, %r24, 3;
	mov.f32 	%f29, 0f00000000;
	mov.u32 	%r31, 0;
	@%p2 bra 	$L__BB2_4;

	sub.s32 	%r30, %r14, %r32;
	mul.wide.s32 	%rd20, %r2, 4;
	add.s64 	%rd32, %rd1, %rd20;
	mul.lo.s32 	%r26, %r3, %r13;
	mul.wide.s32 	%rd21, %r26, 4;
	add.s64 	%rd22, %rd2, %rd21;
	add.s64 	%rd31, %rd22, 8;
	mul.wide.s32 	%rd5, %r15, 4;

$L__BB2_3:
	ld.global.f32 	%f12, [%rd32];
	ld.global.f32 	%f13, [%rd31+-8];
	fma.rn.f32 	%f14, %f13, %f12, %f29;
	add.s64 	%rd23, %rd32, %rd5;
	ld.global.f32 	%f15, [%rd23];
	ld.global.f32 	%f16, [%rd31+-4];
	fma.rn.f32 	%f17, %f16, %f15, %f14;
	add.s64 	%rd24, %rd23, %rd5;
	ld.global.f32 	%f18, [%rd24];
	ld.global.f32 	%f19, [%rd31];
	fma.rn.f32 	%f20, %f19, %f18, %f17;
	add.s64 	%rd25, %rd24, %rd5;
	add.s64 	%rd32, %rd25, %rd5;
	ld.global.f32 	%f21, [%rd25];
	ld.global.f32 	%f22, [%rd31+4];
	fma.rn.f32 	%f29, %f22, %f21, %f20;
	add.s32 	%r31, %r31, 4;
	add.s64 	%rd31, %rd31, 16;
	add.s32 	%r30, %r30, -4;
	setp.ne.s32 	%p3, %r30, 0;
	@%p3 bra 	$L__BB2_3;

$L__BB2_4:
	setp.eq.s32 	%p4, %r32, 0;
	@%p4 bra 	$L__BB2_7;

	mad.lo.s32 	%r27, %r31, %r15, %r2;
	mul.wide.s32 	%rd26, %r27, 4;
	add.s64 	%rd34, %rd1, %rd26;
	mul.wide.s32 	%rd11, %r15, 4;
	mad.lo.s32 	%r28, %r3, %r13, %r31;
	mul.wide.s32 	%rd27, %r28, 4;
	add.s64 	%rd33, %rd2, %rd27;

$L__BB2_6:
	.pragma "nounroll";
	ld.global.f32 	%f23, [%rd34];
	ld.global.f32 	%f24, [%rd33];
	fma.rn.f32 	%f29, %f24, %f23, %f29;
	add.s64 	%rd34, %rd34, %rd11;
	add.s64 	%rd33, %rd33, 4;
	add.s32 	%r32, %r32, -1;
	setp.ne.s32 	%p5, %r32, 0;
	@%p5 bra 	$L__BB2_6;

$L__BB2_7:
	cvta.to.global.u64 	%rd28, %rd17;
	mul.wide.s32 	%rd29, %r1, 4;
	add.s64 	%rd30, %rd28, %rd29;
	st.global.f32 	[%rd30], %f29;
	ret;

}
	// .globl	add_scalar
.visible .entry add_scalar(
	.param .u64 add_scalar_param_0,
	.param .u64 add_scalar_param_1,
	.param .f32 add_scalar_param_2
)
{
	.reg .f32 	%f<4>;
	.reg .b32 	%r<8>;
	.reg .b64 	%rd<8>;


	ld.param.u64 	%rd1, [add_scalar_param_0];
	ld.param.u64 	%rd2, [add_scalar_param_1];
	ld.param.f32 	%f1, [add_scalar_param_2];
	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mov.u32 	%r1, %tid.x;
	mov.u32 	%r2, %tid.y;
	mov.u32 	%r3, %ntid.x;
	mov.u32 	%r4, %ctaid.x;
	mov.u32 	%r5, %ntid.y;
	mad.lo.s32 	%r6, %r5, %r4, %r2;
	mad.lo.s32 	%r7, %r6, %r3, %r1;
	mul.wide.s32 	%rd5, %r7, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f2, [%rd6];
	add.f32 	%f3, %f2, %f1;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;
	ret;

}
	// .globl	sub_scalar
.visible .entry sub_scalar(
	.param .u64 sub_scalar_param_0,
	.param .u64 sub_scalar_param_1,
	.param .f32 sub_scalar_param_2
)
{
	.reg .f32 	%f<4>;
	.reg .b32 	%r<8>;
	.reg .b64 	%rd<8>;


	ld.param.u64 	%rd1, [sub_scalar_param_0];
	ld.param.u64 	%rd2, [sub_scalar_param_1];
	ld.param.f32 	%f1, [sub_scalar_param_2];
	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mov.u32 	%r1, %tid.x;
	mov.u32 	%r2, %tid.y;
	mov.u32 	%r3, %ntid.x;
	mov.u32 	%r4, %ctaid.x;
	mov.u32 	%r5, %ntid.y;
	mad.lo.s32 	%r6, %r5, %r4, %r2;
	mad.lo.s32 	%r7, %r6, %r3, %r1;
	mul.wide.s32 	%rd5, %r7, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f2, [%rd6];
	sub.f32 	%f3, %f2, %f1;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;
	ret;

}
	// .globl	mul_scalar
.visible .entry mul_scalar(
	.param .u64 mul_scalar_param_0,
	.param .u64 mul_scalar_param_1,
	.param .f32 mul_scalar_param_2
)
{
	.reg .f32 	%f<4>;
	.reg .b32 	%r<8>;
	.reg .b64 	%rd<8>;


	ld.param.u64 	%rd1, [mul_scalar_param_0];
	ld.param.u64 	%rd2, [mul_scalar_param_1];
	ld.param.f32 	%f1, [mul_scalar_param_2];
	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mov.u32 	%r1, %tid.x;
	mov.u32 	%r2, %tid.y;
	mov.u32 	%r3, %ntid.x;
	mov.u32 	%r4, %ctaid.x;
	mov.u32 	%r5, %ntid.y;
	mad.lo.s32 	%r6, %r5, %r4, %r2;
	mad.lo.s32 	%r7, %r6, %r3, %r1;
	mul.wide.s32 	%rd5, %r7, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f2, [%rd6];
	mul.f32 	%f3, %f2, %f1;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;
	ret;

}
	// .globl	div_scalar
.visible .entry div_scalar(
	.param .u64 div_scalar_param_0,
	.param .u64 div_scalar_param_1,
	.param .f32 div_scalar_param_2
)
{
	.reg .f32 	%f<4>;
	.reg .b32 	%r<8>;
	.reg .b64 	%rd<8>;


	ld.param.u64 	%rd1, [div_scalar_param_0];
	ld.param.u64 	%rd2, [div_scalar_param_1];
	ld.param.f32 	%f1, [div_scalar_param_2];
	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mov.u32 	%r1, %tid.x;
	mov.u32 	%r2, %tid.y;
	mov.u32 	%r3, %ntid.x;
	mov.u32 	%r4, %ctaid.x;
	mov.u32 	%r5, %ntid.y;
	mad.lo.s32 	%r6, %r5, %r4, %r2;
	mad.lo.s32 	%r7, %r6, %r3, %r1;
	mul.wide.s32 	%rd5, %r7, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f2, [%rd6];
	div.rn.f32 	%f3, %f2, %f1;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;
	ret;

}
	// .globl	scalar_sub
.visible .entry scalar_sub(
	.param .u64 scalar_sub_param_0,
	.param .u64 scalar_sub_param_1,
	.param .f32 scalar_sub_param_2
)
{
	.reg .f32 	%f<4>;
	.reg .b32 	%r<8>;
	.reg .b64 	%rd<8>;


	ld.param.u64 	%rd1, [scalar_sub_param_0];
	ld.param.u64 	%rd2, [scalar_sub_param_1];
	ld.param.f32 	%f1, [scalar_sub_param_2];
	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mov.u32 	%r1, %tid.x;
	mov.u32 	%r2, %tid.y;
	mov.u32 	%r3, %ntid.x;
	mov.u32 	%r4, %ctaid.x;
	mov.u32 	%r5, %ntid.y;
	mad.lo.s32 	%r6, %r5, %r4, %r2;
	mad.lo.s32 	%r7, %r6, %r3, %r1;
	mul.wide.s32 	%rd5, %r7, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f2, [%rd6];
	sub.f32 	%f3, %f1, %f2;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;
	ret;

}
	// .globl	scalar_div
.visible .entry scalar_div(
	.param .u64 scalar_div_param_0,
	.param .u64 scalar_div_param_1,
	.param .f32 scalar_div_param_2
)
{
	.reg .f32 	%f<4>;
	.reg .b32 	%r<8>;
	.reg .b64 	%rd<8>;


	ld.param.u64 	%rd1, [scalar_div_param_0];
	ld.param.u64 	%rd2, [scalar_div_param_1];
	ld.param.f32 	%f1, [scalar_div_param_2];
	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mov.u32 	%r1, %tid.x;
	mov.u32 	%r2, %tid.y;
	mov.u32 	%r3, %ntid.x;
	mov.u32 	%r4, %ctaid.x;
	mov.u32 	%r5, %ntid.y;
	mad.lo.s32 	%r6, %r5, %r4, %r2;
	mad.lo.s32 	%r7, %r6, %r3, %r1;
	mul.wide.s32 	%rd5, %r7, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f2, [%rd6];
	div.rn.f32 	%f3, %f1, %f2;
	add.s64 	%rd7, %rd3, %rd5;
	st.global.f32 	[%rd7], %f3;
	ret;

}