#include <stdio.h>
__global__
void k_lu_f32(int n, float* A, int stride_row, int stride_col){
	for (int i=0;i<n;i++){
		for (int j=i+1;j<n;j++){
			float factor=-A[j*stride_col+i*stride_row]/A[i*stride_col+i*stride_row];
			for (int k=i+1;k<n;k++){
				A[j*stride_col+k*stride_row]=A[j*stride_col+k*stride_row]+A[i*stride_col+k*stride_row]*factor;
			}
			A[j*stride_col+i*stride_row]=-factor;
		}
	}
}


//Solves LX=A, whereas L is a lower triangular matrix and X an unknown Matrix. It is assumed that L is stored inside A=LU with ones on the diagonal
__global__
void k_lu_solve_lower_f32(int n, int m, const float* L, int stride_row_l,int stride_col_l, const float* A, int stride_row_a, int stride_col_a, float* X, int stride_row_x, int stride_col_x){
	const int TILE_SIZE=64;
	int tx=threadIdx.x;
	int bx=blockIdx.x;
	int bpx=(bx*TILE_SIZE+tx)*stride_row_x;
	int bpa=(bx*TILE_SIZE+tx)*stride_row_a;//weitermachen
	if ((bx*TILE_SIZE+tx)<m){
		for (int i=0;i<n;i++){
			float sum=A[bpa+i*stride_col_a];
			for (int j=0;j<i;j++){
				sum-=L[i*stride_col_l+j*stride_row_l]*X[bpx+j*stride_col_x];
			}
			X[bpx+i*stride_col_x]=sum;
		}
	}
}

//Solves UX=A, whereas U is a upprt triangular matrix coming from lu factorization and X an unknown Matrix
__global__
void k_lu_solve_upper_f32(int n, int m,const float* U, int stride_row_u,int stride_col_u, const float* A, int stride_row_a, int stride_col_a, float* X, int stride_row_x, int stride_col_x){
	const int TILE_SIZE=64;
	int tx=threadIdx.x;
	int bx=blockIdx.x;
	int bpx=(bx*TILE_SIZE+tx)*stride_row_x;
	int bpa=(bx*TILE_SIZE+tx)*stride_row_a;
	if ((bx*TILE_SIZE+tx)<m){
		for (int i=n-1;i>=0;i--){
			float sum=A[bpa+i*stride_col_a];
			
			for (int j=n-1;j>i;j--){
				sum-=U[i*stride_col_u+j*stride_row_u]*X[bpx+j*stride_col_x];
				
			}
			sum/=U[i*stride_col_u+i*stride_row_u];
			//printf("sum:%.7f and U:%.g and i%d\n",sum,U[i*stride_col_u+i*stride_row_u], i);
			X[bpx+i*stride_col_x]=sum;
		}
	}
}

//Solves AX=B
__host__
void lu_solve_f32_device(int n, int m, float* A_d, int stride_row_a, int stride_col_a, float* B_d, int stride_row_b, int stride_col_b, float* X_d, int stride_row_x, int stride_col_x){
	float bsmx=64;
	float* Y;
	cudaMalloc((void**)&Y,sizeof(float)*n*m);
	k_lu_solve_lower_f32<<<ceil(m/bsmx),bsmx>>>(n,m,A_d,stride_row_a,stride_col_a,B_d,stride_row_b,stride_col_b,Y,1,m);
	k_lu_solve_upper_f32<<<ceil(m/bsmx),1>>>(n,m,A_d,stride_row_a,stride_col_a,Y,1,m,X_d,stride_row_x,stride_col_x); //TODO: replace with non kernel version
	cudaFree(Y);
}
	.file	"tmpxft_003436fe_00000000-6_lu.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z30__device_stub__Z8k_lu_f32iPfiiiPfii
	.type	_Z30__device_stub__Z8k_lu_f32iPfiiiPfii, @function
