#include <math.h>
#include <iostream>
#include "device_launch_parameters.h"
#include "cuda_runtime.h"
#include "chebyshev.cuh"
#include "kernel.cu"
#include <string>
#include <stdlib.h>
#include <math.h>

#define Im1 0               // i - 1
#define I 1                 // i
#define Ip1 2               // i + 1

using namespace std;

double* chebyshev(int matrix_size, double** Ab, int s, int max_iter) {

	int N = matrix_size;
	dim3 threadsPerBlock(N);
	dim3 blocksPerGrid(1);
	if (N > 512) {
		threadsPerBlock.x = 512;
		blocksPerGrid.x = ceil(double(N) / double(threadsPerBlock.x));
	}

	double delta, x_2_norm, a, w_0, c, L, B, scalar_1, scalar_2;
	double *x_start, *A_vector, *b_vector, *w, *norm_vector, *rows;;										//HOST	
	double *d_x, *d_scalar_1, *d_scalar_2, *d_temp_vector, *d_x_start, *d_A_vector, *d_b_vector;    //DEVICE
	int iteration, k;

	delta = 0.000001;       // accuracy
	x_2_norm = 0;           // second norm for stop criteria
	a = 100;                // alpha
	bool stop = false;      // stop criteria bool

	//ALLOCATE MEMORY ON HOST
	x_start = (double *)malloc(N * sizeof(double));
	A_vector = (double *)malloc(N * N * sizeof(double));
	b_vector = (double *)malloc(N * sizeof(double));
	w = (double *)malloc(N * 2 * sizeof(double));
	norm_vector = (double *)malloc(N * sizeof(double));
	rows = (double *)malloc(3 * N * sizeof(double));

	//ALLOCATE MEMORY ON DEVICE
	cudaMalloc((void **)&d_x, N * 3 * sizeof(double));
	cudaMalloc((void **)&d_scalar_1, sizeof(double));
	cudaMalloc((void **)&d_scalar_2, sizeof(double));
	cudaMalloc((void **)&d_temp_vector, N * sizeof(double));
	cudaMalloc((void **)&d_x_start, N * sizeof(double));
	cudaMalloc((void **)&d_A_vector, N * N * sizeof(double));
	cudaMalloc((void **)&d_b_vector, N * sizeof(double));


	B = Ab[0][0];

	//init x_start, find Beta, copy Ab to vectors
	for (int i = 0; i < N; i++) {
		x_start[i] = 0;
		b_vector[i] = Ab[i][N];
		if (Ab[i][i] > B) {
			B = Ab[i][i];
		}
		for (int j = 0; j < N; j++) {
			A_vector[i * N + j] = Ab[i][j];
		}
	}

	B = 2 * B;

	//COPY FROM HOST TO DEVICE
	cudaMemcpy(d_x_start, x_start, N * sizeof(double), cudaMemcpyHostToDevice);
	cudaMemcpy(d_A_vector, A_vector, N * N * sizeof(double), cudaMemcpyHostToDevice);
	cudaMemcpy(d_b_vector, b_vector, N * sizeof(double), cudaMemcpyHostToDevice);

	//Step 0:
	iteration = 0;
	w_0 = (B - a) / (B + a);
	c = 2 / (B + a);
	L = 2 * (B + a) / (B - a);

	copyVectorToMatRowKernel << <blocksPerGrid, threadsPerBlock >> >(d_x_start, d_x, Im1, N);

	while (iteration < max_iter && stop == false) {
		//Step 1
		k = 0;

		copyVectorToMatRowKernel <<<blocksPerGrid, threadsPerBlock >>>(d_x_start, d_x, I, N);

		w[Im1] = 0;
		w[I] = w_0;

		while (iteration < max_iter) {
			//Step 2
			scalar_1 = w[I] * w[Im1];
			scalar_2 = c * (1 + w[I] * w[Im1]);
			cudaMemcpy(d_scalar_1, &scalar_1, sizeof(double), cudaMemcpyHostToDevice);
			cudaMemcpy(d_scalar_2, &scalar_2, sizeof(double), cudaMemcpyHostToDevice);

			//copies x(i) to temp_vector
			copyMatRowToVectorKernel <<<blocksPerGrid, threadsPerBlock>>>(d_x, d_temp_vector, I, N);

			//multiples A matrix in vector form with x(i) stored in temp_vector
			multMatrixVectorKernel <<<blocksPerGrid, threadsPerBlock>>>(d_A_vector, d_temp_vector, d_temp_vector, N);

			//calculates x(i+1) and sets x(i-1) and x(i)
			calculateXplus1Kernel <<<blocksPerGrid, threadsPerBlock>>>(d_x, d_temp_vector, d_b_vector, d_scalar_1, d_scalar_2, N);

			w[Im1] = w[I];
			w[I] = 1 / (L - w[Im1]);
			
			x_2_norm = 0;

			normKernel <<<blocksPerGrid, threadsPerBlock >> >(d_x, d_temp_vector, N);

			cudaMemcpy(norm_vector, d_temp_vector, N * sizeof(double), cudaMemcpyDeviceToHost);

			for (int i = 0; i < N; i++) {
				x_2_norm += norm_vector[i];
			}
			x_2_norm = sqrt(x_2_norm);
			if (x_2_norm < delta) {
				stop = true;
				break;
			}

			// Step 3
			iteration++;
			k++;
			if (k >= s) {
				copyMatRowToVectorKernel << <blocksPerGrid, threadsPerBlock >> >(d_x, d_x_start, I, N);
				break;
			}
		}
	}
	copyMatRowToVectorKernel << <blocksPerGrid, threadsPerBlock >> >(d_x, d_temp_vector, I, N);
	cudaMemcpy(x_start, d_temp_vector, N * sizeof(double), cudaMemcpyDeviceToHost);

	cudaFree(d_x);
	cudaFree(d_scalar_1);
	cudaFree(d_scalar_2);
	cudaFree(d_temp_vector);
	cudaFree(d_x_start);
	cudaFree(d_A_vector);
	cudaFree(d_b_vector);

	free(A_vector);
	free(b_vector);
	free(w);
	free(norm_vector);
	cout << "iters: " << iteration << endl;
	return x_start;
}




	.file	"tmpxft_00302f25_00000000-6_chebyshev.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3639:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3639:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z47__device_stub__Z22multMatrixVectorKernelPdS_S_iPdS_S_i
	.type	_Z47__device_stub__Z22multMatrixVectorKernelPdS_S_iPdS_S_i, @function