_Z30__device_stub__Z8k_lu_f32iPfiiiPfii:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 24(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 12(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z8k_lu_f32iPfii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z30__device_stub__Z8k_lu_f32iPfiiiPfii, .-_Z30__device_stub__Z8k_lu_f32iPfiiiPfii
	.globl	_Z8k_lu_f32iPfii
	.type	_Z8k_lu_f32iPfii, @function
_Z8k_lu_f32iPfii:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z30__device_stub__Z8k_lu_f32iPfiiiPfii
	.cfi_endproc
.LFE2053:
	.size	_Z8k_lu_f32iPfii, .-_Z8k_lu_f32iPfii
	.globl	_Z54__device_stub__Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii
	.type	_Z54__device_stub__Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii, @function
_Z54__device_stub__Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	240(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 32(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 28(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	248(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	256(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	56(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfii(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 240
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L8:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z54__device_stub__Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii, .-_Z54__device_stub__Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii
	.globl	_Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfii
	.type	_Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfii, @function
_Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfii:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z54__device_stub__Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii
	.cfi_endproc
.LFE2055:
	.size	_Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfii, .-_Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfii
	.globl	_Z54__device_stub__Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii
	.type	_Z54__device_stub__Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii, @function
_Z54__device_stub__Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii:
.LFB2056:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	240(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 32(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 28(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	248(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	256(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	56(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfii(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 240
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L13:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2056:
	.size	_Z54__device_stub__Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii, .-_Z54__device_stub__Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii
	.globl	_Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfii
	.type	_Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfii, @function
_Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfii:
.LFB2057:
	.cfi_startproc
	endbr64
	jmp	_Z54__device_stub__Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii
	.cfi_endproc
.LFE2057:
	.size	_Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfii, .-_Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfii
	.globl	_Z19lu_solve_f32_deviceiiPfiiS_iiS_ii
	.type	_Z19lu_solve_f32_deviceiiPfiiS_iiS_ii, @function
_Z19lu_solve_f32_deviceiiPfiiS_iiS_ii:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%r9, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%r8d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%edi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%esi, %ebx
	movslq	%edi, %rsi
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movl	%ecx, 12(%rsp)
	leaq	24(%rsp), %rdi
	movq	144(%rsp), %r15
	movq	%rdx, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	movslq	%ebx, %rax
	imulq	%rax, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	cvtsi2ssl	%ebx, %xmm0
	mulss	.LC0(%rip), %xmm0
	movl	$67108865, %edx
	salq	$6, %rdx
	movl	$1, 52(%rsp)
	movq	%rdx, 44(%rsp)
	call	ceilf@PLT
	movl	$67108865, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	cvttss2siq	%xmm0, %r14
	salq	$6, %rdx
	movl	$1, %ecx
	movabsq	$4294967297, %rax
	movq	%rax, 36(%rsp)
	movl	40(%rsp), %esi
	movl	%r14d, 32(%rsp)
	movq	32(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	pushq	%rdx
	.cfi_def_cfa_offset 136
	movq	%r13, %r9
	movl	%r12d, %r8d
	movl	%ebx, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 144
	movl	%ebp, %edi
	pushq	$1
	.cfi_def_cfa_offset 152
	pushq	48(%rsp)
	.cfi_def_cfa_offset 160
	movl	168(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 168
	movl	168(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 176
	movl	60(%rsp), %ecx
	movq	48(%rsp), %rdx
	call	_Z54__device_stub__Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii
	addq	$48, %rsp
	.cfi_def_cfa_offset 128
.L19:
	movl	%r14d, 32(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$4294967297, %rdx
	movl	$1, 52(%rsp)
	movq	%rdx, 36(%rsp)
	movq	32(%rsp), %rdi
	movl	40(%rsp), %esi
	movq	%rdx, 44(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L20
	pushq	%rax
	.cfi_def_cfa_offset 136
	movl	168(%rsp), %eax
	movl	%r12d, %r8d
	movl	%ebx, %esi
	movl	%ebp, %edi
	pushq	%rax
	.cfi_def_cfa_offset 144
	movl	168(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 152
	pushq	%r15
	.cfi_def_cfa_offset 160
	pushq	%rbx
	.cfi_def_cfa_offset 168
	pushq	$1
	.cfi_def_cfa_offset 176
	movq	72(%rsp), %r9
	movl	60(%rsp), %ecx
	movq	48(%rsp), %rdx
	call	_Z54__device_stub__Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfiiiiPKfiiS0_iiPfii
	addq	$48, %rsp
	.cfi_def_cfa_offset 128
.L20:
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z19lu_solve_f32_deviceiiPfiiS_iiS_ii, .-_Z19lu_solve_f32_deviceiiPfiiS_iiS_ii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"_Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfii"
.LC2:
	.string	"_Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfii"
.LC3:
	.string	"_Z8k_lu_f32iPfii"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2059:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC1(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z20k_lu_solve_upper_f32iiPKfiiS0_iiPfii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z20k_lu_solve_lower_f32iiPKfiiS0_iiPfii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z8k_lu_f32iPfii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2059:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC0:
	.long	1015021568
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#pragma once
#include <hip/hip_runtime.h>


#ifdef __HIPCC__
__global__
#endif
void k_lu_f32(int n, float* A, int stride_row, int stride_col);

#ifdef __HIPCC__
__host__
#endif
void lu_solve_f32_device(int n, int m, float* A_d, int stride_row_a, int stride_col_a, float* B_d, int stride_row_b, int stride_col_b, float* X_d, int stride_row_x, int stride_col_x);
	.text
	.file	"lu.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