_Z47__device_stub__Z22multMatrixVectorKernelPdS_S_iPdS_S_i:
.LFB3661:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%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	_Z22multMatrixVectorKernelPdS_S_i(%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
.LFE3661:
	.size	_Z47__device_stub__Z22multMatrixVectorKernelPdS_S_iPdS_S_i, .-_Z47__device_stub__Z22multMatrixVectorKernelPdS_S_iPdS_S_i
	.globl	_Z22multMatrixVectorKernelPdS_S_i
	.type	_Z22multMatrixVectorKernelPdS_S_i, @function
_Z22multMatrixVectorKernelPdS_S_i:
.LFB3662:
	.cfi_startproc
	endbr64
	jmp	_Z47__device_stub__Z22multMatrixVectorKernelPdS_S_iPdS_S_i
	.cfi_endproc
.LFE3662:
	.size	_Z22multMatrixVectorKernelPdS_S_i, .-_Z22multMatrixVectorKernelPdS_S_i
	.globl	_Z48__device_stub__Z24copyVectorToMatRowKernelPdS_iiPdS_ii
	.type	_Z48__device_stub__Z24copyVectorToMatRowKernelPdS_iiPdS_ii, @function
_Z48__device_stub__Z24copyVectorToMatRowKernelPdS_iiPdS_ii:
.LFB3663:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%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	.L8
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z24copyVectorToMatRowKernelPdS_ii(%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
.L8:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3663:
	.size	_Z48__device_stub__Z24copyVectorToMatRowKernelPdS_iiPdS_ii, .-_Z48__device_stub__Z24copyVectorToMatRowKernelPdS_iiPdS_ii
	.globl	_Z24copyVectorToMatRowKernelPdS_ii
	.type	_Z24copyVectorToMatRowKernelPdS_ii, @function
_Z24copyVectorToMatRowKernelPdS_ii:
.LFB3664:
	.cfi_startproc
	endbr64
	jmp	_Z48__device_stub__Z24copyVectorToMatRowKernelPdS_iiPdS_ii
	.cfi_endproc
.LFE3664:
	.size	_Z24copyVectorToMatRowKernelPdS_ii, .-_Z24copyVectorToMatRowKernelPdS_ii
	.globl	_Z48__device_stub__Z24copyMatRowToVectorKernelPdS_iiPdS_ii
	.type	_Z48__device_stub__Z24copyMatRowToVectorKernelPdS_iiPdS_ii, @function
_Z48__device_stub__Z24copyMatRowToVectorKernelPdS_iiPdS_ii:
.LFB3665:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%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	.L13
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z24copyMatRowToVectorKernelPdS_ii(%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
.L13:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3665:
	.size	_Z48__device_stub__Z24copyMatRowToVectorKernelPdS_iiPdS_ii, .-_Z48__device_stub__Z24copyMatRowToVectorKernelPdS_iiPdS_ii
	.globl	_Z24copyMatRowToVectorKernelPdS_ii
	.type	_Z24copyMatRowToVectorKernelPdS_ii, @function
_Z24copyMatRowToVectorKernelPdS_ii:
.LFB3666:
	.cfi_startproc
	endbr64
	jmp	_Z48__device_stub__Z24copyMatRowToVectorKernelPdS_iiPdS_ii
	.cfi_endproc
.LFE3666:
	.size	_Z24copyMatRowToVectorKernelPdS_ii, .-_Z24copyMatRowToVectorKernelPdS_ii
	.globl	_Z18copyMatRowToVectorPdS_ii
	.type	_Z18copyMatRowToVectorPdS_ii, @function
_Z18copyMatRowToVectorPdS_ii:
.LFB3635:
	.cfi_startproc
	endbr64
	movabsq	$4294967297, %rax
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	%edx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%ecx, %ebx
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	movq	%rax, 20(%rsp)
	movl	%ecx, %eax
	movl	$1, 12(%rsp)
	movl	$1, 16(%rsp)
	movl	$1, 28(%rsp)
	cmpl	$512, %ecx
	jle	.L20
	cvtsi2sdl	%ecx, %xmm0
	mulsd	.LC0(%rip), %xmm0
	call	ceil@PLT
	cvttsd2siq	%xmm0, %rax
	movl	%eax, 20(%rsp)
	movl	$512, %eax
.L20:
	movl	%eax, 8(%rsp)
	movl	16(%rsp), %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	8(%rsp), %rdx
	movq	20(%rsp), %rdi
	movl	28(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 40
	movl	%ebx, %ecx
	movl	%r13d, %edx
	movq	%r12, %rsi
	popq	%rbx
	.cfi_def_cfa_offset 32
	movq	%rbp, %rdi
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	_Z48__device_stub__Z24copyMatRowToVectorKernelPdS_iiPdS_ii
.L18:
	.cfi_restore_state
	addq	$40, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3635:
	.size	_Z18copyMatRowToVectorPdS_ii, .-_Z18copyMatRowToVectorPdS_ii
	.globl	_Z50__device_stub__Z21calculateXplus1KernelPdS_S_S_S_iPdS_S_S_S_i
	.type	_Z50__device_stub__Z21calculateXplus1KernelPdS_S_S_S_iPdS_S_S_S_i, @function
_Z50__device_stub__Z21calculateXplus1KernelPdS_S_S_S_iPdS_S_S_S_i:
.LFB3667:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movl	%r9d, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z21calculateXplus1KernelPdS_S_S_S_i(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L23:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3667:
	.size	_Z50__device_stub__Z21calculateXplus1KernelPdS_S_S_S_iPdS_S_S_S_i, .-_Z50__device_stub__Z21calculateXplus1KernelPdS_S_S_S_iPdS_S_S_S_i
	.globl	_Z21calculateXplus1KernelPdS_S_S_S_i
	.type	_Z21calculateXplus1KernelPdS_S_S_S_i, @function
_Z21calculateXplus1KernelPdS_S_S_S_i:
.LFB3668:
	.cfi_startproc
	endbr64
	jmp	_Z50__device_stub__Z21calculateXplus1KernelPdS_S_S_S_iPdS_S_S_S_i
	.cfi_endproc
.LFE3668:
	.size	_Z21calculateXplus1KernelPdS_S_S_S_i, .-_Z21calculateXplus1KernelPdS_S_S_S_i
	.globl	_Z33__device_stub__Z10normKernelPdS_iPdS_i
	.type	_Z33__device_stub__Z10normKernelPdS_iPdS_i, @function
_Z33__device_stub__Z10normKernelPdS_iPdS_i:
.LFB3669:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z10normKernelPdS_i(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L28:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L30
	call	__stack_chk_fail@PLT
.L30:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3669:
	.size	_Z33__device_stub__Z10normKernelPdS_iPdS_i, .-_Z33__device_stub__Z10normKernelPdS_iPdS_i
	.globl	_Z10normKernelPdS_i
	.type	_Z10normKernelPdS_i, @function
_Z10normKernelPdS_i:
.LFB3670:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z10normKernelPdS_iPdS_i
	.cfi_endproc
.LFE3670:
	.size	_Z10normKernelPdS_i, .-_Z10normKernelPdS_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC6:
	.string	"iters: "
	.text
	.globl	_Z9chebysheviPPdii
	.type	_Z9chebysheviPPdii, @function
_Z9chebysheviPPdii:
.LFB3636:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movl	%edi, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%edi, %ebx
	subq	$200, %rsp
	.cfi_def_cfa_offset 256
	movl	%edx, 76(%rsp)
	movl	%ecx, 72(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	movabsq	$4294967297, %rax
	movl	$1, 164(%rsp)
	movl	$1, 168(%rsp)
	movq	%rax, 172(%rsp)
	movl	$1, 180(%rsp)
	cmpl	$512, %edi
	jle	.L35
	cvtsi2sdl	%edi, %xmm0
	mulsd	.LC0(%rip), %xmm0
	movl	$512, %r14d
	call	ceil@PLT
	cvttsd2siq	%xmm0, %rax
	movl	%eax, 172(%rsp)
.L35:
	movl	%ebx, %r13d
	movslq	%ebx, %rbp
	imull	%ebx, %r13d
	salq	$3, %rbp
	movq	%rbp, %rdi
	call	malloc@PLT
	movslq	%r13d, %r13
	movq	%rax, (%rsp)
	salq	$3, %r13
	movq	%r13, %rdi
	call	malloc@PLT
	movq	%rbp, %rdi
	movq	%rax, 16(%rsp)
	call	malloc@PLT
	leal	(%rbx,%rbx), %edi
	movslq	%edi, %rdi
	movq	%rax, 24(%rsp)
	salq	$3, %rdi
	call	malloc@PLT
	movq	%rbp, %rdi
	movq	%rax, %r12
	call	malloc@PLT
	leal	(%rbx,%rbx,2), %esi
	leaq	104(%rsp), %rdi
	movslq	%esi, %rsi
	movq	%rax, 32(%rsp)
	salq	$3, %rsi
	call	cudaMalloc@PLT
	leaq	112(%rsp), %rdi
	movl	$8, %esi
	call	cudaMalloc@PLT
	leaq	120(%rsp), %rdi
	movl	$8, %esi
	call	cudaMalloc@PLT
	leaq	128(%rsp), %rdi
	movq	%rbp, %rsi
	call	cudaMalloc@PLT
	leaq	136(%rsp), %rdi
	movq	%rbp, %rsi
	call	cudaMalloc@PLT
	leaq	144(%rsp), %rdi
	movq	%r13, %rsi
	call	cudaMalloc@PLT
	leaq	152(%rsp), %rdi
	movq	%rbp, %rsi
	call	cudaMalloc@PLT
	movq	(%r15), %rax
	movq	16(%rsp), %rcx
	movsd	(%rax), %xmm0
	xorl	%eax, %eax
.L36:
	cmpl	%eax, %ebx
	jle	.L61
	movq	(%rsp), %rdx
	movq	(%r15,%rax,8), %rsi
	movq	$0x000000000, (%rdx,%rax,8)
	movq	24(%rsp), %rdx
	movsd	(%rsi,%rbp), %xmm1
	movsd	%xmm1, (%rdx,%rax,8)
	movsd	(%rsi,%rax,8), %xmm1
	xorl	%edx, %edx
	maxsd	%xmm0, %xmm1
.L38:
	movsd	(%rsi,%rdx,8), %xmm0
	movsd	%xmm0, (%rcx,%rdx,8)
	incq	%rdx
	cmpl	%edx, %ebx
	jg	.L38
	incq	%rax
	addq	%rbp, %rcx
	movaps	%xmm1, %xmm0
	jmp	.L36
.L61:
	addsd	%xmm0, %xmm0
	movq	(%rsp), %rsi
	movl	$1, %ecx
	movq	%rbp, %rdx
	movq	136(%rsp), %rdi
	movsd	%xmm0, 8(%rsp)
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rsi
	movl	$1, %ecx
	movq	%r13, %rdx
	movq	144(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rsi
	movl	$1, %ecx
	movq	%rbp, %rdx
	movq	152(%rsp), %rdi
	call	cudaMemcpy@PLT
	movsd	8(%rsp), %xmm0
	xorl	%r9d, %r9d
	movsd	.LC2(%rip), %xmm2
	movl	%r14d, 160(%rsp)
	movl	168(%rsp), %ecx
	xorl	%r8d, %r8d
	movaps	%xmm0, %xmm1
	addsd	%xmm2, %xmm0
	movq	160(%rsp), %rdx
	movq	172(%rsp), %rdi
	subsd	%xmm2, %xmm1
	movsd	.LC3(%rip), %xmm2
	movl	180(%rsp), %esi
	divsd	%xmm0, %xmm2
	movaps	%xmm1, %xmm3
	divsd	%xmm0, %xmm3
	addsd	%xmm0, %xmm0
	movsd	%xmm2, 56(%rsp)
	divsd	%xmm1, %xmm0
	movsd	%xmm3, 48(%rsp)
	movsd	%xmm0, 64(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L40
	movq	104(%rsp), %rsi
	movq	136(%rsp), %rdi
	movl	%ebx, %ecx
	xorl	%edx, %edx
	call	_Z48__device_stub__Z24copyVectorToMatRowKernelPdS_iiPdS_ii
.L40:
	movb	$0, 47(%rsp)
	xorl	%r13d, %r13d
.L41:
	movl	72(%rsp), %eax
	movl	%r14d, 160(%rsp)
	movl	$0, %r9d
	movl	$0, %r8d
	movl	168(%rsp), %ecx
	movq	172(%rsp), %rdi
	cmpl	%eax, %r13d
	movl	180(%rsp), %esi
	movq	160(%rsp), %rdx
	setge	%al
	orb	%al, 47(%rsp)
	jne	.L62
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L42
	movq	104(%rsp), %rsi
	movq	136(%rsp), %rdi
	movl	%ebx, %ecx
	movl	$1, %edx
	call	_Z48__device_stub__Z24copyVectorToMatRowKernelPdS_iiPdS_ii
.L42:
	movsd	48(%rsp), %xmm4
	movq	$0x000000000, (%r12)
	movl	%r13d, %r15d
	movsd	%xmm4, 8(%r12)
.L53:
	movsd	(%r12), %xmm0
	movq	112(%rsp), %rdi
	leaq	88(%rsp), %rsi
	movl	$1, %ecx
	movsd	8(%r12), %xmm5
	movl	$8, %edx
	mulsd	%xmm5, %xmm0
	movsd	%xmm5, 8(%rsp)
	movsd	%xmm0, 88(%rsp)
	addsd	.LC4(%rip), %xmm0
	mulsd	56(%rsp), %xmm0
	movsd	%xmm0, 96(%rsp)
	call	cudaMemcpy@PLT
	movq	120(%rsp), %rdi
	leaq	96(%rsp), %rsi
	movl	$1, %ecx
	movl	$8, %edx
	call	cudaMemcpy@PLT
	movl	%r14d, 160(%rsp)
	xorl	%r9d, %r9d
	movl	168(%rsp), %ecx
	movq	160(%rsp), %rdx
	movq	172(%rsp), %rdi
	xorl	%r8d, %r8d
	movl	180(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L43
	movq	128(%rsp), %rsi
	movq	104(%rsp), %rdi
	movl	%ebx, %ecx
	movl	$1, %edx
	call	_Z48__device_stub__Z24copyMatRowToVectorKernelPdS_iiPdS_ii
.L43:
	movl	168(%rsp), %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	160(%rsp), %rdx
	movq	172(%rsp), %rdi
	movl	180(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L44
	movq	128(%rsp), %rsi
	movq	144(%rsp), %rdi
	movl	%ebx, %ecx
	movq	%rsi, %rdx
	call	_Z47__device_stub__Z22multMatrixVectorKernelPdS_S_iPdS_S_i
.L44:
	movl	168(%rsp), %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	160(%rsp), %rdx
	movq	172(%rsp), %rdi
	movl	180(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L45
	movq	120(%rsp), %r8
	movq	112(%rsp), %rcx
	movl	%ebx, %r9d
	movq	152(%rsp), %rdx
	movq	128(%rsp), %rsi
	movq	104(%rsp), %rdi
	call	_Z50__device_stub__Z21calculateXplus1KernelPdS_S_S_S_iPdS_S_S_S_i
.L45:
	movsd	8(%rsp), %xmm7
	movsd	64(%rsp), %xmm1
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movsd	.LC4(%rip), %xmm0
	movl	168(%rsp), %ecx
	subsd	%xmm7, %xmm1
	movq	160(%rsp), %rdx
	movq	172(%rsp), %rdi
	movsd	%xmm7, (%r12)
	movl	180(%rsp), %esi
	divsd	%xmm1, %xmm0
	movsd	%xmm0, 8(%r12)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L46
	movq	128(%rsp), %rsi
	movq	104(%rsp), %rdi
	movl	%ebx, %edx
	call	_Z33__device_stub__Z10normKernelPdS_iPdS_i
.L46:
	movq	128(%rsp), %rsi
	movq	32(%rsp), %rdi
	movl	$2, %ecx
	movq	%rbp, %rdx
	call	cudaMemcpy@PLT
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
.L47:
	cmpl	%eax, %ebx
	jle	.L63
	movq	32(%rsp), %rdx
	addsd	(%rdx,%rax,8), %xmm0
	incq	%rax
	jmp	.L47
.L63:
	call	sqrt@PLT
	movsd	.LC5(%rip), %xmm1
	comisd	%xmm0, %xmm1
	ja	.L57
	incl	%r15d
	movl	%r15d, %eax
	subl	%r13d, %eax
	cmpl	%eax, 76(%rsp)
	jg	.L50
	movl	168(%rsp), %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	160(%rsp), %rdx
	movq	172(%rsp), %rdi
	movl	180(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L49
	movq	136(%rsp), %rsi
	movq	104(%rsp), %rdi
	movl	%ebx, %ecx
	movl	$1, %edx
	call	_Z48__device_stub__Z24copyMatRowToVectorKernelPdS_iiPdS_ii
	jmp	.L49
.L50:
	cmpl	%r15d, 72(%rsp)
	jg	.L53
	jmp	.L49
.L57:
	movb	$1, 47(%rsp)
.L49:
	movl	%r15d, %r13d
	jmp	.L41
.L62:
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L55
	movq	128(%rsp), %rsi
	movq	104(%rsp), %rdi
	movl	%ebx, %ecx
	movl	$1, %edx
	call	_Z48__device_stub__Z24copyMatRowToVectorKernelPdS_iiPdS_ii
.L55:
	movq	128(%rsp), %rsi
	movq	(%rsp), %rdi
	movl	$2, %ecx
	movq	%rbp, %rdx
	call	cudaMemcpy@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	movq	112(%rsp), %rdi
	call	cudaFree@PLT
	movq	120(%rsp), %rdi
	call	cudaFree@PLT
	movq	128(%rsp), %rdi
	call	cudaFree@PLT
	movq	136(%rsp), %rdi
	call	cudaFree@PLT
	movq	144(%rsp), %rdi
	call	cudaFree@PLT
	movq	152(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	free@PLT
	movq	24(%rsp), %rdi
	call	free@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	32(%rsp), %rdi
	call	free@PLT
	leaq	.LC6(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r13d, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L56
	call	__stack_chk_fail@PLT
.L56:
	movq	(%rsp), %rax
	addq	$200, %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
.LFE3636:
	.size	_Z9chebysheviPPdii, .-_Z9chebysheviPPdii
	.section	.rodata.str1.1
.LC7:
	.string	"_Z10normKernelPdS_i"
.LC8:
	.string	"_Z21calculateXplus1KernelPdS_S_S_S_i"
.LC9:
	.string	"_Z24copyMatRowToVectorKernelPdS_ii"
.LC10:
	.string	"_Z24copyVectorToMatRowKernelPdS_ii"
.LC11:
	.string	"_Z22multMatrixVectorKernelPdS_S_i"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3672:
	.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	.LC7(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z10normKernelPdS_i(%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	.LC8(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z21calculateXplus1KernelPdS_S_S_S_i(%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	.LC9(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z24copyMatRowToVectorKernelPdS_ii(%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	.LC10(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z24copyVectorToMatRowKernelPdS_ii(%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	.LC11(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z22multMatrixVectorKernelPdS_S_i(%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
.LFE3672:
	.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.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.long	0
	.long	1063256064
	.align 8
.LC2:
	.long	0
	.long	1079574528
	.align 8
.LC3:
	.long	0
	.long	1073741824
	.align 8
.LC4:
	.long	0
	.long	1072693248
	.align 8
.LC5:
	.long	-1598689907
	.long	1051772663
	.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:

#ifndef CHEBYSHEV_CUH_
#define CHEBYSHEV_CUH_

double* chebyshev(int matrix_size, double** Ab, int s, int max_iter);

#endif
	.text
	.file	"chebyshev.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
