#include "Matrix.cuh"
#include <cstring>
#include <fstream>
#include <ctime>
#include <device_functions.h>

#ifdef __CUDACC__
#define cuda_SYNCTHREADS() __syncthreads()
#else
#define cuda_SYNCTHREADS()
#endif

#define Zero ZeroCPU
#define PRINT_LOG false
//#define TARGET_RESIDUE ((double)1.0e-9);

const double TARGET_RESIDUE = 1.0e-6;

Matrix::Matrix(int cols, int rows) : cols(cols), rows(rows)
{
	if (PRINT_LOG) printf("Matrix constructor\n");
	cudaMallocManaged(&mat, cols * rows * sizeof(double));
}

unsigned Matrix::getRows() const
{
	return rows;
}

unsigned Matrix::getCols() const
{
	return cols;
}

Matrix::Matrix(int cols, int rows, double* mat) : cols(cols), rows(rows), mat(mat)
{
	if (PRINT_LOG) printf("Matrix constructor\n");
	//cudaMallocManaged(&mat, cols * rows * sizeof(double));
}

Matrix::Matrix(const Matrix& a)
{
	if (PRINT_LOG) printf("Matrix copy constructor\n");
	rows = a.rows;
	cols = a.cols;
	cudaMallocManaged(&mat, cols * rows * sizeof(double));
	std::memcpy(mat, a.mat, cols * rows * sizeof(double));
}

void Matrix::operator=(const Matrix& a)
{
	if (PRINT_LOG) printf("Matrix assignment operator\n");
	rows = a.rows;
	cols = a.cols;
	cudaFree(mat);
	cudaMallocManaged(&mat, cols * rows * sizeof(double));
	std::memcpy(mat, a.mat, cols * rows * sizeof(double));
}

Matrix Matrix::Stub()
{
	return Matrix(1, 1);
}

Matrix Matrix::ZeroCPU(int cols, int rows)
{
	double* mat;
	cudaMallocManaged(&mat, cols * rows * sizeof(double));
	cudaDeviceSynchronize();
	for (long i = 0; i < cols * rows; i++)
	{
		mat[i] = 0.0f;
	}
	return Matrix(cols, rows, mat);
}

Matrix Matrix::OneCPU(int cols, int rows)
{
	double* mat;
	cudaMallocManaged(&mat, cols * rows * sizeof(double));
	for (long i = 0; i < cols * rows; i++)
	{
		mat[i] = 1.0f;
	}
	return Matrix(cols, rows, mat);
}

__global__ void ZeroGPUKernel(const int n, double* A)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;
	for (int i = index; i < n; i += stride)
	{
		A[index] = 0.0f;
	}
}

Matrix Matrix::ZeroGPU(int cols, int rows)
{
	double* mat;
	cudaMallocManaged(&mat, cols * rows * sizeof(double));
	int blockCount = (cols * rows + BLOCK_SIZE - 1) / BLOCK_SIZE;
	ZeroGPUKernel <<<blockCount, BLOCK_SIZE >>>(cols * rows, mat);
	cudaDeviceSynchronize();
	return Matrix(cols, rows, mat);
}

Matrix Matrix::IdentityCPU(int cols, int rows)
{
	if (cols != rows) throw "Identity matrix must be square";
	auto ret = Zero(cols, rows);
	for (int i = 0; i < cols; ++i)
	{
		ret.mat[i * cols + i] = 1.0f;
	}
	return ret;
}

Matrix Matrix::FromFile(std::string path)
{
	std::fstream reader;
	int cols, rows;
	reader.open(path, std::ios::in);
	reader.seekp(0);
	reader >> cols;
	reader >> rows;
	double* mat;
	cudaMallocManaged(&mat, cols * rows * sizeof(double));
	for (int i = 0; i < cols * rows; ++i)
	{
		reader >> mat[i];
	}
	reader.close();
	return Matrix(cols, rows, mat);
}

Matrix Matrix::Jacobi(const Matrix& A, const Matrix& b)
{
	auto LU = A;
	auto invD = (LU.separateDiagonal());
	auto x = ZeroCPU(1, A.getRows());
	invD.inverseDiagonalInPlaceCPU();
	auto M = -invD * LU;
	auto temp = invD * b;
	double res = 1;
	int counter = 0;

	do
	{
		x = (M * x + temp);
		//if (counter++ == 9)
		//{
		//	counter = 0;
		res = (A * x - b).vectorEuclideanNorm();
		//	printf("res: %f\n", res);
		//}
		counter++;
	}
	while (res > TARGET_RESIDUE);
	printf("res: %d \n", counter);

	return x;
}

Matrix Matrix::JacobiOptimal(const Matrix& A, const Matrix& b)
{
	// 25% czasu wykonania (80000us) prawdopodobnie kopiowanie pamieci z device na host i z powrotem
	//auto LU = A;
	//->
	auto LU = Matrix(A.cols, A.rows);
	copyGPU(LU, A);
	//32x wzrost wydajnosci

	//auto invD = (LU.separateDiagonal());
	//invD.inverseDiagonalInPlaceCPU();
	auto invD = Matrix(A.cols, A.rows);
	separateDiagonalAndInverseGPU(invD, LU);

	auto x = ZeroGPU(1, A.getRows());

	//auto temp1 = invD * b;
	auto temp1 = Matrix(1, A.rows);
	refMul(temp1, invD, b);

	//auto M = -invD * LU;
	//auto M = Matrix(A.cols, A.rows);
	auto M = Matrix(A.cols, A.rows);
	additiveInverseInPlaceGPU(invD);
	refMulDiag(M, invD, LU);


	double res = 100;
	int counter = 9;

	auto memmul = Matrix(1, A.rows);

	auto _Amulx = Matrix(1, A.rows);
	auto resVector = Matrix(1, A.rows);

	do
	{
		refMul(memmul, M, x);
		refAdd(x, memmul, temp1);
		//x = (M * x + temp);
		if (counter++ == 9)
		{
			counter = 0;
			refMul(_Amulx, A, x);
			refSub(resVector, _Amulx, b);
			res = resVector.vectorEuclideanNorm();
			//printf("res: %f\n", res);
		}
	}
	while (res > TARGET_RESIDUE);
	return x;
}

Matrix Matrix::ForwardSubstitution(const Matrix& A, const Matrix& b)
{
	if (!(A.cols == A.rows && A.rows == b.rows)) throw "Incorrect dimensions";
	auto x = Matrix(1, A.getRows());

	for (int i = 0; i < x.rows; ++i)
	{
		double sum = 0;
		for (int j = 0; j < i; ++j)
		{
			sum += A.mat[i * A.cols + j] * x.mat[j];
		}
		x.mat[i] = (b.mat[i] - sum) / A.mat[i * A.cols + i];
	}
	return x;
}

Matrix Matrix::BackwardSubstitution(const Matrix& A, const Matrix& b)
{
	if (!(A.cols == A.rows && A.rows == b.rows)) throw "Incorrect dimensions";
	auto x = Matrix(1, A.getRows());

	x.mat[0] = b.mat[0] / A.mat[0];

	for (int i = x.rows - 1; i >= 0; --i)
	{
		double sum = 0;
		for (int j = i + 1; j < A.cols; ++j)
		{
			sum += A.mat[i * A.cols + j] * x.mat[j];
		}
		x.mat[i] = (b.mat[i] - sum) / A.mat[i * A.cols + i];
	}
	return x;
}

Matrix Matrix::GaussSeidel(const Matrix& A, const Matrix& b)
{
	auto DL = -(A.lowerCPU() + A.diagonalCPU());
	auto U = A.upperCPU();
	auto x = ZeroCPU(1, A.getRows());
	auto temp = Matrix::ForwardSubstitution(DL, b);
	double res = 1;
	int counter = 0;

	do
	{
		//x = -(Matrix::ForwardSubstitution(DL, U * x)) + temp;
		x = (Matrix::ForwardSubstitution(DL, U * x)) + temp;
		//if (counter++ == 9)
		//{
		counter++;
		res = (A * (-x) - b).vectorEuclideanNorm();
		//}
		//printf("res: %f \n", res);
		//(x).print();
	}
	while (res > TARGET_RESIDUE);
	printf("res: %d \n", counter);
	return -x;
}

Matrix Matrix::GaussSeidelOptimal(const Matrix& A, const Matrix& b)
{
	//auto DL = (A.lowerCPU() + A.diagonalCPU());
	//auto U = A.upperCPU();
	auto DL = Matrix(A.cols, A.rows);
	auto U = Matrix(A.cols, A.rows);
	copyGPU(DL, A);
	separateUpperGPU(U, DL);

	//auto DL = (A.lowerCPU() + A.diagonalCPU());
	//auto U = A.upperCPU();

	auto x = ZeroCPU(1, A.getRows());
	auto temp = Matrix::ForwardSubstitution(DL, b);

	auto memmul = Matrix(1, A.rows);
	auto memforwardsub = Matrix(1, A.rows);


	auto memmulres = Matrix(1, A.rows);
	auto resVector = Matrix(1, A.rows);
	double res;
	int counter = 9;

	do
	{
		//x = -(Matrix::ForwardSubstitution(DL, U * x)) + temp;
		refMul(memmul, U, x);

		forwardSubstitutionGPU(memforwardsub, DL, memmul);
		//memforwardsub = Matrix::ForwardSubstitution(DL, memmul);

		//double xd = maxError(memforwardsub, memforwardsub2);

		additiveInverseInPlaceGPU(memforwardsub);
		refAdd(x, memforwardsub, temp);

		//x = memforwardsub + temp;
		if (counter++ == 9)
		{
			counter = 0;
			refMul(memmulres, A, x);
			refSub(resVector, memmulres, b);
			res = resVector.vectorEuclideanNorm();
		}
		//printf("res: %f \n", res);
		//(x).print();
	}
	while (res > TARGET_RESIDUE);
	return x;
}

Matrix Matrix::LUMehtod(const Matrix& A, const Matrix& b)
{
	Matrix L = Matrix::Stub();
	Matrix U = Matrix::Stub();

	Matrix::doolitle(L, U, A);

	auto y = Matrix::ForwardSubstitution(L, b);

	return Matrix::BackwardSubstitution(U, y);
}

Matrix Matrix::LUMehtodOptimal(const Matrix& A, const Matrix& b)
{
	Matrix L = Matrix::Stub();
	Matrix U = Matrix::Stub();

	Matrix::doolitle(L, U, A);

	auto y = Matrix::ForwardSubstitution(L, b);

	return Matrix::BackwardSubstitution(U, y);
}

void Matrix::doolitle(Matrix& L, Matrix& U, const Matrix& A)
{
	if (A.cols != A.rows) throw "Matrix is not square";
	L = OneCPU(A.cols, A.rows).diagonalCPU();
	U = ZeroCPU(A.cols, A.rows);
	for (int j = 0; j < A.cols; ++j)
	{
		for (int i = 0; i <= j; ++i)
		{
			double sum = 0;
			for (int k = 0; k < i; ++k)
			{
				sum += L.mat[i * L.cols + k] * U.mat[k * U.cols + j];
			}
			U.mat[i * U.cols + j] = A.mat[i * U.cols + j] - sum;
		}

		for (int i = j + 1; i < A.cols; ++i)
		{
			double sum = 0;
			for (int k = 0; k < j; ++k)
			{
				sum += L.mat[i * L.cols + k] * U.mat[k * U.cols + j];
			}
			L.mat[i * U.cols + j] = 1 / U.mat[j * U.cols + j] * (A.mat[i * U.cols + j] - sum);
		}
	}
}

__global__ void doolitleKernel(const int n, double* A, double* B)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		A[j] = B[j];
	}
}


void Matrix::doolitleGPU(Matrix& L, Matrix& U, const Matrix& A)
{
	int blockCount = (A.rows * A.cols + BLOCK_SIZE - 1) / BLOCK_SIZE;
	//doolitleKernel <<< blockCount, BLOCK_SIZE >>> (A.rows * A.cols, A.mat);
	cudaDeviceSynchronize();
}

void Matrix::createTest(Matrix& A, Matrix& b, Matrix& x, int size)
{
	srand(time(NULL));

	const int constrange = 100;
	const auto r = [](int range)-> double { return (double)(rand() % 20000) / 100 - 100; };
	x = Matrix(1, size);
	A = Matrix(size, size);
	b = Matrix(1, size);
	for (int i = 0; i < size; ++i)
	{
		x.mat[i] = r(100);
	}

	for (int i = 0; i < size; ++i)
	{
		double sum = 0;
		for (int j = 0; j < size; ++j)
		{
			if (i != j)
			{
				A.mat[i * size + j] = r(100);
				sum += fabs(A.mat[i * size + j]);
			}
			double randomized = r(100);
			if (randomized > 0)
			{
				A.mat[i * size + i] = sum + r(10);
			}
			else
			{
				A.mat[i * size + i] = -sum + r(10);
			}
		}
	}

	for (int i = 0; i < size; ++i)
	{
		double sum = 0;
		for (int j = 0; j < size; ++j)
		{
			sum += A.mat[i * size + j] * x.mat[j];
		}
		b.mat[i] = sum;
	}
}

void Matrix::createTask(Matrix& A, Matrix& b, const int size)
{
	//const int size = 994;
	const int a1 = 5 + 7;
	const int a2 = -1;
	const int a3 = a2;
	const int inSin(1 + 1);
	A = Matrix::ZeroCPU(size, size);
	b = Matrix(1, size);

	for (int i = 0; i < size; ++i)
	{
		A.mat[size * i + i] = a1;
		if (size * i + i - 1 >= 0)
			A.mat[size * i + i - 1] = a2;
		if (size * i + i - 2 >= 0)
			A.mat[size * i + i - 2] = a3;
		if (size * i + i + 1 < size * size)
			A.mat[size * i + i + 1] = a2;
		if (size * i + i + 2 < size * size)
			A.mat[size * i + i + 2] = a3;
	}

	for (int i = 0; i < size; ++i)
	{
		b.mat[i] = sin(i * inSin);
	}
}

void Matrix::createTaskC(Matrix& A, Matrix& b)
{
	const int size = 994;
	const int a1 = 3;
	const int a2 = -1;
	const int a3 = a2;
	const int inSin(1 + 1);
	A = Matrix::ZeroCPU(size, size);
	b = Matrix(1, size);

	for (int i = 0; i < size; ++i)
	{
		A.mat[size * i + i] = a1;
		if (size * i + i - 1 >= 0)
			A.mat[size * i + i - 1] = a2;
		if (size * i + i - 2 >= 0)
			A.mat[size * i + i - 2] = a3;
		if (size * i + i + 1 < size * size)
			A.mat[size * i + i + 1] = a2;
		if (size * i + i + 2 < size * size)
			A.mat[size * i + i + 2] = a3;
	}

	for (int i = 0; i < size; ++i)
	{
		b.mat[i] = sin(i * inSin);
	}
}

double Matrix::maxError(Matrix& x, Matrix& r)
{
	if (x.rows * x.cols != r.rows * r.cols) throw "Matrices are not the same size";
	double max = 0;
	for (int i = 0; i < x.rows * x.cols; ++i)
	{
		if (fabs(x.mat[i] - r.mat[i]) > max)
			max = fabs(x.mat[i] - r.mat[i]);
	}
	return max;
}

__global__ void copyKernel(const int n, double* A, double* B)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		A[j] = B[j];
	}
}

void Matrix::copyGPU(Matrix& a, const Matrix& b)
{
	int blockCount = (a.cols * a.rows + BLOCK_SIZE - 1) / BLOCK_SIZE;
	copyKernel <<< blockCount, BLOCK_SIZE >>>(a.cols * a.rows, a.mat, b.mat);
	cudaDeviceSynchronize();
}

__global__ void separateDiagonalKernel(const int n, double* d, double* A)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		d[j * n + j] = 1 / A[j * n + j];
		A[j * n + j] = 0;
	}
}

void Matrix::separateDiagonalAndInverseGPU(Matrix& d, Matrix& A)
{
	int blockCount = (A.cols + BLOCK_SIZE - 1) / BLOCK_SIZE;
	separateDiagonalKernel <<< blockCount, BLOCK_SIZE >>>(A.cols, d.mat, A.mat);
	cudaDeviceSynchronize();
}

__global__ void separateUpperKernel(const int n, const int cols, double* U, double* A)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		int row = j / cols;
		int col = j % cols;
		if (col > row)
		{
			U[j] = A[j];
			A[j] = 0;
		}
	}
}

void Matrix::separateUpperGPU(Matrix& U, Matrix& A)
{
	int blockCount = (A.cols + BLOCK_SIZE - 1) / BLOCK_SIZE;
	separateUpperKernel <<< blockCount, BLOCK_SIZE >>>(A.cols * A.rows, A.cols, U.mat, A.mat);
	cudaDeviceSynchronize();
}

__global__ void additiveInverseInPlaceKernel(const int n, double* A)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		A[j] = -A[j];
	}
}

void Matrix::additiveInverseInPlaceGPU(Matrix& A)
{
	int blockCount = (A.rows * A.cols + BLOCK_SIZE - 1) / BLOCK_SIZE;
	additiveInverseInPlaceKernel <<< blockCount, BLOCK_SIZE >>>(A.rows * A.cols, A.mat);
	cudaDeviceSynchronize();
}

__global__ void forwardSubstitutionKernel(const int n, double* A, double* b, double* x)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		double sum = 0;
		for (int i = 0; i < n; i++)
		{
			if (i == j)
			{
				x[j] = (b[j] - sum) / A[j * n + j];
			}
			cuda_SYNCTHREADS();
			if (i < j)
			{
				sum += A[j * n + i] * x[i];
			}
		}
	}
}

void Matrix::forwardSubstitutionGPU(Matrix& result, const Matrix& A, const Matrix& b)
{
	int blockCount = 1;
	int blockSize = pow(2, ceil(log2f(A.cols)));
	forwardSubstitutionKernel <<< blockCount, blockSize >>>(A.cols, A.mat, b.mat, result.mat);
	cudaDeviceSynchronize();
}

void Matrix::backwardSubstitutionGPU(Matrix& result, const Matrix& A, const Matrix& b)
{
}

void Matrix::toFile(std::string path)
{
	std::fstream writer;
	writer.open(path, std::ios::out);
	writer.seekg(0);
	writer << cols << ' ' << rows << '\n';
	for (int i = 0; i < rows; ++i)
	{
		for (int j = 0; j < cols; ++j)
		{
			writer << mat[i * cols + j] << ' ';
		}
		writer << "\n";
	}
	writer.close();
}

Matrix Matrix::separateDiagonal()
{
	if (cols != rows) throw "Matrix is not square";
	auto ret = Zero(cols, rows);
	for (int i = 0; i < cols; ++i)
	{
		ret.mat[i * cols + i] = mat[i * cols + i];
		mat[i * cols + i] = 0.0f;
	}
	return ret;
}

Matrix Matrix::diagonalCPU() const
{
	if (cols != rows) throw "Matrix is not square";
	auto ret = Zero(cols, rows);
	for (int i = 0; i < cols; ++i)
	{
		ret.mat[i * cols + i] = mat[i * cols + i];
	}
	return ret;
}

Matrix Matrix::lowerCPU() const
{
	if (cols != rows) throw "Matrix is not square";
	auto ret = Zero(cols, rows);
	for (int j = 0; j < cols; ++j)
	{
		for (int i = 0; i < j; ++i)
		{
			ret.mat[j * cols + i] = mat[j * cols + i];
		}
	}
	return ret;
}

Matrix Matrix::upperCPU() const
{
	if (cols != rows) throw "Matrix is not square";
	auto ret = Zero(cols, rows);
	for (int j = 0; j < cols; ++j)
	{
		for (int i = j + 1; i < cols; ++i)
		{
			ret.mat[j * cols + i] = mat[j * cols + i];
		}
	}
	return ret;
}

void Matrix::inverseDiagonalInPlaceCPU()
{
	if (cols != rows) throw "Matrix is not square";
	for (int i = 0; i < cols; ++i)
	{
		if (mat[i * cols + i] == 0) throw "0 on diagonal";
		mat[i * cols + i] = 1 / mat[i * cols + i];
	}
}

void Matrix::transposeVectorInPlace()
{
	unsigned int tmp = cols;
	cols = rows;
	rows = tmp;
}

double Matrix::vectorEuclideanNorm()
{
	if (cols != 1 && rows != 1) throw "Matrix is not a vector";
	double sum = 0;
	for (int i = 0; i < cols * rows; ++i)
	{
		sum += mat[i] * mat[i];
	}
	return sqrt(sum);
}

Matrix Matrix::lu()
{
	throw "Not implemented";
}

void Matrix::print() const
{
	for (int i = 0; i < rows; ++i)
	{
		for (int j = 0; j < cols; ++j)
		{
			printf("%f ", mat[i * cols + j]);
		}
		printf("\n");
	}
	printf("\n");
}

Matrix::~Matrix()
{
	if (PRINT_LOG) printf("Matrix destructor\n");
	cudaFree(mat);
	//free(mat);
}

__global__ void mulKernel(const int commonDim, const int cols, const int n, double* A, double* B, double* C)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		int row = j / cols;
		int col = j % cols;
		C[j] = 0;
		for (int i = 0; i < commonDim; i++)
		{
			C[j] += A[row * commonDim + i] * B[i * cols + col];
		}
	}
}

void Matrix::refMul(Matrix& result, const Matrix& a, const Matrix& b)
{
	int blockCount = (a.rows * b.cols + BLOCK_SIZE - 1) / BLOCK_SIZE;
	mulKernel <<< blockCount, BLOCK_SIZE >>>(a.cols, b.cols, b.cols * a.rows, a.mat, b.mat, result.mat);
	cudaDeviceSynchronize();
}

__global__ void mulDiagKernel(const int commonDim, const int cols, const int n, double* A, double* B, double* C)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		int row = j / cols;
		int col = j % cols;
		C[j] = A[row * commonDim + row] * B[row * commonDim + col];
	}
}

void Matrix::refMulDiag(Matrix& result, const Matrix& a, const Matrix& b)
{
	int blockCount = (a.rows * b.cols + BLOCK_SIZE - 1) / BLOCK_SIZE;
	mulDiagKernel << < blockCount, BLOCK_SIZE >> >(a.cols, b.cols, b.cols * a.rows, a.mat, b.mat, result.mat);
	cudaDeviceSynchronize();
}

Matrix operator*(const Matrix& a, const Matrix& b)
{
	if (a.cols != b.rows) throw "wrong dimensions for multiplication";
	double* mat;
	cudaMallocManaged(&mat, b.cols * a.rows * sizeof(double));
	int blockCount = (a.rows * b.cols + BLOCK_SIZE - 1) / BLOCK_SIZE;
	if (PRINT_LOG) printf("Matrix multiplication on %d blocks x %d threads\n", blockCount, BLOCK_SIZE);
	mulKernel <<< blockCount, BLOCK_SIZE >>>(a.cols, b.cols, b.cols * a.rows, a.mat, b.mat, mat);
	cudaDeviceSynchronize();
	return Matrix(b.cols, a.rows, mat);
}

__global__ void addKernel(const int n, double* A, double* B, double* C)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		C[j] = A[j] + B[j];
	}
}

void Matrix::refAdd(Matrix& result, const Matrix& a, const Matrix& b)
{
	int blockCount = (a.cols * a.rows + BLOCK_SIZE - 1) / BLOCK_SIZE;
	addKernel <<< blockCount, BLOCK_SIZE >>>(a.cols * a.rows, a.mat, b.mat, result.mat);
	cudaDeviceSynchronize();
}

Matrix operator+(const Matrix& a, const Matrix& b)
{
	if (a.cols != b.cols || a.rows != b.rows) throw "dimensions must equal for addition";
	double* mat;
	cudaMallocManaged(&mat, a.cols * a.rows * sizeof(double));
	int blockCount = (a.cols * a.rows + BLOCK_SIZE - 1) / BLOCK_SIZE;
	if (PRINT_LOG) printf("Matrix addition on %d blocks x %d threads\n", blockCount, BLOCK_SIZE);
	addKernel <<< blockCount, BLOCK_SIZE >>>(a.cols * a.rows, a.mat, b.mat, mat);
	cudaDeviceSynchronize();
	return Matrix(a.cols, a.rows, mat);
}

__global__ void subKernel(const int n, double* A, double* B, double* C)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		C[j] = A[j] - B[j];
	}
}

void Matrix::refSub(Matrix& result, const Matrix& a, const Matrix& b)
{
	int blockCount = (a.cols * a.rows + BLOCK_SIZE - 1) / BLOCK_SIZE;
	subKernel <<< blockCount, BLOCK_SIZE >> >(a.cols * a.rows, a.mat, b.mat, result.mat);
	cudaDeviceSynchronize();
}

Matrix operator-(const Matrix& a, const Matrix& b)
{
	if (a.cols != b.cols || a.rows != b.rows) throw "dimensions must equal for addition";
	double* mat;
	cudaMallocManaged(&mat, a.cols * a.rows * sizeof(double));
	int blockCount = (a.cols * a.rows + BLOCK_SIZE - 1) / BLOCK_SIZE;
	if (PRINT_LOG) printf("Matrix addition on %d blocks x %d threads\n", blockCount, BLOCK_SIZE);
	subKernel <<< blockCount, BLOCK_SIZE >>>(a.cols * a.rows, a.mat, b.mat, mat);
	cudaDeviceSynchronize();
	return Matrix(a.cols, a.rows, mat);
}

__global__ void additiveInverseKernel(const int n, double* A, double* B)
{
	int index = blockIdx.x * blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for (int j = index; j < n; j += stride)
	{
		A[j] = -B[j];
	}
}

Matrix operator-(const Matrix& a)
{
	double* mat;
	cudaMallocManaged(&mat, a.cols * a.rows * sizeof(double));
	int blockCount = (a.cols * a.rows + BLOCK_SIZE - 1) / BLOCK_SIZE;
	additiveInverseKernel <<<blockCount, BLOCK_SIZE >>>(a.cols * a.rows, mat, a.mat);
	cudaDeviceSynchronize();
	return Matrix(a.cols, a.rows, mat);
}

	.file	"tmpxft_002e372c_00000000-6_Matrix.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3830:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3830:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0, @function
_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0:
.LFB4581:
	.cfi_startproc
	movl	$1, %edx
	jmp	cudaMallocManaged@PLT
	.cfi_endproc
.LFE4581:
	.size	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0, .-_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	.type	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0, @function
_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0:
.LFB4582:
	.cfi_startproc
	jmp	cudaLaunchKernel@PLT
	.cfi_endproc
.LFE4582:
	.size	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0, .-_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	.align 2
	.type	_ZNSt13basic_fstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0, @function
_ZNSt13basic_fstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0:
.LFB4583:
	.cfi_startproc
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	addq	$24, %rdi
	call	_ZNSt13basic_filebufIcSt11char_traitsIcEE4openEPKcSt13_Ios_Openmode@PLT
	testq	%rax, %rax
	movq	(%rbx), %rax
	jne	.L5
	addq	-24(%rax), %rbx
	movl	$4, %esi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_remember_state
	.cfi_def_cfa_offset 8
	jmp	_ZNSt9basic_iosIcSt11char_traitsIcEE8setstateESt12_Ios_Iostate@PLT
.L5:
	.cfi_restore_state
	addq	-24(%rax), %rbx
	xorl	%esi, %esi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	_ZNSt9basic_iosIcSt11char_traitsIcEE5clearESt12_Ios_Iostate@PLT
	.cfi_endproc
.LFE4583:
	.size	_ZNSt13basic_fstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0, .-_ZNSt13basic_fstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0
	.align 2
	.type	_ZZN6Matrix10createTestERS_S0_S0_iENKUliE_clEi.isra.0, @function
_ZZN6Matrix10createTestERS_S0_S0_iENKUliE_clEi.isra.0:
.LFB4585:
	.cfi_startproc
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	rand@PLT
	movl	$20000, %ecx
	movsd	.LC0(%rip), %xmm1
	cltd
	idivl	%ecx
	cvtsi2sdl	%edx, %xmm0
	popq	%rdx
	.cfi_def_cfa_offset 8
	divsd	%xmm1, %xmm0
	subsd	%xmm1, %xmm0
	ret
	.cfi_endproc
.LFE4585:
	.size	_ZZN6Matrix10createTestERS_S0_S0_iENKUliE_clEi.isra.0, .-_ZZN6Matrix10createTestERS_S0_S0_iENKUliE_clEi.isra.0
	.align 2
	.globl	_ZN6MatrixC2Eii
	.type	_ZN6MatrixC2Eii, @function
_ZN6MatrixC2Eii:
.LFB3767:
	.cfi_startproc
	endbr64
	movl	%edx, (%rdi)
	imull	%esi, %edx
	addq	$8, %rdi
	movl	%esi, -4(%rdi)
	movslq	%edx, %rsi
	salq	$3, %rsi
	jmp	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	.cfi_endproc
.LFE3767:
	.size	_ZN6MatrixC2Eii, .-_ZN6MatrixC2Eii
	.globl	_ZN6MatrixC1Eii
	.set	_ZN6MatrixC1Eii,_ZN6MatrixC2Eii
	.align 2
	.globl	_ZNK6Matrix7getRowsEv
	.type	_ZNK6Matrix7getRowsEv, @function
_ZNK6Matrix7getRowsEv:
.LFB3769:
	.cfi_startproc
	endbr64
	movl	(%rdi), %eax
	ret
	.cfi_endproc
.LFE3769:
	.size	_ZNK6Matrix7getRowsEv, .-_ZNK6Matrix7getRowsEv
	.align 2
	.globl	_ZNK6Matrix7getColsEv
	.type	_ZNK6Matrix7getColsEv, @function
_ZNK6Matrix7getColsEv:
.LFB3770:
	.cfi_startproc
	endbr64
	movl	4(%rdi), %eax
	ret
	.cfi_endproc
.LFE3770:
	.size	_ZNK6Matrix7getColsEv, .-_ZNK6Matrix7getColsEv
	.align 2
	.globl	_ZN6MatrixC2EiiPd
	.type	_ZN6MatrixC2EiiPd, @function
_ZN6MatrixC2EiiPd:
.LFB3772:
	.cfi_startproc
	endbr64
	movl	%edx, (%rdi)
	movl	%esi, 4(%rdi)
	movq	%rcx, 8(%rdi)
	ret
	.cfi_endproc
.LFE3772:
	.size	_ZN6MatrixC2EiiPd, .-_ZN6MatrixC2EiiPd
	.globl	_ZN6MatrixC1EiiPd
	.set	_ZN6MatrixC1EiiPd,_ZN6MatrixC2EiiPd
	.align 2
	.globl	_ZN6MatrixC2ERKS_
	.type	_ZN6MatrixC2ERKS_, @function
_ZN6MatrixC2ERKS_:
.LFB3775:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	leaq	8(%rdi), %rdi
	pushq	%rax
	.cfi_def_cfa_offset 32
	movl	4(%rbp), %eax
	movl	(%rsi), %esi
	movl	%eax, -4(%rdi)
	movl	%esi, -8(%rdi)
	imull	%eax, %esi
	salq	$3, %rsi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movl	4(%rbx), %ecx
	imull	(%rbx), %ecx
	movq	8(%rbx), %rax
	movq	8(%rbp), %rsi
	salq	$3, %rcx
	movq	%rax, %rdi
	rep movsb
	popq	%rdx
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3775:
	.size	_ZN6MatrixC2ERKS_, .-_ZN6MatrixC2ERKS_
	.globl	_ZN6MatrixC1ERKS_
	.set	_ZN6MatrixC1ERKS_,_ZN6MatrixC2ERKS_
	.align 2
	.globl	_ZN6MatrixaSERKS_
	.type	_ZN6MatrixaSERKS_, @function
_ZN6MatrixaSERKS_:
.LFB3777:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	pushq	%rax
	.cfi_def_cfa_offset 32
	movl	(%rsi), %eax
	movl	%eax, (%rdi)
	movl	4(%rsi), %eax
	movl	%eax, 4(%rdi)
	movq	8(%rdi), %rdi
	call	cudaFree@PLT
	movl	4(%rbx), %esi
	imull	(%rbx), %esi
	leaq	8(%rbx), %rdi
	salq	$3, %rsi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movl	4(%rbx), %ecx
	imull	(%rbx), %ecx
	movq	8(%rbx), %rax
	movq	8(%rbp), %rsi
	salq	$3, %rcx
	movq	%rax, %rdi
	rep movsb
	popq	%rdx
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3777:
	.size	_ZN6MatrixaSERKS_, .-_ZN6MatrixaSERKS_
	.align 2
	.globl	_ZN6Matrix4StubEv
	.type	_ZN6Matrix4StubEv, @function
_ZN6Matrix4StubEv:
.LFB3778:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	$1, %edx
	movq	%rdi, %rbx
	movl	$1, %esi
	call	_ZN6MatrixC1Eii
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3778:
	.size	_ZN6Matrix4StubEv, .-_ZN6Matrix4StubEv
	.align 2
	.globl	_ZN6Matrix7ZeroCPUEii
	.type	_ZN6Matrix7ZeroCPUEii, @function
_ZN6Matrix7ZeroCPUEii:
.LFB3779:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	%esi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	imull	%edx, %r13d
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	movslq	%r13d, %r13
	leaq	0(,%r13,8), %rsi
	subq	$24, %rsp
	.cfi_def_cfa_offset 64
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	call	cudaDeviceSynchronize@PLT
	movq	(%rsp), %rdx
	xorl	%eax, %eax
.L21:
	cmpq	%rax, %r13
	jle	.L25
	movq	$0x000000000, (%rdx,%rax,8)
	incq	%rax
	jmp	.L21
.L25:
	movl	%r12d, (%rbx)
	movl	%ebp, 4(%rbx)
	movq	%rdx, 8(%rbx)
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L23
	call	__stack_chk_fail@PLT
.L23:
	addq	$24, %rsp
	.cfi_def_cfa_offset 40
	movq	%rbx, %rax
	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
.LFE3779:
	.size	_ZN6Matrix7ZeroCPUEii, .-_ZN6Matrix7ZeroCPUEii
	.align 2
	.globl	_ZN6Matrix6OneCPUEii
	.type	_ZN6Matrix6OneCPUEii, @function
_ZN6Matrix6OneCPUEii:
.LFB3780:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	%esi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	imull	%edx, %r13d
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	movslq	%r13d, %r13
	leaq	0(,%r13,8), %rsi
	subq	$24, %rsp
	.cfi_def_cfa_offset 64
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	(%rsp), %rdx
	movsd	.LC2(%rip), %xmm0
	xorl	%eax, %eax
.L27:
	cmpq	%rax, %r13
	jle	.L31
	movsd	%xmm0, (%rdx,%rax,8)
	incq	%rax
	jmp	.L27
.L31:
	movl	%r12d, (%rbx)
	movl	%ebp, 4(%rbx)
	movq	%rdx, 8(%rbx)
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	addq	$24, %rsp
	.cfi_def_cfa_offset 40
	movq	%rbx, %rax
	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
.LFE3780:
	.size	_ZN6Matrix6OneCPUEii, .-_ZN6Matrix6OneCPUEii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"Identity matrix must be square"
	.text
	.align 2
	.globl	_ZN6Matrix11IdentityCPUEii
	.type	_ZN6Matrix11IdentityCPUEii, @function
_ZN6Matrix11IdentityCPUEii:
.LFB3782:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	cmpl	%edx, %esi
	je	.L33
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC3(%rip), %rax
	movq	%rax, (%rdi)
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L39
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	call	__cxa_throw@PLT
.L33:
	movl	%esi, %ebx
	movl	%esi, %edx
	movq	%rdi, %rbp
	call	_ZN6Matrix7ZeroCPUEii
	movslq	%ebx, %rax
	movsd	.LC2(%rip), %xmm0
	xorl	%edx, %edx
	leaq	8(,%rax,8), %rcx
	xorl	%eax, %eax
.L35:
	cmpl	%ebx, %eax
	jge	.L40
	movq	8(%rbp), %rsi
	incl	%eax
	movsd	%xmm0, (%rsi,%rdx)
	addq	%rcx, %rdx
	jmp	.L35
.L40:
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L37
.L39:
	call	__stack_chk_fail@PLT
.L37:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3782:
	.size	_ZN6Matrix11IdentityCPUEii, .-_ZN6Matrix11IdentityCPUEii
	.section	.rodata.str1.1
.LC4:
	.string	"Incorrect dimensions"
	.text
	.align 2
	.globl	_ZN6Matrix19ForwardSubstitutionERKS_S1_
	.type	_ZN6Matrix19ForwardSubstitutionERKS_S1_, @function
_ZN6Matrix19ForwardSubstitutionERKS_S1_:
.LFB3786:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movl	(%rsi), %edx
	cmpl	%edx, 4(%rsi)
	jne	.L42
	cmpl	(%r12), %edx
	je	.L43
.L42:
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	movq	%rax, %rdi
	leaq	.LC4(%rip), %rax
	movq	%rax, (%rdi)
	call	__cxa_throw@PLT
.L43:
	movq	%rdi, %rbx
	movq	%rsi, %rbp
	movl	$1, %esi
	call	_ZN6MatrixC1Eii
	movl	(%rbx), %r8d
	xorl	%eax, %eax
.L44:
	movl	%eax, %r10d
	cmpl	%r8d, %eax
	jnb	.L41
	movl	4(%rbp), %edx
	movq	8(%rbp), %rdi
	xorl	%ecx, %ecx
	xorps	%xmm1, %xmm1
	movq	8(%rbx), %rsi
	movl	%edx, %r9d
	imull	%eax, %r9d
.L47:
	cmpl	%eax, %ecx
	jge	.L49
	leal	(%r9,%rcx), %r11d
	movsd	(%rdi,%r11,8), %xmm0
	mulsd	(%rsi,%rcx,8), %xmm0
	incq	%rcx
	addsd	%xmm0, %xmm1
	jmp	.L47
.L49:
	movq	8(%r12), %rcx
	incl	%edx
	imull	%r10d, %edx
	movsd	(%rcx,%rax,8), %xmm0
	subsd	%xmm1, %xmm0
	divsd	(%rdi,%rdx,8), %xmm0
	movsd	%xmm0, (%rsi,%rax,8)
	incq	%rax
	jmp	.L44
.L41:
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3786:
	.size	_ZN6Matrix19ForwardSubstitutionERKS_S1_, .-_ZN6Matrix19ForwardSubstitutionERKS_S1_
	.align 2
	.globl	_ZN6Matrix20BackwardSubstitutionERKS_S1_
	.type	_ZN6Matrix20BackwardSubstitutionERKS_S1_, @function
_ZN6Matrix20BackwardSubstitutionERKS_S1_:
.LFB3787:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movl	(%rsi), %edx
	cmpl	%edx, 4(%rsi)
	jne	.L51
	cmpl	(%r12), %edx
	je	.L52
.L51:
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	movq	%rax, %rdi
	leaq	.LC4(%rip), %rax
	movq	%rax, (%rdi)
	call	__cxa_throw@PLT
.L52:
	movq	%rdi, %rbx
	movq	%rsi, %rbp
	movl	$1, %esi
	call	_ZN6MatrixC1Eii
	movl	(%rbx), %eax
	movq	8(%r12), %rdi
	movq	8(%rbp), %rcx
	movq	8(%rbx), %rsi
	decl	%eax
	movsd	(%rdi), %xmm0
	cltq
	divsd	(%rcx), %xmm0
	movsd	%xmm0, (%rsi)
.L53:
	testl	%eax, %eax
	js	.L58
	movl	4(%rbp), %r9d
	movl	%eax, %r10d
	leaq	1(%rax), %rdx
	xorps	%xmm1, %xmm1
	movl	%r9d, %r8d
	imull	%eax, %r8d
.L54:
	cmpl	%r9d, %edx
	jnb	.L59
	leal	(%r8,%rdx), %r11d
	movsd	(%rcx,%r11,8), %xmm0
	mulsd	(%rsi,%rdx,8), %xmm0
	incq	%rdx
	addsd	%xmm0, %xmm1
	jmp	.L54
.L59:
	movsd	(%rdi,%rax,8), %xmm0
	leal	(%r8,%r10), %edx
	subsd	%xmm1, %xmm0
	divsd	(%rcx,%rdx,8), %xmm0
	movsd	%xmm0, (%rsi,%rax,8)
	decq	%rax
	jmp	.L53
.L58:
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3787:
	.size	_ZN6Matrix20BackwardSubstitutionERKS_S1_, .-_ZN6Matrix20BackwardSubstitutionERKS_S1_
	.align 2
	.globl	_ZN6Matrix11doolitleGPUERS_S0_RKS_
	.type	_ZN6Matrix11doolitleGPUERS_S0_RKS_, @function
_ZN6Matrix11doolitleGPUERS_S0_RKS_:
.LFB3793:
	.cfi_startproc
	endbr64
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3793:
	.size	_ZN6Matrix11doolitleGPUERS_S0_RKS_, .-_ZN6Matrix11doolitleGPUERS_S0_RKS_
	.section	.rodata.str1.1
.LC5:
	.string	"Matrices are not the same size"
	.text
	.align 2
	.globl	_ZN6Matrix8maxErrorERS_S0_
	.type	_ZN6Matrix8maxErrorERS_S0_, @function
_ZN6Matrix8maxErrorERS_S0_:
.LFB3800:
	.cfi_startproc
	endbr64
	movl	(%rdi), %edx
	movl	(%rsi), %ecx
	imull	4(%rdi), %edx
	imull	4(%rsi), %ecx
	cmpl	%ecx, %edx
	jne	.L67
	movq	.LC6(%rip), %xmm2
	xorl	%eax, %eax
	xorps	%xmm1, %xmm1
	jmp	.L62
.L67:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	movq	%rax, %rdi
	leaq	.LC5(%rip), %rax
	movq	%rax, (%rdi)
	call	__cxa_throw@PLT
.L62:
	.cfi_def_cfa_offset 8
	cmpl	%edx, %eax
	jnb	.L70
	movq	8(%rdi), %r8
	movq	8(%rsi), %rcx
	movsd	(%r8,%rax,8), %xmm0
	subsd	(%rcx,%rax,8), %xmm0
	incq	%rax
	andps	%xmm2, %xmm0
	maxsd	%xmm1, %xmm0
	movaps	%xmm0, %xmm1
	jmp	.L62
.L70:
	movaps	%xmm1, %xmm0
	ret
	.cfi_endproc
.LFE3800:
	.size	_ZN6Matrix8maxErrorERS_S0_, .-_ZN6Matrix8maxErrorERS_S0_
	.align 2
	.globl	_ZN6Matrix23backwardSubstitutionGPUERS_RKS_S2_
	.type	_ZN6Matrix23backwardSubstitutionGPUERS_RKS_S2_, @function
_ZN6Matrix23backwardSubstitutionGPUERS_RKS_S2_:
.LFB3806:
	.cfi_startproc
	endbr64
	ret
	.cfi_endproc
.LFE3806:
	.size	_ZN6Matrix23backwardSubstitutionGPUERS_RKS_S2_, .-_ZN6Matrix23backwardSubstitutionGPUERS_RKS_S2_
	.section	.rodata.str1.1
.LC7:
	.string	"Matrix is not square"
	.text
	.align 2
	.globl	_ZN6Matrix16separateDiagonalEv
	.type	_ZN6Matrix16separateDiagonalEv, @function
_ZN6Matrix16separateDiagonalEv:
.LFB3808:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movl	4(%rsi), %esi
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	cmpl	(%rbx), %esi
	je	.L73
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC7(%rip), %rax
	movq	%rax, (%rdi)
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L79
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	call	__cxa_throw@PLT
.L73:
	movl	%esi, %edx
	movq	%rdi, %rbp
	call	_ZN6Matrix7ZeroCPUEii
	movl	4(%rbx), %ecx
	xorl	%edx, %edx
	xorl	%eax, %eax
	leal	1(%rcx), %r8d
.L75:
	cmpl	%ecx, %eax
	je	.L80
	movq	8(%rbx), %rsi
	movl	%edx, %edi
	movq	8(%rbp), %r9
	incl	%eax
	addl	%r8d, %edx
	leaq	(%rsi,%rdi,8), %rsi
	movsd	(%rsi), %xmm0
	movsd	%xmm0, (%r9,%rdi,8)
	movq	$0x000000000, (%rsi)
	jmp	.L75
.L80:
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L77
.L79:
	call	__stack_chk_fail@PLT
.L77:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3808:
	.size	_ZN6Matrix16separateDiagonalEv, .-_ZN6Matrix16separateDiagonalEv
	.align 2
	.globl	_ZNK6Matrix11diagonalCPUEv
	.type	_ZNK6Matrix11diagonalCPUEv, @function
_ZNK6Matrix11diagonalCPUEv:
.LFB3809:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movl	4(%rsi), %esi
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	cmpl	(%rbx), %esi
	je	.L82
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC7(%rip), %rax
	movq	%rax, (%rdi)
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L88
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	call	__cxa_throw@PLT
.L82:
	movl	%esi, %edx
	movq	%rdi, %rbp
	call	_ZN6Matrix7ZeroCPUEii
	movl	4(%rbx), %ecx
	xorl	%edx, %edx
	xorl	%eax, %eax
	leal	1(%rcx), %edi
.L84:
	cmpl	%ecx, %eax
	je	.L89
	movq	8(%rbx), %r8
	movl	%edx, %esi
	incl	%eax
	addl	%edi, %edx
	movsd	(%r8,%rsi,8), %xmm0
	movq	8(%rbp), %r8
	movsd	%xmm0, (%r8,%rsi,8)
	jmp	.L84
.L89:
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L86
.L88:
	call	__stack_chk_fail@PLT
.L86:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3809:
	.size	_ZNK6Matrix11diagonalCPUEv, .-_ZNK6Matrix11diagonalCPUEv
	.align 2
	.globl	_ZNK6Matrix8lowerCPUEv
	.type	_ZNK6Matrix8lowerCPUEv, @function
_ZNK6Matrix8lowerCPUEv:
.LFB3810:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movl	4(%rsi), %esi
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	cmpl	(%rbx), %esi
	je	.L91
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC7(%rip), %rax
	movq	%rax, (%rdi)
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L99
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	call	__cxa_throw@PLT
.L91:
	movl	%esi, %edx
	movq	%rdi, %rbp
	call	_ZN6Matrix7ZeroCPUEii
	movl	4(%rbx), %esi
	xorl	%ecx, %ecx
	xorl	%eax, %eax
.L93:
	cmpl	%esi, %eax
	je	.L100
	xorl	%edx, %edx
.L95:
	cmpl	%edx, %eax
	je	.L101
	movq	8(%rbx), %r8
	leal	(%rdx,%rcx), %edi
	incl	%edx
	movsd	(%r8,%rdi,8), %xmm0
	movq	8(%rbp), %r8
	movsd	%xmm0, (%r8,%rdi,8)
	jmp	.L95
.L101:
	incl	%eax
	addl	%esi, %ecx
	jmp	.L93
.L100:
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L96
.L99:
	call	__stack_chk_fail@PLT
.L96:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3810:
	.size	_ZNK6Matrix8lowerCPUEv, .-_ZNK6Matrix8lowerCPUEv
	.align 2
	.globl	_ZNK6Matrix8upperCPUEv
	.type	_ZNK6Matrix8upperCPUEv, @function
_ZNK6Matrix8upperCPUEv:
.LFB3811:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movl	4(%rsi), %esi
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	cmpl	(%rbx), %esi
	je	.L103
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC7(%rip), %rax
	movq	%rax, (%rdi)
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L111
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	call	__cxa_throw@PLT
.L103:
	movl	%esi, %edx
	movq	%rdi, %rbp
	call	_ZN6Matrix7ZeroCPUEii
	movl	4(%rbx), %ecx
	xorl	%esi, %esi
	xorl	%eax, %eax
.L105:
	cmpl	%ecx, %eax
	je	.L112
	incl	%eax
	movl	%eax, %edx
.L106:
	cmpl	%ecx, %edx
	je	.L113
	movq	8(%rbx), %r8
	leal	(%rdx,%rsi), %edi
	incl	%edx
	movsd	(%r8,%rdi,8), %xmm0
	movq	8(%rbp), %r8
	movsd	%xmm0, (%r8,%rdi,8)
	jmp	.L106
.L113:
	addl	%ecx, %esi
	jmp	.L105
.L112:
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L109
.L111:
	call	__stack_chk_fail@PLT
.L109:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3811:
	.size	_ZNK6Matrix8upperCPUEv, .-_ZNK6Matrix8upperCPUEv
	.section	.rodata.str1.1
.LC8:
	.string	"0 on diagonal"
	.text
	.align 2
	.globl	_ZN6Matrix25inverseDiagonalInPlaceCPUEv
	.type	_ZN6Matrix25inverseDiagonalInPlaceCPUEv, @function
_ZN6Matrix25inverseDiagonalInPlaceCPUEv:
.LFB3812:
	.cfi_startproc
	endbr64
	pushq	%rdx
	.cfi_def_cfa_offset 16
	movl	4(%rdi), %eax
	cmpl	(%rdi), %eax
	jne	.L115
	movsd	.LC2(%rip), %xmm2
	leal	1(%rax), %r8d
	xorl	%ecx, %ecx
	xorl	%edx, %edx
	xorps	%xmm1, %xmm1
	jmp	.L116
.L115:
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC7(%rip), %rax
	jmp	.L122
.L117:
	movaps	%xmm2, %xmm3
	incl	%edx
	divsd	%xmm0, %xmm3
	movsd	%xmm3, (%rsi)
.L116:
	cmpl	%eax, %edx
	je	.L123
	movq	8(%rdi), %rsi
	movl	%ecx, %r9d
	addl	%r8d, %ecx
	leaq	(%rsi,%r9,8), %rsi
	movsd	(%rsi), %xmm0
	ucomisd	%xmm1, %xmm0
	jp	.L117
	jne	.L117
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC8(%rip), %rax
.L122:
	movq	%rax, (%rdi)
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	call	__cxa_throw@PLT
.L123:
	popq	%rax
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3812:
	.size	_ZN6Matrix25inverseDiagonalInPlaceCPUEv, .-_ZN6Matrix25inverseDiagonalInPlaceCPUEv
	.align 2
	.globl	_ZN6Matrix22transposeVectorInPlaceEv
	.type	_ZN6Matrix22transposeVectorInPlaceEv, @function
_ZN6Matrix22transposeVectorInPlaceEv:
.LFB3813:
	.cfi_startproc
	endbr64
	movl	4(%rdi), %eax
	movl	(%rdi), %edx
	movl	%eax, (%rdi)
	movl	%edx, 4(%rdi)
	ret
	.cfi_endproc
.LFE3813:
	.size	_ZN6Matrix22transposeVectorInPlaceEv, .-_ZN6Matrix22transposeVectorInPlaceEv
	.section	.rodata.str1.1
.LC9:
	.string	"Matrix is not a vector"
	.text
	.align 2
	.globl	_ZN6Matrix19vectorEuclideanNormEv
	.type	_ZN6Matrix19vectorEuclideanNormEv, @function
_ZN6Matrix19vectorEuclideanNormEv:
.LFB3814:
	.cfi_startproc
	endbr64
	movl	4(%rdi), %eax
	movl	(%rdi), %edx
	cmpl	$1, %eax
	jne	.L126
.L128:
	imull	%edx, %eax
	xorps	%xmm1, %xmm1
	xorl	%edx, %edx
	jmp	.L127
.L126:
	cmpl	$1, %edx
	je	.L128
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	movq	%rax, %rdi
	leaq	.LC9(%rip), %rax
	movq	%rax, (%rdi)
	call	__cxa_throw@PLT
.L127:
	.cfi_def_cfa_offset 8
	cmpl	%eax, %edx
	jnb	.L136
	movq	8(%rdi), %rcx
	movsd	(%rcx,%rdx,8), %xmm0
	incq	%rdx
	mulsd	%xmm0, %xmm0
	addsd	%xmm0, %xmm1
	jmp	.L127
.L136:
	movaps	%xmm1, %xmm0
	jmp	sqrt@PLT
	.cfi_endproc
.LFE3814:
	.size	_ZN6Matrix19vectorEuclideanNormEv, .-_ZN6Matrix19vectorEuclideanNormEv
	.section	.rodata.str1.1
.LC10:
	.string	"Not implemented"
	.section	.text.unlikely,"ax",@progbits
	.align 2
	.globl	_ZN6Matrix2luEv
	.type	_ZN6Matrix2luEv, @function
_ZN6Matrix2luEv:
.LFB3815:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	movl	$8, %edi
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	__cxa_allocate_exception@PLT
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	movq	%rax, %rdi
	leaq	.LC10(%rip), %rax
	movq	%rax, (%rdi)
	call	__cxa_throw@PLT
	.cfi_endproc
.LFE3815:
	.size	_ZN6Matrix2luEv, .-_ZN6Matrix2luEv
	.section	.rodata.str1.1
.LC11:
	.string	"%f "
.LC12:
	.string	"\n"
	.text
	.align 2
	.globl	_ZNK6Matrix5printEv
	.type	_ZNK6Matrix5printEv, @function
_ZNK6Matrix5printEv:
.LFB3816:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC11(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	pushq	%rcx
	.cfi_def_cfa_offset 48
.L140:
	cmpl	(%rbx), %ebp
	jnb	.L145
	xorl	%r12d, %r12d
.L142:
	movl	4(%rbx), %eax
	cmpl	%eax, %r12d
	jnb	.L146
	imull	%ebp, %eax
	movq	%r13, %rsi
	movl	$2, %edi
	leal	(%rax,%r12), %edx
	movq	8(%rbx), %rax
	incl	%r12d
	movsd	(%rax,%rdx,8), %xmm0
	movb	$1, %al
	call	__printf_chk@PLT
	jmp	.L142
.L146:
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incl	%ebp
	call	__printf_chk@PLT
	jmp	.L140
.L145:
	popq	%rdx
	.cfi_def_cfa_offset 40
	leaq	.LC12(%rip), %rsi
	popq	%rbx
	.cfi_def_cfa_offset 32
	movl	$2, %edi
	popq	%rbp
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE3816:
	.size	_ZNK6Matrix5printEv, .-_ZNK6Matrix5printEv
	.align 2
	.globl	_ZN6MatrixD2Ev
	.type	_ZN6MatrixD2Ev, @function
_ZN6MatrixD2Ev:
.LFB3818:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3818
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	8(%rdi), %rdi
	call	cudaFree@PLT
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3818:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA3818:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3818-.LLSDACSB3818
.LLSDACSB3818:
.LLSDACSE3818:
	.text
	.size	_ZN6MatrixD2Ev, .-_ZN6MatrixD2Ev
	.globl	_ZN6MatrixD1Ev
	.set	_ZN6MatrixD1Ev,_ZN6MatrixD2Ev
	.align 2
	.globl	_ZN6Matrix8doolitleERS_S0_RKS_
	.type	_ZN6Matrix8doolitleERS_S0_RKS_, @function
_ZN6Matrix8doolitleERS_S0_RKS_:
.LFB3792:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3792
	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
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movl	4(%rdx), %esi
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	cmpl	(%rdx), %esi
	je	.L150
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC7(%rip), %rax
	movq	%rax, (%rdi)
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L173
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
.LEHB0:
	call	__cxa_throw@PLT
.L150:
	leaq	24(%rsp), %r14
	movq	%rdi, %rbp
	leaq	40(%rsp), %r13
	movq	%rdx, %rbx
	movq	%r14, %rdi
	movl	%esi, %edx
	call	_ZN6Matrix6OneCPUEii
.LEHE0:
	movq	%r14, %rsi
	movq	%r13, %rdi
.LEHB1:
	call	_ZNK6Matrix11diagonalCPUEv
.LEHE1:
	movq	%r13, %rsi
	movq	%rbp, %rdi
.LEHB2:
	call	_ZN6MatrixaSERKS_
.LEHE2:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	movq	%r14, %rdi
	call	_ZN6MatrixD1Ev
	movl	4(%rbx), %esi
	movl	(%rbx), %edx
	movq	%r13, %rdi
.LEHB3:
	call	_ZN6Matrix7ZeroCPUEii
.LEHE3:
	movq	%r13, %rsi
	movq	%r12, %rdi
.LEHB4:
	call	_ZN6MatrixaSERKS_
.LEHE4:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	movl	4(%rbx), %eax
	movsd	.LC2(%rip), %xmm2
	movl	%eax, 12(%rsp)
	xorl	%eax, %eax
.L152:
	cmpl	%eax, 12(%rsp)
	je	.L149
	movq	8(%rbx), %r10
	movl	4(%r12), %edx
	movl	%eax, %r8d
	xorl	%edi, %edi
	movq	8(%r12), %rcx
.L160:
	movl	%eax, %r9d
	xorl	%esi, %esi
	xorps	%xmm0, %xmm0
.L155:
	cmpl	%edi, %esi
	je	.L175
	movl	4(%rbp), %r11d
	movq	8(%rbp), %r13
	imull	%edi, %r11d
	leal	(%r11,%rsi), %r14d
	movl	%r9d, %r11d
	incl	%esi
	addl	%edx, %r9d
	movsd	0(%r13,%r14,8), %xmm1
	mulsd	(%rcx,%r11,8), %xmm1
	addsd	%xmm1, %xmm0
	jmp	.L155
.L175:
	movl	%r8d, %edi
	addl	%edx, %r8d
	movsd	(%r10,%rdi,8), %xmm1
	subsd	%xmm0, %xmm1
	movsd	%xmm1, (%rcx,%rdi,8)
	leal	1(%rsi), %edi
	cmpl	%eax, %edi
	jle	.L160
	leal	1(%rdx), %esi
	leal	1(%rax), %r8d
	imull	%eax, %esi
	movl	%r8d, %edi
	leaq	(%rcx,%rsi,8), %r13
	movl	%edx, %esi
	imull	%r8d, %esi
	addl	%eax, %esi
	movl	%esi, 8(%rsp)
.L156:
	cmpl	%edi, 12(%rsp)
	je	.L167
	movq	8(%rbp), %r11
	movl	%eax, %r9d
	xorl	%esi, %esi
	xorps	%xmm1, %xmm1
.L158:
	cmpl	%eax, %esi
	je	.L176
	movl	4(%rbp), %r14d
	imull	%edi, %r14d
	leal	(%r14,%rsi), %r15d
	movl	%r9d, %r14d
	incl	%esi
	addl	%edx, %r9d
	movsd	(%r11,%r15,8), %xmm0
	mulsd	(%rcx,%r14,8), %xmm0
	addsd	%xmm0, %xmm1
	jmp	.L158
.L176:
	movl	8(%rsp), %esi
	incl	%edi
	addl	%edx, 8(%rsp)
	movsd	(%r10,%rsi,8), %xmm0
	subsd	%xmm1, %xmm0
	movaps	%xmm2, %xmm1
	divsd	0(%r13), %xmm1
	mulsd	%xmm1, %xmm0
	movsd	%xmm0, (%r11,%rsi,8)
	jmp	.L156
.L167:
	movl	%r8d, %eax
	jmp	.L152
.L169:
	endbr64
	movq	%rax, %rbx
.L161:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L162
.L168:
	endbr64
	movq	%rax, %rbx
.L162:
	movq	%r14, %rdi
	jmp	.L174
.L170:
	endbr64
	movq	%rax, %rbx
.L164:
	movq	%r13, %rdi
.L174:
	call	_ZN6MatrixD1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L173
	movq	%rbx, %rdi
.LEHB5:
	call	_Unwind_Resume@PLT
.LEHE5:
.L149:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L166
.L173:
	call	__stack_chk_fail@PLT
.L166:
	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
.LFE3792:
	.section	.gcc_except_table
.LLSDA3792:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3792-.LLSDACSB3792
.LLSDACSB3792:
	.uleb128 .LEHB0-.LFB3792
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB3792
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L168-.LFB3792
	.uleb128 0
	.uleb128 .LEHB2-.LFB3792
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L169-.LFB3792
	.uleb128 0
	.uleb128 .LEHB3-.LFB3792
	.uleb128 .LEHE3-.LEHB3
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB4-.LFB3792
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L170-.LFB3792
	.uleb128 0
	.uleb128 .LEHB5-.LFB3792
	.uleb128 .LEHE5-.LEHB5
	.uleb128 0
	.uleb128 0
.LLSDACSE3792:
	.text
	.size	_ZN6Matrix8doolitleERS_S0_RKS_, .-_ZN6Matrix8doolitleERS_S0_RKS_
	.align 2
	.globl	_ZN6Matrix8LUMehtodERKS_S1_
	.type	_ZN6Matrix8LUMehtodERKS_S1_, @function
_ZN6Matrix8LUMehtodERKS_S1_:
.LFB3790:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3790
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movq	%rsi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
	subq	$64, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rbp
	leaq	24(%rsp), %r12
	movq	%rbp, %rdi
.LEHB6:
	call	_ZN6Matrix4StubEv
.LEHE6:
	movq	%r12, %rdi
.LEHB7:
	call	_ZN6Matrix4StubEv
.LEHE7:
	movq	%r13, %rdx
	movq	%r12, %rsi
	movq	%rbp, %rdi
.LEHB8:
	call	_ZN6Matrix8doolitleERS_S0_RKS_
	leaq	40(%rsp), %r13
	movq	%r14, %rdx
	movq	%rbp, %rsi
	movq	%r13, %rdi
	call	_ZN6Matrix19ForwardSubstitutionERKS_S1_
.LEHE8:
	movq	%r13, %rdx
	movq	%r12, %rsi
	movq	%rbx, %rdi
.LEHB9:
	call	_ZN6Matrix20BackwardSubstitutionERKS_S1_
.LEHE9:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	movq	%r12, %rdi
	call	_ZN6MatrixD1Ev
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L182
	jmp	.L186
.L185:
	endbr64
	movq	%rax, %rbx
.L178:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L179
.L184:
	endbr64
	movq	%rax, %rbx
.L179:
	movq	%r12, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L180
.L183:
	endbr64
	movq	%rax, %rbx
.L180:
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L186
	movq	%rbx, %rdi
.LEHB10:
	call	_Unwind_Resume@PLT
.LEHE10:
.L186:
	call	__stack_chk_fail@PLT
.L182:
	addq	$64, %rsp
	.cfi_def_cfa_offset 48
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3790:
	.section	.gcc_except_table
.LLSDA3790:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3790-.LLSDACSB3790
.LLSDACSB3790:
	.uleb128 .LEHB6-.LFB3790
	.uleb128 .LEHE6-.LEHB6
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB7-.LFB3790
	.uleb128 .LEHE7-.LEHB7
	.uleb128 .L183-.LFB3790
	.uleb128 0
	.uleb128 .LEHB8-.LFB3790
	.uleb128 .LEHE8-.LEHB8
	.uleb128 .L184-.LFB3790
	.uleb128 0
	.uleb128 .LEHB9-.LFB3790
	.uleb128 .LEHE9-.LEHB9
	.uleb128 .L185-.LFB3790
	.uleb128 0
	.uleb128 .LEHB10-.LFB3790
	.uleb128 .LEHE10-.LEHB10
	.uleb128 0
	.uleb128 0
.LLSDACSE3790:
	.text
	.size	_ZN6Matrix8LUMehtodERKS_S1_, .-_ZN6Matrix8LUMehtodERKS_S1_
	.align 2
	.globl	_ZN6Matrix15LUMehtodOptimalERKS_S1_
	.type	_ZN6Matrix15LUMehtodOptimalERKS_S1_, @function
_ZN6Matrix15LUMehtodOptimalERKS_S1_:
.LFB4579:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	call	_ZN6Matrix8LUMehtodERKS_S1_
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L189
	call	__stack_chk_fail@PLT
.L189:
	addq	$16, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4579:
	.size	_ZN6Matrix15LUMehtodOptimalERKS_S1_, .-_ZN6Matrix15LUMehtodOptimalERKS_S1_
	.align 2
	.globl	_ZN6Matrix10createTestERS_S0_S0_i
	.type	_ZN6Matrix10createTestERS_S0_S0_i, @function
_ZN6Matrix10createTestERS_S0_S0_i:
.LFB3794:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3794
	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
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	xorl	%edi, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%ecx, %ebx
	subq	$104, %rsp
	.cfi_def_cfa_offset 160
	movq	%rsi, 24(%rsp)
	leaq	72(%rsp), %r13
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	movl	$1, %esi
	movq	%r13, %rdi
	movl	%ebx, %edx
.LEHB11:
	call	_ZN6MatrixC1Eii
.LEHE11:
	movq	%r13, %rsi
	movq	%r12, %rdi
.LEHB12:
	call	_ZN6MatrixaSERKS_
.LEHE12:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	movl	%ebx, %esi
	movq	%r13, %rdi
	movl	%ebx, %edx
.LEHB13:
	call	_ZN6MatrixC1Eii
.LEHE13:
	movq	%r13, %rsi
	movq	%rbp, %rdi
.LEHB14:
	call	_ZN6MatrixaSERKS_
.LEHE14:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	movl	$1, %esi
	movq	%r13, %rdi
	movl	%ebx, %edx
.LEHB15:
	call	_ZN6MatrixC1Eii
.LEHE15:
	movq	24(%rsp), %rdi
	movq	%r13, %rsi
.LEHB16:
	call	_ZN6MatrixaSERKS_
.LEHE16:
	movq	%r13, %rdi
	xorl	%r13d, %r13d
	call	_ZN6MatrixD1Ev
.L192:
	cmpl	%r13d, %ebx
	jle	.L224
	movq	8(%r12), %rax
	leaq	(%rax,%r13,8), %r14
	incq	%r13
	call	_ZZN6Matrix10createTestERS_S0_S0_iENKUliE_clEi.isra.0
	movsd	%xmm0, (%r14)
	jmp	.L192
.L224:
	leal	1(%rbx), %eax
	xorl	%r15d, %r15d
	xorl	%r14d, %r14d
	movl	%eax, 36(%rsp)
	xorl	%eax, %eax
	movl	%eax, 32(%rsp)
.L194:
	movslq	%ebx, %rax
	cmpl	%ebx, %r14d
	movq	%rax, 40(%rsp)
	jge	.L225
	movslq	32(%rsp), %rax
	xorl	%r13d, %r13d
	movq	$0x000000000, 8(%rsp)
	salq	$3, %rax
	movq	%rax, 16(%rsp)
.L199:
	cmpl	%r13d, %r14d
	je	.L195
	movq	8(%rbp), %rdx
	leal	(%r15,%r13), %eax
	cltq
	leaq	(%rdx,%rax,8), %rdx
	movq	%rax, 56(%rsp)
	movq	%rdx, 48(%rsp)
	call	_ZZN6Matrix10createTestERS_S0_S0_iENKUliE_clEi.isra.0
	movq	48(%rsp), %rdx
	movq	56(%rsp), %rax
	movsd	%xmm0, (%rdx)
	movq	8(%rbp), %rdx
	movsd	(%rdx,%rax,8), %xmm0
	andps	.LC6(%rip), %xmm0
	addsd	8(%rsp), %xmm0
	movsd	%xmm0, 8(%rsp)
.L195:
	call	_ZZN6Matrix10createTestERS_S0_S0_iENKUliE_clEi.isra.0
	movaps	%xmm0, %xmm1
	xorps	%xmm0, %xmm0
	comisd	%xmm0, %xmm1
	jbe	.L217
	call	_ZZN6Matrix10createTestERS_S0_S0_iENKUliE_clEi.isra.0
	movq	8(%rbp), %rax
	addsd	8(%rsp), %xmm0
	jmp	.L220
.L217:
	call	_ZZN6Matrix10createTestERS_S0_S0_iENKUliE_clEi.isra.0
	subsd	8(%rsp), %xmm0
	movq	8(%rbp), %rax
.L220:
	movq	16(%rsp), %rcx
	incq	%r13
	movsd	%xmm0, (%rax,%rcx)
	cmpq	%r13, 40(%rsp)
	jne	.L199
	movl	36(%rsp), %ecx
	incl	%r14d
	addl	%ecx, 32(%rsp)
	addl	%ebx, %r15d
	jmp	.L194
.L225:
	movq	%rax, %rcx
	xorl	%esi, %esi
	xorl	%edx, %edx
	salq	$3, %rcx
.L201:
	cmpl	%edx, %ebx
	jle	.L226
	movq	8(%rbp), %r8
	movq	8(%r12), %rdi
	xorl	%eax, %eax
	xorps	%xmm0, %xmm0
	addq	%rsi, %r8
.L202:
	movsd	(%r8,%rax,8), %xmm1
	mulsd	(%rdi,%rax,8), %xmm1
	incq	%rax
	addsd	%xmm1, %xmm0
	cmpl	%eax, %ebx
	jg	.L202
	movq	24(%rsp), %rax
	addq	%rcx, %rsi
	movq	8(%rax), %rax
	movsd	%xmm0, (%rax,%rdx,8)
	incq	%rdx
	jmp	.L201
.L211:
	endbr64
	jmp	.L223
.L212:
	endbr64
	jmp	.L223
.L213:
	endbr64
.L223:
	movq	%rax, %rbx
.L208:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L222
	movq	%rbx, %rdi
.LEHB17:
	call	_Unwind_Resume@PLT
.LEHE17:
.L226:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L210
.L222:
	call	__stack_chk_fail@PLT
.L210:
	addq	$104, %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
.LFE3794:
	.section	.gcc_except_table
.LLSDA3794:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3794-.LLSDACSB3794
.LLSDACSB3794:
	.uleb128 .LEHB11-.LFB3794
	.uleb128 .LEHE11-.LEHB11
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB12-.LFB3794
	.uleb128 .LEHE12-.LEHB12
	.uleb128 .L211-.LFB3794
	.uleb128 0
	.uleb128 .LEHB13-.LFB3794
	.uleb128 .LEHE13-.LEHB13
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB14-.LFB3794
	.uleb128 .LEHE14-.LEHB14
	.uleb128 .L212-.LFB3794
	.uleb128 0
	.uleb128 .LEHB15-.LFB3794
	.uleb128 .LEHE15-.LEHB15
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB16-.LFB3794
	.uleb128 .LEHE16-.LEHB16
	.uleb128 .L213-.LFB3794
	.uleb128 0
	.uleb128 .LEHB17-.LFB3794
	.uleb128 .LEHE17-.LEHB17
	.uleb128 0
	.uleb128 0
.LLSDACSE3794:
	.text
	.size	_ZN6Matrix10createTestERS_S0_S0_i, .-_ZN6Matrix10createTestERS_S0_S0_i
	.align 2
	.globl	_ZN6Matrix10createTaskERS_S0_i
	.type	_ZN6Matrix10createTaskERS_S0_i, @function
_ZN6Matrix10createTaskERS_S0_i:
.LFB3798:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3798
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movq	%rdi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rsi, %r12
	movl	%edx, %esi
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%edx, %ebx
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rbp
	movq	%rbp, %rdi
.LEHB18:
	call	_ZN6Matrix7ZeroCPUEii
.LEHE18:
	movq	%rbp, %rsi
	movq	%r13, %rdi
.LEHB19:
	call	_ZN6MatrixaSERKS_
.LEHE19:
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movl	$1, %esi
	movq	%rbp, %rdi
	movl	%ebx, %edx
.LEHB20:
	call	_ZN6MatrixC1Eii
.LEHE20:
	movq	%rbp, %rsi
	movq	%r12, %rdi
.LEHB21:
	call	_ZN6MatrixaSERKS_
.LEHE21:
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movl	%ebx, %edi
	movslq	%ebx, %rax
	leal	1(%rbx), %r9d
	imull	%ebx, %edi
	movsd	.LC13(%rip), %xmm1
	xorl	%edx, %edx
	xorl	%esi, %esi
	movsd	.LC14(%rip), %xmm0
	leaq	8(,%rax,8), %r8
	movl	$16, %eax
.L228:
	cmpl	%ebx, %esi
	jge	.L252
	movq	8(%r13), %rcx
	movsd	%xmm1, -16(%rcx,%rax)
	testl	%edx, %edx
	je	.L230
	movsd	%xmm0, -24(%rcx,%rax)
	cmpl	$1, %edx
	je	.L230
	movsd	%xmm0, -32(%rcx,%rax)
.L230:
	leal	1(%rdx), %r10d
	cmpl	%r10d, %edi
	jle	.L232
	movsd	%xmm0, -8(%rcx,%rax)
.L232:
	leal	2(%rdx), %r10d
	cmpl	%r10d, %edi
	jle	.L233
	movsd	%xmm0, (%rcx,%rax)
.L233:
	incl	%esi
	addq	%r8, %rax
	addl	%r9d, %edx
	jmp	.L228
.L252:
	xorl	%ebp, %ebp
.L235:
	cmpl	%ebp, %ebx
	jle	.L253
	leal	(%rbp,%rbp), %eax
	cvtsi2sdl	%eax, %xmm0
	call	sin@PLT
	movq	8(%r12), %rax
	movsd	%xmm0, (%rax,%rbp,8)
	incq	%rbp
	jmp	.L235
.L242:
	endbr64
	jmp	.L251
.L243:
	endbr64
.L251:
	movq	%rax, %rbx
.L239:
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L250
	movq	%rbx, %rdi
.LEHB22:
	call	_Unwind_Resume@PLT
.LEHE22:
.L253:
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L241
.L250:
	call	__stack_chk_fail@PLT
.L241:
	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
.LFE3798:
	.section	.gcc_except_table
.LLSDA3798:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3798-.LLSDACSB3798
.LLSDACSB3798:
	.uleb128 .LEHB18-.LFB3798
	.uleb128 .LEHE18-.LEHB18
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB19-.LFB3798
	.uleb128 .LEHE19-.LEHB19
	.uleb128 .L242-.LFB3798
	.uleb128 0
	.uleb128 .LEHB20-.LFB3798
	.uleb128 .LEHE20-.LEHB20
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB21-.LFB3798
	.uleb128 .LEHE21-.LEHB21
	.uleb128 .L243-.LFB3798
	.uleb128 0
	.uleb128 .LEHB22-.LFB3798
	.uleb128 .LEHE22-.LEHB22
	.uleb128 0
	.uleb128 0
.LLSDACSE3798:
	.text
	.size	_ZN6Matrix10createTaskERS_S0_i, .-_ZN6Matrix10createTaskERS_S0_i
	.align 2
	.globl	_ZN6Matrix11createTaskCERS_S0_
	.type	_ZN6Matrix11createTaskCERS_S0_, @function
_ZN6Matrix11createTaskCERS_S0_:
.LFB3799:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3799
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movl	$994, %edx
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rsi, %rbp
	movl	$994, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rbx
	movq	%rbx, %rdi
.LEHB23:
	call	_ZN6Matrix7ZeroCPUEii
.LEHE23:
	movq	%rbx, %rsi
	movq	%r12, %rdi
.LEHB24:
	call	_ZN6MatrixaSERKS_
.LEHE24:
	movq	%rbx, %rdi
	call	_ZN6MatrixD1Ev
	movl	$1, %esi
	movq	%rbx, %rdi
	movl	$994, %edx
.LEHB25:
	call	_ZN6MatrixC1Eii
.LEHE25:
	movq	%rbx, %rsi
	movq	%rbp, %rdi
.LEHB26:
	call	_ZN6MatrixaSERKS_
.LEHE26:
	movq	%rbx, %rdi
	call	_ZN6MatrixD1Ev
	movq	8(%r12), %rax
	movsd	.LC15(%rip), %xmm1
	xorl	%edx, %edx
	movsd	.LC14(%rip), %xmm0
.L257:
	movsd	%xmm1, (%rax)
	testl	%edx, %edx
	je	.L255
	movsd	%xmm0, -8(%rax)
	movsd	%xmm0, -16(%rax)
	cmpl	$993, %edx
	jne	.L255
	movq	8(%rbp), %rbp
	xorl	%ebx, %ebx
	jmp	.L256
.L255:
	movsd	%xmm0, 8(%rax)
	incl	%edx
	addq	$7960, %rax
	movsd	%xmm0, -7944(%rax)
	jmp	.L257
.L256:
	leal	(%rbx,%rbx), %eax
	cvtsi2sdl	%eax, %xmm0
	call	sin@PLT
	movsd	%xmm0, 0(%rbp,%rbx,8)
	incq	%rbx
	cmpq	$994, %rbx
	jne	.L256
	jmp	.L274
.L263:
	endbr64
	jmp	.L273
.L264:
	endbr64
.L273:
	movq	%rax, %rbp
.L260:
	movq	%rbx, %rdi
	call	_ZN6MatrixD1Ev
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L272
	movq	%rbp, %rdi
.LEHB27:
	call	_Unwind_Resume@PLT
.LEHE27:
.L274:
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L262
.L272:
	call	__stack_chk_fail@PLT
.L262:
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3799:
	.section	.gcc_except_table
.LLSDA3799:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3799-.LLSDACSB3799
.LLSDACSB3799:
	.uleb128 .LEHB23-.LFB3799
	.uleb128 .LEHE23-.LEHB23
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB24-.LFB3799
	.uleb128 .LEHE24-.LEHB24
	.uleb128 .L263-.LFB3799
	.uleb128 0
	.uleb128 .LEHB25-.LFB3799
	.uleb128 .LEHE25-.LEHB25
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB26-.LFB3799
	.uleb128 .LEHE26-.LEHB26
	.uleb128 .L264-.LFB3799
	.uleb128 0
	.uleb128 .LEHB27-.LFB3799
	.uleb128 .LEHE27-.LEHB27
	.uleb128 0
	.uleb128 0
.LLSDACSE3799:
	.text
	.size	_ZN6Matrix11createTaskCERS_S0_, .-_ZN6Matrix11createTaskCERS_S0_
	.globl	_Z34__device_stub__Z13ZeroGPUKerneliPdiPd
	.type	_Z34__device_stub__Z13ZeroGPUKerneliPdiPd, @function
_Z34__device_stub__Z13ZeroGPUKerneliPdiPd:
.LFB3852:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movl	%edi, 12(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	32(%rsp), %rdi
	leaq	44(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 56(%rsp)
	movq	%rsp, %rax
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L275
	pushq	24(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z13ZeroGPUKerneliPd(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 112
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	72(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L275:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L277
	call	__stack_chk_fail@PLT
.L277:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3852:
	.size	_Z34__device_stub__Z13ZeroGPUKerneliPdiPd, .-_Z34__device_stub__Z13ZeroGPUKerneliPdiPd
	.globl	_Z13ZeroGPUKerneliPd
	.type	_Z13ZeroGPUKerneliPd, @function
_Z13ZeroGPUKerneliPd:
.LFB3853:
	.cfi_startproc
	endbr64
	jmp	_Z34__device_stub__Z13ZeroGPUKerneliPdiPd
	.cfi_endproc
.LFE3853:
	.size	_Z13ZeroGPUKerneliPd, .-_Z13ZeroGPUKerneliPd
	.align 2
	.globl	_ZN6Matrix7ZeroGPUEii
	.type	_ZN6Matrix7ZeroGPUEii, @function
_ZN6Matrix7ZeroGPUEii:
.LFB3781:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	%esi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	imull	%edx, %r13d
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	movslq	%r13d, %rsi
	salq	$3, %rsi
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movl	$256, %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leal	255(%r13), %eax
	cltd
	idivl	%ecx
	movl	$16777217, %edx
	movl	$1, %ecx
	salq	$8, %rdx
	movl	%eax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 20(%rsp)
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L281
	movq	8(%rsp), %rsi
	movl	%r13d, %edi
	call	_Z34__device_stub__Z13ZeroGPUKerneliPdiPd
.L281:
	call	cudaDeviceSynchronize@PLT
	movq	8(%rsp), %rax
	movl	%r12d, (%rbx)
	movl	%ebp, 4(%rbx)
	movq	%rax, 8(%rbx)
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L282
	call	__stack_chk_fail@PLT
.L282:
	addq	$56, %rsp
	.cfi_def_cfa_offset 40
	movq	%rbx, %rax
	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
.LFE3781:
	.size	_ZN6Matrix7ZeroGPUEii, .-_ZN6Matrix7ZeroGPUEii
	.globl	_Z37__device_stub__Z14doolitleKerneliPdS_iPdS_
	.type	_Z37__device_stub__Z14doolitleKerneliPdS_iPdS_, @function
_Z37__device_stub__Z14doolitleKerneliPdS_iPdS_:
.LFB3854:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L284
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z14doolitleKerneliPdS_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L284:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L286
	call	__stack_chk_fail@PLT
.L286:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3854:
	.size	_Z37__device_stub__Z14doolitleKerneliPdS_iPdS_, .-_Z37__device_stub__Z14doolitleKerneliPdS_iPdS_
	.globl	_Z14doolitleKerneliPdS_
	.type	_Z14doolitleKerneliPdS_, @function
_Z14doolitleKerneliPdS_:
.LFB3855:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z14doolitleKerneliPdS_iPdS_
	.cfi_endproc
.LFE3855:
	.size	_Z14doolitleKerneliPdS_, .-_Z14doolitleKerneliPdS_
	.globl	_Z33__device_stub__Z10copyKerneliPdS_iPdS_
	.type	_Z33__device_stub__Z10copyKerneliPdS_iPdS_, @function
_Z33__device_stub__Z10copyKerneliPdS_iPdS_:
.LFB3856:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L289
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z10copyKerneliPdS_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L289:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L291
	call	__stack_chk_fail@PLT
.L291:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3856:
	.size	_Z33__device_stub__Z10copyKerneliPdS_iPdS_, .-_Z33__device_stub__Z10copyKerneliPdS_iPdS_
	.globl	_Z10copyKerneliPdS_
	.type	_Z10copyKerneliPdS_, @function
_Z10copyKerneliPdS_:
.LFB3857:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z10copyKerneliPdS_iPdS_
	.cfi_endproc
.LFE3857:
	.size	_Z10copyKerneliPdS_, .-_Z10copyKerneliPdS_
	.align 2
	.globl	_ZN6Matrix7copyGPUERS_RKS_
	.type	_ZN6Matrix7copyGPUERS_RKS_, @function
_ZN6Matrix7copyGPUERS_RKS_:
.LFB3801:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	$16777217, %edx
	movq	%rsi, %rbp
	xorl	%r9d, %r9d
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	xorl	%r8d, %r8d
	salq	$8, %rdx
	movl	$1, %ecx
	subq	$40, %rsp
	.cfi_def_cfa_offset 64
	movl	4(%rdi), %eax
	imull	(%rdi), %eax
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L295
	movl	4(%rbx), %edi
	imull	(%rbx), %edi
	movq	8(%rbp), %rdx
	movq	8(%rbx), %rsi
	call	_Z33__device_stub__Z10copyKerneliPdS_iPdS_
.L295:
	addq	$40, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3801:
	.size	_ZN6Matrix7copyGPUERS_RKS_, .-_ZN6Matrix7copyGPUERS_RKS_
	.globl	_Z45__device_stub__Z22separateDiagonalKerneliPdS_iPdS_
	.type	_Z45__device_stub__Z22separateDiagonalKerneliPdS_iPdS_, @function
_Z45__device_stub__Z22separateDiagonalKerneliPdS_iPdS_:
.LFB3858:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L297
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z22separateDiagonalKerneliPdS_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L297:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L299
	call	__stack_chk_fail@PLT
.L299:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3858:
	.size	_Z45__device_stub__Z22separateDiagonalKerneliPdS_iPdS_, .-_Z45__device_stub__Z22separateDiagonalKerneliPdS_iPdS_
	.globl	_Z22separateDiagonalKerneliPdS_
	.type	_Z22separateDiagonalKerneliPdS_, @function
_Z22separateDiagonalKerneliPdS_:
.LFB3859:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z22separateDiagonalKerneliPdS_iPdS_
	.cfi_endproc
.LFE3859:
	.size	_Z22separateDiagonalKerneliPdS_, .-_Z22separateDiagonalKerneliPdS_
	.align 2
	.globl	_ZN6Matrix29separateDiagonalAndInverseGPUERS_S0_
	.type	_ZN6Matrix29separateDiagonalAndInverseGPUERS_S0_, @function
_ZN6Matrix29separateDiagonalAndInverseGPUERS_S0_:
.LFB3802:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	$16777217, %edx
	movq	%rdi, %rbp
	xorl	%r9d, %r9d
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	xorl	%r8d, %r8d
	salq	$8, %rdx
	movl	$1, %ecx
	subq	$40, %rsp
	.cfi_def_cfa_offset 64
	movl	4(%rsi), %eax
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L303
	movq	8(%rbx), %rdx
	movq	8(%rbp), %rsi
	movl	4(%rbx), %edi
	call	_Z45__device_stub__Z22separateDiagonalKerneliPdS_iPdS_
.L303:
	addq	$40, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3802:
	.size	_ZN6Matrix29separateDiagonalAndInverseGPUERS_S0_, .-_ZN6Matrix29separateDiagonalAndInverseGPUERS_S0_
	.globl	_Z43__device_stub__Z19separateUpperKerneliiPdS_iiPdS_
	.type	_Z43__device_stub__Z19separateUpperKerneliiPdS_iiPdS_, @function
_Z43__device_stub__Z19separateUpperKerneliiPdS_iiPdS_:
.LFB3860:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L305
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z19separateUpperKerneliiPdS_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L305:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L307
	call	__stack_chk_fail@PLT
.L307:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3860:
	.size	_Z43__device_stub__Z19separateUpperKerneliiPdS_iiPdS_, .-_Z43__device_stub__Z19separateUpperKerneliiPdS_iiPdS_
	.globl	_Z19separateUpperKerneliiPdS_
	.type	_Z19separateUpperKerneliiPdS_, @function
_Z19separateUpperKerneliiPdS_:
.LFB3861:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z19separateUpperKerneliiPdS_iiPdS_
	.cfi_endproc
.LFE3861:
	.size	_Z19separateUpperKerneliiPdS_, .-_Z19separateUpperKerneliiPdS_
	.align 2
	.globl	_ZN6Matrix16separateUpperGPUERS_S0_
	.type	_ZN6Matrix16separateUpperGPUERS_S0_, @function
_ZN6Matrix16separateUpperGPUERS_S0_:
.LFB3803:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	$16777217, %edx
	movq	%rdi, %rbp
	xorl	%r9d, %r9d
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	xorl	%r8d, %r8d
	salq	$8, %rdx
	movl	$1, %ecx
	subq	$40, %rsp
	.cfi_def_cfa_offset 64
	movl	4(%rsi), %eax
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L311
	movl	4(%rbx), %esi
	movl	(%rbx), %edi
	movq	8(%rbx), %rcx
	movq	8(%rbp), %rdx
	imull	%esi, %edi
	call	_Z43__device_stub__Z19separateUpperKerneliiPdS_iiPdS_
.L311:
	addq	$40, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3803:
	.size	_ZN6Matrix16separateUpperGPUERS_S0_, .-_ZN6Matrix16separateUpperGPUERS_S0_
	.globl	_Z49__device_stub__Z28additiveInverseInPlaceKerneliPdiPd
	.type	_Z49__device_stub__Z28additiveInverseInPlaceKerneliPdiPd, @function
_Z49__device_stub__Z28additiveInverseInPlaceKerneliPdiPd:
.LFB3862:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movl	%edi, 12(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	32(%rsp), %rdi
	leaq	44(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 56(%rsp)
	movq	%rsp, %rax
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L313
	pushq	24(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z28additiveInverseInPlaceKerneliPd(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 112
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	72(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L313:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L315
	call	__stack_chk_fail@PLT
.L315:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3862:
	.size	_Z49__device_stub__Z28additiveInverseInPlaceKerneliPdiPd, .-_Z49__device_stub__Z28additiveInverseInPlaceKerneliPdiPd
	.globl	_Z28additiveInverseInPlaceKerneliPd
	.type	_Z28additiveInverseInPlaceKerneliPd, @function
_Z28additiveInverseInPlaceKerneliPd:
.LFB3863:
	.cfi_startproc
	endbr64
	jmp	_Z49__device_stub__Z28additiveInverseInPlaceKerneliPdiPd
	.cfi_endproc
.LFE3863:
	.size	_Z28additiveInverseInPlaceKerneliPd, .-_Z28additiveInverseInPlaceKerneliPd
	.align 2
	.globl	_ZN6Matrix25additiveInverseInPlaceGPUERS_
	.type	_ZN6Matrix25additiveInverseInPlaceGPUERS_, @function
_ZN6Matrix25additiveInverseInPlaceGPUERS_:
.LFB3804:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	$16777217, %edx
	movq	%rdi, %rbx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$8, %rdx
	movl	$1, %ecx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	movl	(%rdi), %eax
	imull	4(%rdi), %eax
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L319
	movl	(%rbx), %edi
	imull	4(%rbx), %edi
	movq	8(%rbx), %rsi
	call	_Z49__device_stub__Z28additiveInverseInPlaceKerneliPdiPd
.L319:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3804:
	.size	_ZN6Matrix25additiveInverseInPlaceGPUERS_, .-_ZN6Matrix25additiveInverseInPlaceGPUERS_
	.globl	_Z50__device_stub__Z25forwardSubstitutionKerneliPdS_S_iPdS_S_
	.type	_Z50__device_stub__Z25forwardSubstitutionKerneliPdS_S_iPdS_S_, @function
_Z50__device_stub__Z25forwardSubstitutionKerneliPdS_S_iPdS_S_:
.LFB3864:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L321
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z25forwardSubstitutionKerneliPdS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L321:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L323
	call	__stack_chk_fail@PLT
.L323:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3864:
	.size	_Z50__device_stub__Z25forwardSubstitutionKerneliPdS_S_iPdS_S_, .-_Z50__device_stub__Z25forwardSubstitutionKerneliPdS_S_iPdS_S_
	.globl	_Z25forwardSubstitutionKerneliPdS_S_
	.type	_Z25forwardSubstitutionKerneliPdS_S_, @function
_Z25forwardSubstitutionKerneliPdS_S_:
.LFB3865:
	.cfi_startproc
	endbr64
	jmp	_Z50__device_stub__Z25forwardSubstitutionKerneliPdS_S_iPdS_S_
	.cfi_endproc
.LFE3865:
	.size	_Z25forwardSubstitutionKerneliPdS_S_, .-_Z25forwardSubstitutionKerneliPdS_S_
	.align 2
	.globl	_ZN6Matrix22forwardSubstitutionGPUERS_RKS_S2_
	.type	_ZN6Matrix22forwardSubstitutionGPUERS_RKS_S2_, @function
_ZN6Matrix22forwardSubstitutionGPUERS_RKS_S2_:
.LFB3805:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	4(%rsi), %eax
	cvtsi2ssq	%rax, %xmm0
	call	log2f@PLT
	call	ceilf@PLT
	cvtss2sd	%xmm0, %xmm1
	movsd	.LC16(%rip), %xmm0
	call	pow@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	cvttsd2sil	%xmm0, %eax
	movq	%rdi, 24(%rsp)
	movl	28(%rsp), %ecx
	movl	%eax, 20(%rsp)
	movq	20(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L329
	movq	8(%r12), %rcx
	movq	8(%rbp), %rdx
	movq	8(%rbx), %rsi
	movl	4(%rbx), %edi
	call	_Z50__device_stub__Z25forwardSubstitutionKerneliPdS_S_iPdS_S_
.L329:
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3805:
	.size	_ZN6Matrix22forwardSubstitutionGPUERS_RKS_S2_, .-_ZN6Matrix22forwardSubstitutionGPUERS_RKS_S2_
	.globl	_Z35__device_stub__Z9mulKerneliiiPdS_S_iiiPdS_S_
	.type	_Z35__device_stub__Z9mulKerneliiiPdS_S_iiiPdS_S_, @function
_Z35__device_stub__Z9mulKerneliiiPdS_S_iiiPdS_S_:
.LFB3866:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	76(%rsp), %rsi
	movl	%edx, 36(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 24(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 16(%rsp)
	movq	%r9, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 88(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L331
	pushq	56(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z9mulKerneliiiPdS_S_(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 176
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	104(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L331:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L333
	call	__stack_chk_fail@PLT
.L333:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3866:
	.size	_Z35__device_stub__Z9mulKerneliiiPdS_S_iiiPdS_S_, .-_Z35__device_stub__Z9mulKerneliiiPdS_S_iiiPdS_S_
	.globl	_Z9mulKerneliiiPdS_S_
	.type	_Z9mulKerneliiiPdS_S_, @function
_Z9mulKerneliiiPdS_S_:
.LFB3867:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z9mulKerneliiiPdS_S_iiiPdS_S_
	.cfi_endproc
.LFE3867:
	.size	_Z9mulKerneliiiPdS_S_, .-_Z9mulKerneliiiPdS_S_
	.align 2
	.globl	_ZN6Matrix6refMulERS_RKS_S2_
	.type	_ZN6Matrix6refMulERS_RKS_S2_, @function
_ZN6Matrix6refMulERS_RKS_S2_:
.LFB3820:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdi, %r12
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdx, %rbp
	movl	$1, %ecx
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	(%rsi), %eax
	imull	4(%rdx), %eax
	movl	$16777217, %edx
	salq	$8, %rdx
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L337
	movl	4(%rbp), %esi
	movl	(%rbx), %edx
	movq	8(%rbx), %rcx
	movl	4(%rbx), %edi
	imull	%esi, %edx
	movq	8(%r12), %r9
	movq	8(%rbp), %r8
	call	_Z35__device_stub__Z9mulKerneliiiPdS_S_iiiPdS_S_
.L337:
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3820:
	.size	_ZN6Matrix6refMulERS_RKS_S2_, .-_ZN6Matrix6refMulERS_RKS_S2_
	.section	.rodata.str1.1
.LC17:
	.string	"wrong dimensions for multiplication"
	.text
	.globl	_ZmlRK6MatrixS1_
	.type	_ZmlRK6MatrixS1_, @function
_ZmlRK6MatrixS1_:
.LFB3822:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$48, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	movl	(%rdx), %eax
	cmpl	%eax, 4(%rsi)
	je	.L340
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC17(%rip), %rax
	movq	%rax, (%rdi)
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L345
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	call	__cxa_throw@PLT
.L340:
	movq	%rsi, %rbx
	movl	4(%rdx), %esi
	movq	%rdx, %rbp
	movq	%rdi, %r12
	imull	(%rbx), %esi
	leaq	8(%rsp), %rdi
	salq	$3, %rsi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movl	(%rbx), %eax
	imull	4(%rbp), %eax
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$8, %rdx
	movl	$1, %ecx
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 20(%rsp)
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L342
	movl	4(%rbp), %esi
	movl	(%rbx), %edx
	movq	8(%rbx), %rcx
	movl	4(%rbx), %edi
	imull	%esi, %edx
	movq	8(%rsp), %r9
	movq	8(%rbp), %r8
	call	_Z35__device_stub__Z9mulKerneliiiPdS_S_iiiPdS_S_
.L342:
	call	cudaDeviceSynchronize@PLT
	movl	4(%rbp), %eax
	movl	(%rbx), %edx
	movl	%eax, 4(%r12)
	movq	8(%rsp), %rax
	movl	%edx, (%r12)
	movq	%rax, 8(%r12)
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L343
.L345:
	call	__stack_chk_fail@PLT
.L343:
	addq	$48, %rsp
	.cfi_def_cfa_offset 32
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3822:
	.size	_ZmlRK6MatrixS1_, .-_ZmlRK6MatrixS1_
	.globl	_Z40__device_stub__Z13mulDiagKerneliiiPdS_S_iiiPdS_S_
	.type	_Z40__device_stub__Z13mulDiagKerneliiiPdS_S_iiiPdS_S_, @function
_Z40__device_stub__Z13mulDiagKerneliiiPdS_S_iiiPdS_S_:
.LFB3868:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	76(%rsp), %rsi
	movl	%edx, 36(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 24(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 16(%rsp)
	movq	%r9, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 88(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L346
	pushq	56(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z13mulDiagKerneliiiPdS_S_(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 176
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	104(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L346:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L348
	call	__stack_chk_fail@PLT
.L348:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3868:
	.size	_Z40__device_stub__Z13mulDiagKerneliiiPdS_S_iiiPdS_S_, .-_Z40__device_stub__Z13mulDiagKerneliiiPdS_S_iiiPdS_S_
	.globl	_Z13mulDiagKerneliiiPdS_S_
	.type	_Z13mulDiagKerneliiiPdS_S_, @function
_Z13mulDiagKerneliiiPdS_S_:
.LFB3869:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z13mulDiagKerneliiiPdS_S_iiiPdS_S_
	.cfi_endproc
.LFE3869:
	.size	_Z13mulDiagKerneliiiPdS_S_, .-_Z13mulDiagKerneliiiPdS_S_
	.align 2
	.globl	_ZN6Matrix10refMulDiagERS_RKS_S2_
	.type	_ZN6Matrix10refMulDiagERS_RKS_S2_, @function
_ZN6Matrix10refMulDiagERS_RKS_S2_:
.LFB3821:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdi, %r12
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdx, %rbp
	movl	$1, %ecx
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	(%rsi), %eax
	imull	4(%rdx), %eax
	movl	$16777217, %edx
	salq	$8, %rdx
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L352
	movl	4(%rbp), %esi
	movl	(%rbx), %edx
	movq	8(%rbx), %rcx
	movl	4(%rbx), %edi
	imull	%esi, %edx
	movq	8(%r12), %r9
	movq	8(%rbp), %r8
	call	_Z40__device_stub__Z13mulDiagKerneliiiPdS_S_iiiPdS_S_
.L352:
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3821:
	.size	_ZN6Matrix10refMulDiagERS_RKS_S2_, .-_ZN6Matrix10refMulDiagERS_RKS_S2_
	.globl	_Z33__device_stub__Z9addKerneliPdS_S_iPdS_S_
	.type	_Z33__device_stub__Z9addKerneliPdS_S_iPdS_S_, @function
_Z33__device_stub__Z9addKerneliPdS_S_iPdS_S_:
.LFB3870:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L354
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z9addKerneliPdS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L354:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L356
	call	__stack_chk_fail@PLT
.L356:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3870:
	.size	_Z33__device_stub__Z9addKerneliPdS_S_iPdS_S_, .-_Z33__device_stub__Z9addKerneliPdS_S_iPdS_S_
	.globl	_Z9addKerneliPdS_S_
	.type	_Z9addKerneliPdS_S_, @function
_Z9addKerneliPdS_S_:
.LFB3871:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z9addKerneliPdS_S_iPdS_S_
	.cfi_endproc
.LFE3871:
	.size	_Z9addKerneliPdS_S_, .-_Z9addKerneliPdS_S_
	.align 2
	.globl	_ZN6Matrix6refAddERS_RKS_S2_
	.type	_ZN6Matrix6refAddERS_RKS_S2_, @function
_ZN6Matrix6refAddERS_RKS_S2_:
.LFB3823:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdi, %r12
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdx, %rbp
	movl	$16777217, %edx
	movl	$1, %ecx
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rsi, %rbx
	salq	$8, %rdx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	4(%rsi), %eax
	imull	(%rsi), %eax
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L360
	movl	4(%rbx), %edi
	imull	(%rbx), %edi
	movq	8(%r12), %rcx
	movq	8(%rbp), %rdx
	movq	8(%rbx), %rsi
	call	_Z33__device_stub__Z9addKerneliPdS_S_iPdS_S_
.L360:
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3823:
	.size	_ZN6Matrix6refAddERS_RKS_S2_, .-_ZN6Matrix6refAddERS_RKS_S2_
	.section	.rodata.str1.1
.LC18:
	.string	"dimensions must equal for addition"
	.text
	.globl	_ZplRK6MatrixS1_
	.type	_ZplRK6MatrixS1_, @function
_ZplRK6MatrixS1_:
.LFB3824:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$48, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	movq	(%rdx), %rax
	cmpq	%rax, (%rsi)
	je	.L363
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC18(%rip), %rax
	movq	%rax, (%rdi)
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L368
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	call	__cxa_throw@PLT
.L363:
	movq	%rsi, %rbx
	movl	4(%rsi), %esi
	movq	%rdi, %rbp
	leaq	8(%rsp), %rdi
	imull	(%rbx), %esi
	movq	%rdx, %r12
	salq	$3, %rsi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movl	4(%rbx), %eax
	imull	(%rbx), %eax
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$8, %rdx
	movl	$1, %ecx
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 20(%rsp)
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L365
	movl	4(%rbx), %edi
	imull	(%rbx), %edi
	movq	8(%r12), %rdx
	movq	8(%rbx), %rsi
	movq	8(%rsp), %rcx
	call	_Z33__device_stub__Z9addKerneliPdS_S_iPdS_S_
.L365:
	call	cudaDeviceSynchronize@PLT
	movq	(%rbx), %rax
	movq	%rax, 0(%rbp)
	movq	8(%rsp), %rax
	movq	%rax, 8(%rbp)
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L366
.L368:
	call	__stack_chk_fail@PLT
.L366:
	addq	$48, %rsp
	.cfi_def_cfa_offset 32
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3824:
	.size	_ZplRK6MatrixS1_, .-_ZplRK6MatrixS1_
	.globl	_Z33__device_stub__Z9subKerneliPdS_S_iPdS_S_
	.type	_Z33__device_stub__Z9subKerneliPdS_S_iPdS_S_, @function
_Z33__device_stub__Z9subKerneliPdS_S_iPdS_S_:
.LFB3872:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L369
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z9subKerneliPdS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L369:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L371
	call	__stack_chk_fail@PLT
.L371:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3872:
	.size	_Z33__device_stub__Z9subKerneliPdS_S_iPdS_S_, .-_Z33__device_stub__Z9subKerneliPdS_S_iPdS_S_
	.globl	_Z9subKerneliPdS_S_
	.type	_Z9subKerneliPdS_S_, @function
_Z9subKerneliPdS_S_:
.LFB3873:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z9subKerneliPdS_S_iPdS_S_
	.cfi_endproc
.LFE3873:
	.size	_Z9subKerneliPdS_S_, .-_Z9subKerneliPdS_S_
	.align 2
	.globl	_ZN6Matrix6refSubERS_RKS_S2_
	.type	_ZN6Matrix6refSubERS_RKS_S2_, @function
_ZN6Matrix6refSubERS_RKS_S2_:
.LFB3825:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rdi, %r12
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdx, %rbp
	movl	$16777217, %edx
	movl	$1, %ecx
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rsi, %rbx
	salq	$8, %rdx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	movl	4(%rsi), %eax
	imull	(%rsi), %eax
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 8(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 12(%rsp)
	movq	8(%rsp), %rdi
	movl	16(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L375
	movl	4(%rbx), %edi
	imull	(%rbx), %edi
	movq	8(%r12), %rcx
	movq	8(%rbp), %rdx
	movq	8(%rbx), %rsi
	call	_Z33__device_stub__Z9subKerneliPdS_S_iPdS_S_
.L375:
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	cudaDeviceSynchronize@PLT
	.cfi_endproc
.LFE3825:
	.size	_ZN6Matrix6refSubERS_RKS_S2_, .-_ZN6Matrix6refSubERS_RKS_S2_
	.align 2
	.globl	_ZN6Matrix13JacobiOptimalERKS_S1_
	.type	_ZN6Matrix13JacobiOptimalERKS_S1_, @function
_ZN6Matrix13JacobiOptimalERKS_S1_:
.LFB3785:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3785
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	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
	movq	%rsi, %rbx
	subq	$184, %rsp
	.cfi_def_cfa_offset 240
	movl	4(%rsi), %esi
	movl	(%rbx), %edx
	movq	%rdi, 16(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
.LEHB28:
	call	_ZN6MatrixC1Eii
.LEHE28:
	movq	24(%rsp), %rdi
	movq	%rbx, %rsi
.LEHB29:
	call	_ZN6Matrix7copyGPUERS_RKS_
	leaq	72(%rsp), %rax
	movl	4(%rbx), %esi
	movl	(%rbx), %edx
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	_ZN6MatrixC1Eii
.LEHE29:
	movq	24(%rsp), %rsi
	movq	8(%rsp), %rdi
.LEHB30:
	call	_ZN6Matrix29separateDiagonalAndInverseGPUERS_S0_
	movl	(%rbx), %edx
	movq	16(%rsp), %rdi
	movl	$1, %esi
	call	_ZN6Matrix7ZeroGPUEii
.LEHE30:
	leaq	88(%rsp), %rax
	movl	(%rbx), %edx
	movl	$1, %esi
	movq	%rax, %rdi
	movq	%rax, 32(%rsp)
.LEHB31:
	call	_ZN6MatrixC1Eii
.LEHE31:
	movq	8(%rsp), %rsi
	movq	32(%rsp), %rdi
	movq	%r14, %rdx
.LEHB32:
	call	_ZN6Matrix6refMulERS_RKS_S2_
	leaq	104(%rsp), %rax
	movl	4(%rbx), %esi
	movl	(%rbx), %edx
	movq	%rax, %rdi
	movq	%rax, 40(%rsp)
	call	_ZN6MatrixC1Eii
.LEHE32:
	movq	8(%rsp), %rdi
.LEHB33:
	call	_ZN6Matrix25additiveInverseInPlaceGPUERS_
	movq	24(%rsp), %rdx
	movq	8(%rsp), %rsi
	movq	40(%rsp), %rdi
	call	_ZN6Matrix10refMulDiagERS_RKS_S2_
	leaq	120(%rsp), %rbp
	movl	(%rbx), %edx
	movl	$1, %esi
	movq	%rbp, %rdi
	call	_ZN6MatrixC1Eii
.LEHE33:
	leaq	136(%rsp), %r12
	movl	(%rbx), %edx
	movl	$1, %esi
	movq	%r12, %rdi
.LEHB34:
	call	_ZN6MatrixC1Eii
.LEHE34:
	leaq	152(%rsp), %r13
	movl	(%rbx), %edx
	movl	$1, %esi
	movq	%r13, %rdi
.LEHB35:
	call	_ZN6MatrixC1Eii
.LEHE35:
	movl	$9, %r15d
.L381:
	movq	16(%rsp), %rdx
	movq	40(%rsp), %rsi
	movq	%rbp, %rdi
.LEHB36:
	call	_ZN6Matrix6refMulERS_RKS_S2_
	movq	32(%rsp), %rdx
	movq	16(%rsp), %rdi
	movq	%rbp, %rsi
	call	_ZN6Matrix6refAddERS_RKS_S2_
	leal	1(%r15), %eax
	cmpl	$9, %r15d
	jne	.L378
	movq	16(%rsp), %rdx
	movq	%rbx, %rsi
	movq	%r12, %rdi
	call	_ZN6Matrix6refMulERS_RKS_S2_
	movq	%r14, %rdx
	movq	%r12, %rsi
	movq	%r13, %rdi
	call	_ZN6Matrix6refSubERS_RKS_S2_
	movq	%r13, %rdi
	call	_ZN6Matrix19vectorEuclideanNormEv
.LEHE36:
	comisd	.LC19(%rip), %xmm0
	jbe	.L401
	xorl	%eax, %eax
.L378:
	movl	%eax, %r15d
	jmp	.L381
.L401:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	movq	%r12, %rdi
	call	_ZN6MatrixD1Ev
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movq	40(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	32(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	8(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	24(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L391
	jmp	.L402
.L399:
	endbr64
	movq	%rax, %rbx
.L382:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L383
.L398:
	endbr64
	movq	%rax, %rbx
.L383:
	movq	%r12, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L384
.L397:
	endbr64
	movq	%rax, %rbx
.L384:
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L385
.L396:
	endbr64
	movq	%rax, %rbx
.L385:
	movq	40(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L386
.L395:
	endbr64
	movq	%rax, %rbx
.L386:
	movq	32(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L387
.L394:
	endbr64
	movq	%rax, %rbx
.L387:
	movq	16(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L388
.L393:
	endbr64
	movq	%rax, %rbx
.L388:
	movq	8(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L389
.L392:
	endbr64
	movq	%rax, %rbx
.L389:
	movq	24(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L402
	movq	%rbx, %rdi
.LEHB37:
	call	_Unwind_Resume@PLT
.LEHE37:
.L402:
	call	__stack_chk_fail@PLT
.L391:
	movq	16(%rsp), %rax
	addq	$184, %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
.LFE3785:
	.section	.gcc_except_table
.LLSDA3785:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3785-.LLSDACSB3785
.LLSDACSB3785:
	.uleb128 .LEHB28-.LFB3785
	.uleb128 .LEHE28-.LEHB28
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB29-.LFB3785
	.uleb128 .LEHE29-.LEHB29
	.uleb128 .L392-.LFB3785
	.uleb128 0
	.uleb128 .LEHB30-.LFB3785
	.uleb128 .LEHE30-.LEHB30
	.uleb128 .L393-.LFB3785
	.uleb128 0
	.uleb128 .LEHB31-.LFB3785
	.uleb128 .LEHE31-.LEHB31
	.uleb128 .L394-.LFB3785
	.uleb128 0
	.uleb128 .LEHB32-.LFB3785
	.uleb128 .LEHE32-.LEHB32
	.uleb128 .L395-.LFB3785
	.uleb128 0
	.uleb128 .LEHB33-.LFB3785
	.uleb128 .LEHE33-.LEHB33
	.uleb128 .L396-.LFB3785
	.uleb128 0
	.uleb128 .LEHB34-.LFB3785
	.uleb128 .LEHE34-.LEHB34
	.uleb128 .L397-.LFB3785
	.uleb128 0
	.uleb128 .LEHB35-.LFB3785
	.uleb128 .LEHE35-.LEHB35
	.uleb128 .L398-.LFB3785
	.uleb128 0
	.uleb128 .LEHB36-.LFB3785
	.uleb128 .LEHE36-.LEHB36
	.uleb128 .L399-.LFB3785
	.uleb128 0
	.uleb128 .LEHB37-.LFB3785
	.uleb128 .LEHE37-.LEHB37
	.uleb128 0
	.uleb128 0
.LLSDACSE3785:
	.text
	.size	_ZN6Matrix13JacobiOptimalERKS_S1_, .-_ZN6Matrix13JacobiOptimalERKS_S1_
	.align 2
	.globl	_ZN6Matrix18GaussSeidelOptimalERKS_S1_
	.type	_ZN6Matrix18GaussSeidelOptimalERKS_S1_, @function
_ZN6Matrix18GaussSeidelOptimalERKS_S1_:
.LFB3789:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3789
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	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
	movq	%rsi, %rbx
	subq	$184, %rsp
	.cfi_def_cfa_offset 240
	movl	4(%rsi), %esi
	movl	(%rbx), %edx
	movq	%rdi, 16(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
.LEHB38:
	call	_ZN6MatrixC1Eii
.LEHE38:
	leaq	72(%rsp), %rax
	movl	4(%rbx), %esi
	movl	(%rbx), %edx
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
.LEHB39:
	call	_ZN6MatrixC1Eii
.LEHE39:
	movq	8(%rsp), %rdi
	movq	%rbx, %rsi
.LEHB40:
	call	_ZN6Matrix7copyGPUERS_RKS_
	movq	8(%rsp), %rsi
	movq	24(%rsp), %rdi
	call	_ZN6Matrix16separateUpperGPUERS_S0_
	movl	(%rbx), %edx
	movq	16(%rsp), %rdi
	movl	$1, %esi
	call	_ZN6Matrix7ZeroCPUEii
.LEHE40:
	leaq	88(%rsp), %rax
	movq	8(%rsp), %rsi
	movq	%r14, %rdx
	movq	%rax, %rdi
	movq	%rax, 40(%rsp)
.LEHB41:
	call	_ZN6Matrix19ForwardSubstitutionERKS_S1_
.LEHE41:
	leaq	104(%rsp), %rax
	movl	(%rbx), %edx
	movl	$1, %esi
	movq	%rax, %rdi
	movq	%rax, 32(%rsp)
.LEHB42:
	call	_ZN6MatrixC1Eii
.LEHE42:
	leaq	120(%rsp), %rbp
	movl	(%rbx), %edx
	movl	$1, %esi
	movq	%rbp, %rdi
.LEHB43:
	call	_ZN6MatrixC1Eii
.LEHE43:
	leaq	136(%rsp), %r12
	movl	(%rbx), %edx
	movl	$1, %esi
	movq	%r12, %rdi
.LEHB44:
	call	_ZN6MatrixC1Eii
.LEHE44:
	leaq	152(%rsp), %r13
	movl	(%rbx), %edx
	movl	$1, %esi
	movq	%r13, %rdi
.LEHB45:
	call	_ZN6MatrixC1Eii
.LEHE45:
	movl	$9, %r15d
.L408:
	movq	16(%rsp), %rdx
	movq	24(%rsp), %rsi
	movq	32(%rsp), %rdi
.LEHB46:
	call	_ZN6Matrix6refMulERS_RKS_S2_
	movq	32(%rsp), %rdx
	movq	8(%rsp), %rsi
	movq	%rbp, %rdi
	call	_ZN6Matrix22forwardSubstitutionGPUERS_RKS_S2_
	movq	%rbp, %rdi
	call	_ZN6Matrix25additiveInverseInPlaceGPUERS_
	movq	40(%rsp), %rdx
	movq	16(%rsp), %rdi
	movq	%rbp, %rsi
	call	_ZN6Matrix6refAddERS_RKS_S2_
	leal	1(%r15), %eax
	cmpl	$9, %r15d
	jne	.L405
	movq	16(%rsp), %rdx
	movq	%rbx, %rsi
	movq	%r12, %rdi
	call	_ZN6Matrix6refMulERS_RKS_S2_
	movq	%r14, %rdx
	movq	%r12, %rsi
	movq	%r13, %rdi
	call	_ZN6Matrix6refSubERS_RKS_S2_
	movq	%r13, %rdi
	call	_ZN6Matrix19vectorEuclideanNormEv
.LEHE46:
	comisd	.LC19(%rip), %xmm0
	jbe	.L428
	xorl	%eax, %eax
.L405:
	movl	%eax, %r15d
	jmp	.L408
.L428:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	movq	%r12, %rdi
	call	_ZN6MatrixD1Ev
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movq	32(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	40(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	24(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	8(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L418
	jmp	.L429
.L426:
	endbr64
	movq	%rax, %rbx
.L409:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L410
.L425:
	endbr64
	movq	%rax, %rbx
.L410:
	movq	%r12, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L411
.L424:
	endbr64
	movq	%rax, %rbx
.L411:
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L412
.L423:
	endbr64
	movq	%rax, %rbx
.L412:
	movq	32(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L413
.L422:
	endbr64
	movq	%rax, %rbx
.L413:
	movq	40(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L414
.L421:
	endbr64
	movq	%rax, %rbx
.L414:
	movq	16(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L415
.L420:
	endbr64
	movq	%rax, %rbx
.L415:
	movq	24(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L416
.L419:
	endbr64
	movq	%rax, %rbx
.L416:
	movq	8(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L429
	movq	%rbx, %rdi
.LEHB47:
	call	_Unwind_Resume@PLT
.LEHE47:
.L429:
	call	__stack_chk_fail@PLT
.L418:
	movq	16(%rsp), %rax
	addq	$184, %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
.LFE3789:
	.section	.gcc_except_table
.LLSDA3789:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3789-.LLSDACSB3789
.LLSDACSB3789:
	.uleb128 .LEHB38-.LFB3789
	.uleb128 .LEHE38-.LEHB38
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB39-.LFB3789
	.uleb128 .LEHE39-.LEHB39
	.uleb128 .L419-.LFB3789
	.uleb128 0
	.uleb128 .LEHB40-.LFB3789
	.uleb128 .LEHE40-.LEHB40
	.uleb128 .L420-.LFB3789
	.uleb128 0
	.uleb128 .LEHB41-.LFB3789
	.uleb128 .LEHE41-.LEHB41
	.uleb128 .L421-.LFB3789
	.uleb128 0
	.uleb128 .LEHB42-.LFB3789
	.uleb128 .LEHE42-.LEHB42
	.uleb128 .L422-.LFB3789
	.uleb128 0
	.uleb128 .LEHB43-.LFB3789
	.uleb128 .LEHE43-.LEHB43
	.uleb128 .L423-.LFB3789
	.uleb128 0
	.uleb128 .LEHB44-.LFB3789
	.uleb128 .LEHE44-.LEHB44
	.uleb128 .L424-.LFB3789
	.uleb128 0
	.uleb128 .LEHB45-.LFB3789
	.uleb128 .LEHE45-.LEHB45
	.uleb128 .L425-.LFB3789
	.uleb128 0
	.uleb128 .LEHB46-.LFB3789
	.uleb128 .LEHE46-.LEHB46
	.uleb128 .L426-.LFB3789
	.uleb128 0
	.uleb128 .LEHB47-.LFB3789
	.uleb128 .LEHE47-.LEHB47
	.uleb128 0
	.uleb128 0
.LLSDACSE3789:
	.text
	.size	_ZN6Matrix18GaussSeidelOptimalERKS_S1_, .-_ZN6Matrix18GaussSeidelOptimalERKS_S1_
	.globl	_ZmiRK6MatrixS1_
	.type	_ZmiRK6MatrixS1_, @function
_ZmiRK6MatrixS1_:
.LFB3826:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$48, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	movq	(%rdx), %rax
	cmpq	%rax, (%rsi)
	je	.L432
	movl	$8, %edi
	call	__cxa_allocate_exception@PLT
	movq	%rax, %rdi
	leaq	.LC18(%rip), %rax
	movq	%rax, (%rdi)
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L437
	xorl	%edx, %edx
	leaq	_ZTIPKc(%rip), %rsi
	call	__cxa_throw@PLT
.L432:
	movq	%rsi, %rbx
	movl	4(%rsi), %esi
	movq	%rdi, %rbp
	leaq	8(%rsp), %rdi
	imull	(%rbx), %esi
	movq	%rdx, %r12
	salq	$3, %rsi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movl	4(%rbx), %eax
	imull	(%rbx), %eax
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$8, %rdx
	movl	$1, %ecx
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 20(%rsp)
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L434
	movl	4(%rbx), %edi
	imull	(%rbx), %edi
	movq	8(%r12), %rdx
	movq	8(%rbx), %rsi
	movq	8(%rsp), %rcx
	call	_Z33__device_stub__Z9subKerneliPdS_S_iPdS_S_
.L434:
	call	cudaDeviceSynchronize@PLT
	movq	(%rbx), %rax
	movq	%rax, 0(%rbp)
	movq	8(%rsp), %rax
	movq	%rax, 8(%rbp)
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L435
.L437:
	call	__stack_chk_fail@PLT
.L435:
	addq	$48, %rsp
	.cfi_def_cfa_offset 32
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3826:
	.size	_ZmiRK6MatrixS1_, .-_ZmiRK6MatrixS1_
	.globl	_Z44__device_stub__Z21additiveInverseKerneliPdS_iPdS_
	.type	_Z44__device_stub__Z21additiveInverseKerneliPdS_iPdS_, @function
_Z44__device_stub__Z21additiveInverseKerneliPdS_iPdS_:
.LFB3874:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L438
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z21additiveInverseKerneliPdS_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L438:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L440
	call	__stack_chk_fail@PLT
.L440:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3874:
	.size	_Z44__device_stub__Z21additiveInverseKerneliPdS_iPdS_, .-_Z44__device_stub__Z21additiveInverseKerneliPdS_iPdS_
	.globl	_Z21additiveInverseKerneliPdS_
	.type	_Z21additiveInverseKerneliPdS_, @function
_Z21additiveInverseKerneliPdS_:
.LFB3875:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z21additiveInverseKerneliPdS_iPdS_
	.cfi_endproc
.LFE3875:
	.size	_Z21additiveInverseKerneliPdS_, .-_Z21additiveInverseKerneliPdS_
	.globl	_ZngRK6Matrix
	.type	_ZngRK6Matrix, @function
_ZngRK6Matrix:
.LFB3827:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rsi, %rbx
	subq	$56, %rsp
	.cfi_def_cfa_offset 80
	movl	4(%rsi), %esi
	imull	(%rbx), %esi
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	salq	$3, %rsi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movl	4(%rbx), %eax
	imull	(%rbx), %eax
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$8, %rdx
	movl	$1, %ecx
	addl	$255, %eax
	shrl	$8, %eax
	movl	%eax, 16(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 20(%rsp)
	movq	16(%rsp), %rdi
	movl	24(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L444
	movl	4(%rbx), %edi
	imull	(%rbx), %edi
	movq	8(%rbx), %rdx
	movq	8(%rsp), %rsi
	call	_Z44__device_stub__Z21additiveInverseKerneliPdS_iPdS_
.L444:
	call	cudaDeviceSynchronize@PLT
	movq	(%rbx), %rax
	movq	%rax, 0(%rbp)
	movq	8(%rsp), %rax
	movq	%rax, 8(%rbp)
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L445
	call	__stack_chk_fail@PLT
.L445:
	addq	$56, %rsp
	.cfi_def_cfa_offset 24
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3827:
	.size	_ZngRK6Matrix, .-_ZngRK6Matrix
	.section	.rodata.str1.1
.LC20:
	.string	"res: %d \n"
	.text
	.align 2
	.globl	_ZN6Matrix6JacobiERKS_S1_
	.type	_ZN6Matrix6JacobiERKS_S1_, @function
_ZN6Matrix6JacobiERKS_S1_:
.LFB3784:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3784
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$168, %rsp
	.cfi_def_cfa_offset 224
	movq	%rdi, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
.LEHB48:
	call	_ZN6MatrixC1ERKS_
.LEHE48:
	leaq	72(%rsp), %rax
	movq	24(%rsp), %rsi
	movq	%rax, %rdi
	movq	%rax, 16(%rsp)
.LEHB49:
	call	_ZN6Matrix16separateDiagonalEv
.LEHE49:
	movl	(%r12), %edx
	movq	8(%rsp), %rdi
	movl	$1, %esi
.LEHB50:
	call	_ZN6Matrix7ZeroCPUEii
.LEHE50:
	movq	16(%rsp), %rdi
.LEHB51:
	call	_ZN6Matrix25inverseDiagonalInPlaceCPUEv
	leaq	136(%rsp), %rbx
	movq	16(%rsp), %rsi
	movq	%rbx, %rdi
	call	_ZngRK6Matrix
.LEHE51:
	leaq	88(%rsp), %rax
	movq	24(%rsp), %rdx
	movq	%rbx, %rsi
	movq	%rax, %rdi
	movq	%rax, 32(%rsp)
.LEHB52:
	call	_ZmlRK6MatrixS1_
.LEHE52:
	movq	%rbx, %rdi
	leaq	104(%rsp), %r13
	call	_ZN6MatrixD1Ev
	movq	16(%rsp), %rsi
	movq	%r14, %rdx
	movq	%r13, %rdi
.LEHB53:
	call	_ZmlRK6MatrixS1_
.LEHE53:
	xorl	%r15d, %r15d
	leaq	120(%rsp), %rbp
.L448:
	movq	8(%rsp), %rdx
	movq	32(%rsp), %rsi
	movq	%rbp, %rdi
.LEHB54:
	call	_ZmlRK6MatrixS1_
.LEHE54:
	movq	%r13, %rdx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
.LEHB55:
	call	_ZplRK6MatrixS1_
.LEHE55:
	movq	8(%rsp), %rdi
	movq	%rbx, %rsi
.LEHB56:
	call	_ZN6MatrixaSERKS_
.LEHE56:
	movq	%rbx, %rdi
	call	_ZN6MatrixD1Ev
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movq	8(%rsp), %rdx
	movq	%r12, %rsi
	movq	%rbp, %rdi
.LEHB57:
	call	_ZmlRK6MatrixS1_
.LEHE57:
	movq	%r14, %rdx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
.LEHB58:
	call	_ZmiRK6MatrixS1_
.LEHE58:
	movq	%rbx, %rdi
.LEHB59:
	call	_ZN6Matrix19vectorEuclideanNormEv
.LEHE59:
	movq	%rbx, %rdi
	movsd	%xmm0, 40(%rsp)
	incl	%r15d
	call	_ZN6MatrixD1Ev
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movsd	40(%rsp), %xmm0
	comisd	.LC19(%rip), %xmm0
	ja	.L448
	movl	%r15d, %edx
	leaq	.LC20(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
.LEHB60:
	call	__printf_chk@PLT
.LEHE60:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	movq	32(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	16(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	24(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L460
	jmp	.L472
.L464:
	endbr64
	movq	%rax, %r12
.L449:
	movq	%rbx, %rdi
	jmp	.L474
.L467:
	endbr64
	jmp	.L476
.L466:
	endbr64
	jmp	.L475
.L470:
	endbr64
.L476:
	movq	%rax, %r12
.L454:
	movq	%rbx, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L455
.L469:
	endbr64
.L475:
	movq	%rax, %r12
.L455:
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L453
.L468:
	endbr64
	movq	%rax, %r12
.L453:
	movq	%r13, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L456
.L465:
	endbr64
	movq	%rax, %r12
.L456:
	movq	32(%rsp), %rdi
.L474:
	call	_ZN6MatrixD1Ev
	jmp	.L450
.L463:
	endbr64
	movq	%rax, %r12
.L450:
	movq	8(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L457
.L462:
	endbr64
	movq	%rax, %r12
.L457:
	movq	16(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L458
.L461:
	endbr64
	movq	%rax, %r12
.L458:
	movq	24(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L472
	movq	%r12, %rdi
.LEHB61:
	call	_Unwind_Resume@PLT
.LEHE61:
.L472:
	call	__stack_chk_fail@PLT
.L460:
	movq	8(%rsp), %rax
	addq	$168, %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
.LFE3784:
	.section	.gcc_except_table
.LLSDA3784:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3784-.LLSDACSB3784
.LLSDACSB3784:
	.uleb128 .LEHB48-.LFB3784
	.uleb128 .LEHE48-.LEHB48
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB49-.LFB3784
	.uleb128 .LEHE49-.LEHB49
	.uleb128 .L461-.LFB3784
	.uleb128 0
	.uleb128 .LEHB50-.LFB3784
	.uleb128 .LEHE50-.LEHB50
	.uleb128 .L462-.LFB3784
	.uleb128 0
	.uleb128 .LEHB51-.LFB3784
	.uleb128 .LEHE51-.LEHB51
	.uleb128 .L463-.LFB3784
	.uleb128 0
	.uleb128 .LEHB52-.LFB3784
	.uleb128 .LEHE52-.LEHB52
	.uleb128 .L464-.LFB3784
	.uleb128 0
	.uleb128 .LEHB53-.LFB3784
	.uleb128 .LEHE53-.LEHB53
	.uleb128 .L465-.LFB3784
	.uleb128 0
	.uleb128 .LEHB54-.LFB3784
	.uleb128 .LEHE54-.LEHB54
	.uleb128 .L468-.LFB3784
	.uleb128 0
	.uleb128 .LEHB55-.LFB3784
	.uleb128 .LEHE55-.LEHB55
	.uleb128 .L466-.LFB3784
	.uleb128 0
	.uleb128 .LEHB56-.LFB3784
	.uleb128 .LEHE56-.LEHB56
	.uleb128 .L467-.LFB3784
	.uleb128 0
	.uleb128 .LEHB57-.LFB3784
	.uleb128 .LEHE57-.LEHB57
	.uleb128 .L468-.LFB3784
	.uleb128 0
	.uleb128 .LEHB58-.LFB3784
	.uleb128 .LEHE58-.LEHB58
	.uleb128 .L469-.LFB3784
	.uleb128 0
	.uleb128 .LEHB59-.LFB3784
	.uleb128 .LEHE59-.LEHB59
	.uleb128 .L470-.LFB3784
	.uleb128 0
	.uleb128 .LEHB60-.LFB3784
	.uleb128 .LEHE60-.LEHB60
	.uleb128 .L468-.LFB3784
	.uleb128 0
	.uleb128 .LEHB61-.LFB3784
	.uleb128 .LEHE61-.LEHB61
	.uleb128 0
	.uleb128 0
.LLSDACSE3784:
	.text
	.size	_ZN6Matrix6JacobiERKS_S1_, .-_ZN6Matrix6JacobiERKS_S1_
	.align 2
	.globl	_ZN6Matrix11GaussSeidelERKS_S1_
	.type	_ZN6Matrix11GaussSeidelERKS_S1_, @function
_ZN6Matrix11GaussSeidelERKS_S1_:
.LFB3788:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3788
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rsi, %r13
	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
	subq	$184, %rsp
	.cfi_def_cfa_offset 240
	movq	%rdi, 32(%rsp)
	leaq	136(%rsp), %rbp
	leaq	120(%rsp), %rbx
	movq	%rbp, %rdi
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
.LEHB62:
	call	_ZNK6Matrix11diagonalCPUEv
.LEHE62:
	movq	%r13, %rsi
	movq	%rbx, %rdi
.LEHB63:
	call	_ZNK6Matrix8lowerCPUEv
.LEHE63:
	leaq	152(%rsp), %r12
	movq	%rbp, %rdx
	movq	%rbx, %rsi
	movq	%r12, %rdi
.LEHB64:
	call	_ZplRK6MatrixS1_
.LEHE64:
	leaq	56(%rsp), %rax
	movq	%r12, %rsi
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
.LEHB65:
	call	_ZngRK6Matrix
.LEHE65:
	movq	%r12, %rdi
	call	_ZN6MatrixD1Ev
	movq	%rbx, %rdi
	call	_ZN6MatrixD1Ev
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	leaq	72(%rsp), %rax
	movq	%r13, %rsi
	movq	%rax, %rdi
	movq	%rax, 16(%rsp)
.LEHB66:
	call	_ZNK6Matrix8upperCPUEv
.LEHE66:
	leaq	88(%rsp), %rax
	movl	0(%r13), %edx
	movl	$1, %esi
	movq	%rax, %rdi
	movq	%rax, (%rsp)
.LEHB67:
	call	_ZN6Matrix7ZeroCPUEii
.LEHE67:
	leaq	104(%rsp), %rax
	movq	8(%rsp), %rsi
	movq	%r14, %rdx
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
.LEHB68:
	call	_ZN6Matrix19ForwardSubstitutionERKS_S1_
.LEHE68:
	xorl	%r15d, %r15d
.L478:
	movq	(%rsp), %rdx
	movq	16(%rsp), %rsi
	movq	%rbx, %rdi
.LEHB69:
	call	_ZmlRK6MatrixS1_
.LEHE69:
	movq	8(%rsp), %rsi
	movq	%rbx, %rdx
	movq	%rbp, %rdi
.LEHB70:
	call	_ZN6Matrix19ForwardSubstitutionERKS_S1_
.LEHE70:
	movq	24(%rsp), %rdx
	movq	%rbp, %rsi
	movq	%r12, %rdi
.LEHB71:
	call	_ZplRK6MatrixS1_
.LEHE71:
	movq	(%rsp), %rdi
	movq	%r12, %rsi
.LEHB72:
	call	_ZN6MatrixaSERKS_
.LEHE72:
	movq	%r12, %rdi
	incl	%r15d
	call	_ZN6MatrixD1Ev
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movq	%rbx, %rdi
	call	_ZN6MatrixD1Ev
	movq	(%rsp), %rsi
	movq	%rbx, %rdi
.LEHB73:
	call	_ZngRK6Matrix
.LEHE73:
	movq	%rbx, %rdx
	movq	%r13, %rsi
	movq	%rbp, %rdi
.LEHB74:
	call	_ZmlRK6MatrixS1_
.LEHE74:
	movq	%r14, %rdx
	movq	%rbp, %rsi
	movq	%r12, %rdi
.LEHB75:
	call	_ZmiRK6MatrixS1_
.LEHE75:
	movq	%r12, %rdi
.LEHB76:
	call	_ZN6Matrix19vectorEuclideanNormEv
.LEHE76:
	movq	%r12, %rdi
	movsd	%xmm0, 40(%rsp)
	call	_ZN6MatrixD1Ev
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movq	%rbx, %rdi
	call	_ZN6MatrixD1Ev
	movsd	40(%rsp), %xmm0
	comisd	.LC19(%rip), %xmm0
	ja	.L478
	movl	%r15d, %edx
	leaq	.LC20(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
.LEHB77:
	call	__printf_chk@PLT
	movq	(%rsp), %rsi
	movq	32(%rsp), %rdi
	call	_ZngRK6Matrix
.LEHE77:
	movq	24(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	16(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	8(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L494
	jmp	.L509
.L497:
	endbr64
	movq	%rax, %r13
.L479:
	movq	%r12, %rdi
	movq	%r13, %r12
	call	_ZN6MatrixD1Ev
	jmp	.L480
.L496:
	endbr64
	movq	%rax, %r12
.L480:
	movq	%rbx, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L481
.L495:
	endbr64
	movq	%rax, %r12
.L481:
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	movq	%r12, %rdi
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L512
	jmp	.L509
.L503:
	endbr64
	jmp	.L515
.L502:
	endbr64
	jmp	.L514
.L501:
	endbr64
	jmp	.L513
.L507:
	endbr64
.L515:
	movq	%rax, %r13
.L487:
	movq	%r12, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L488
.L506:
	endbr64
.L514:
	movq	%rax, %r13
.L488:
	movq	%rbp, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L489
.L505:
	endbr64
.L513:
	movq	%rax, %r13
.L489:
	movq	%rbx, %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L486
.L504:
	endbr64
	movq	%rax, %r13
.L486:
	movq	24(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L490
.L500:
	endbr64
	movq	%rax, %r13
.L490:
	movq	(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L491
.L499:
	endbr64
	movq	%rax, %r13
.L491:
	movq	16(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	jmp	.L492
.L498:
	endbr64
	movq	%rax, %r13
.L492:
	movq	8(%rsp), %rdi
	call	_ZN6MatrixD1Ev
	movq	%r13, %rdi
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L509
.L512:
.LEHB78:
	call	_Unwind_Resume@PLT
.LEHE78:
.L509:
	call	__stack_chk_fail@PLT
.L494:
	movq	32(%rsp), %rax
	addq	$184, %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
.LFE3788:
	.section	.gcc_except_table
.LLSDA3788:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3788-.LLSDACSB3788
.LLSDACSB3788:
	.uleb128 .LEHB62-.LFB3788
	.uleb128 .LEHE62-.LEHB62
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB63-.LFB3788
	.uleb128 .LEHE63-.LEHB63
	.uleb128 .L495-.LFB3788
	.uleb128 0
	.uleb128 .LEHB64-.LFB3788
	.uleb128 .LEHE64-.LEHB64
	.uleb128 .L496-.LFB3788
	.uleb128 0
	.uleb128 .LEHB65-.LFB3788
	.uleb128 .LEHE65-.LEHB65
	.uleb128 .L497-.LFB3788
	.uleb128 0
	.uleb128 .LEHB66-.LFB3788
	.uleb128 .LEHE66-.LEHB66
	.uleb128 .L498-.LFB3788
	.uleb128 0
	.uleb128 .LEHB67-.LFB3788
	.uleb128 .LEHE67-.LEHB67
	.uleb128 .L499-.LFB3788
	.uleb128 0
	.uleb128 .LEHB68-.LFB3788
	.uleb128 .LEHE68-.LEHB68
	.uleb128 .L500-.LFB3788
	.uleb128 0
	.uleb128 .LEHB69-.LFB3788
	.uleb128 .LEHE69-.LEHB69
	.uleb128 .L504-.LFB3788
	.uleb128 0
	.uleb128 .LEHB70-.LFB3788
	.uleb128 .LEHE70-.LEHB70
	.uleb128 .L501-.LFB3788
	.uleb128 0
	.uleb128 .LEHB71-.LFB3788
	.uleb128 .LEHE71-.LEHB71
	.uleb128 .L502-.LFB3788
	.uleb128 0
	.uleb128 .LEHB72-.LFB3788
	.uleb128 .LEHE72-.LEHB72
	.uleb128 .L503-.LFB3788
	.uleb128 0
	.uleb128 .LEHB73-.LFB3788
	.uleb128 .LEHE73-.LEHB73
	.uleb128 .L504-.LFB3788
	.uleb128 0
	.uleb128 .LEHB74-.LFB3788
	.uleb128 .LEHE74-.LEHB74
	.uleb128 .L505-.LFB3788
	.uleb128 0
	.uleb128 .LEHB75-.LFB3788
	.uleb128 .LEHE75-.LEHB75
	.uleb128 .L506-.LFB3788
	.uleb128 0
	.uleb128 .LEHB76-.LFB3788
	.uleb128 .LEHE76-.LEHB76
	.uleb128 .L507-.LFB3788
	.uleb128 0
	.uleb128 .LEHB77-.LFB3788
	.uleb128 .LEHE77-.LEHB77
	.uleb128 .L504-.LFB3788
	.uleb128 0
	.uleb128 .LEHB78-.LFB3788
	.uleb128 .LEHE78-.LEHB78
	.uleb128 0
	.uleb128 0
.LLSDACSE3788:
	.text
	.size	_ZN6Matrix11GaussSeidelERKS_S1_, .-_ZN6Matrix11GaussSeidelERKS_S1_
	.section	.rodata.str1.1
.LC21:
	.string	"_Z21additiveInverseKerneliPdS_"
.LC22:
	.string	"_Z9subKerneliPdS_S_"
.LC23:
	.string	"_Z9addKerneliPdS_S_"
.LC24:
	.string	"_Z13mulDiagKerneliiiPdS_S_"
.LC25:
	.string	"_Z9mulKerneliiiPdS_S_"
.LC26:
	.string	"_Z25forwardSubstitutionKerneliPdS_S_"
.LC27:
	.string	"_Z28additiveInverseInPlaceKerneliPd"
.LC28:
	.string	"_Z19separateUpperKerneliiPdS_"
.LC29:
	.string	"_Z22separateDiagonalKerneliPdS_"
.LC30:
	.string	"_Z10copyKerneliPdS_"
.LC31:
	.string	"_Z14doolitleKerneliPdS_"
.LC32:
	.string	"_Z13ZeroGPUKerneliPd"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3877:
	.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	.LC21(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z21additiveInverseKerneliPdS_(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC22(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z9subKerneliPdS_S_(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC23(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z9addKerneliPdS_S_(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC24(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z13mulDiagKerneliiiPdS_S_(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC25(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z9mulKerneliiiPdS_S_(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC26(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z25forwardSubstitutionKerneliPdS_S_(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC27(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z28additiveInverseInPlaceKerneliPd(%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
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC28(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z19separateUpperKerneliiPdS_(%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
	leaq	.LC29(%rip), %rdx
	pushq	$0
	.cfi_def_cfa_offset 24
	movq	%rdx, %rcx
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	_Z22separateDiagonalKerneliPdS_(%rip), %rsi
	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	.LC30(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z10copyKerneliPdS_(%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	.LC31(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z14doolitleKerneliPdS_(%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	.LC32(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z13ZeroGPUKerneliPd(%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
.LFE3877:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.text
	.align 2
	.globl	_ZN6Matrix8FromFileENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
	.type	_ZN6Matrix8FromFileENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE, @function
_ZN6Matrix8FromFileENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE:
.LFB3783:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3783
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	subq	$560, %rsp
	.cfi_def_cfa_offset 592
	movq	%fs:40, %rax
	movq	%rax, 552(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rbx
	movq	%rbx, %rdi
.LEHB79:
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEEC1Ev@PLT
.LEHE79:
	movq	(%r12), %rsi
	movl	$8, %edx
	movq	%rbx, %rdi
.LEHB80:
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0
	leaq	40(%rsp), %rdi
	xorl	%esi, %esi
	xorl	%edx, %edx
	call	_ZNSo5seekpESt4fposI11__mbstate_tE@PLT
	leaq	8(%rsp), %rsi
	movq	%rbx, %rdi
	call	_ZNSirsERi@PLT
	leaq	12(%rsp), %rsi
	movq	%rbx, %rdi
	call	_ZNSirsERi@PLT
	movl	8(%rsp), %esi
	imull	12(%rsp), %esi
	leaq	16(%rsp), %rdi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	xorl	%r12d, %r12d
.L519:
	movl	8(%rsp), %eax
	imull	12(%rsp), %eax
	cmpl	%r12d, %eax
	jle	.L529
	movq	16(%rsp), %rax
	movq	%rbx, %rdi
	leaq	(%rax,%r12,8), %rsi
	call	_ZNSi10_M_extractIdEERSiRT_@PLT
	incq	%r12
	jmp	.L519
.L529:
	movq	%rbx, %rdi
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEE5closeEv@PLT
.LEHE80:
	movl	12(%rsp), %eax
	movq	%rbx, %rdi
	movl	%eax, 0(%rbp)
	movl	8(%rsp), %eax
	movl	%eax, 4(%rbp)
	movq	16(%rsp), %rax
	movq	%rax, 8(%rbp)
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	je	.L523
	jmp	.L527
.L525:
	endbr64
	movq	%rax, %rbp
.L521:
	movq	%rbx, %rdi
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	552(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L527
	movq	%rbp, %rdi
.LEHB81:
	call	_Unwind_Resume@PLT
.LEHE81:
.L527:
	call	__stack_chk_fail@PLT
.L523:
	addq	$560, %rsp
	.cfi_def_cfa_offset 32
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3783:
	.section	.gcc_except_table
.LLSDA3783:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3783-.LLSDACSB3783
.LLSDACSB3783:
	.uleb128 .LEHB79-.LFB3783
	.uleb128 .LEHE79-.LEHB79
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB80-.LFB3783
	.uleb128 .LEHE80-.LEHB80
	.uleb128 .L525-.LFB3783
	.uleb128 0
	.uleb128 .LEHB81-.LFB3783
	.uleb128 .LEHE81-.LEHB81
	.uleb128 0
	.uleb128 0
.LLSDACSE3783:
	.text
	.size	_ZN6Matrix8FromFileENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE, .-_ZN6Matrix8FromFileENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
	.align 2
	.globl	_ZN6Matrix6toFileENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
	.type	_ZN6Matrix6toFileENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE, @function
_ZN6Matrix6toFileENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE:
.LFB3807:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3807
	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
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdi, %rbx
	subq	$552, %rsp
	.cfi_def_cfa_offset 608
	movq	%fs:40, %rax
	movq	%rax, 536(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rbp
	movq	%rbp, %rdi
.LEHB82:
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEEC1Ev@PLT
.LEHE82:
	movq	(%r12), %rsi
	movl	$16, %edx
	movq	%rbp, %rdi
.LEHB83:
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEE4openERKNSt7__cxx1112basic_stringIcS1_SaIcEEESt13_Ios_Openmode.isra.0
	xorl	%esi, %esi
	xorl	%edx, %edx
	movq	%rbp, %rdi
	call	_ZNSi5seekgESt4fposI11__mbstate_tE@PLT
	leaq	24(%rsp), %r14
	movl	4(%rbx), %esi
	movq	%r14, %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	%rax, %rdi
	movl	$32, %esi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	movl	(%rbx), %esi
	movq	%rax, %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	%rax, %rdi
	movl	$10, %esi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	xorl	%r12d, %r12d
	leaq	.LC12(%rip), %r15
.L531:
	cmpl	(%rbx), %r12d
	jnb	.L543
	xorl	%r13d, %r13d
.L533:
	movl	4(%rbx), %eax
	cmpl	%eax, %r13d
	jnb	.L544
	imull	%r12d, %eax
	movq	%r14, %rdi
	leal	(%rax,%r13), %edx
	movq	8(%rbx), %rax
	movsd	(%rax,%rdx,8), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movl	$32, %esi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	incl	%r13d
	jmp	.L533
.L544:
	movl	$1, %edx
	movq	%r15, %rsi
	movq	%r14, %rdi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
	incl	%r12d
	jmp	.L531
.L543:
	movq	%rbp, %rdi
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEE5closeEv@PLT
.LEHE83:
	movq	%rbp, %rdi
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	536(%rsp), %rax
	subq	%fs:40, %rax
	je	.L536
	jmp	.L541
.L539:
	endbr64
	movq	%rax, %rbx
.L534:
	movq	%rbp, %rdi
	call	_ZNSt13basic_fstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	536(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L541
	movq	%rbx, %rdi
.LEHB84:
	call	_Unwind_Resume@PLT
.LEHE84:
.L541:
	call	__stack_chk_fail@PLT
.L536:
	addq	$552, %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
.LFE3807:
	.section	.gcc_except_table
.LLSDA3807:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3807-.LLSDACSB3807
.LLSDACSB3807:
	.uleb128 .LEHB82-.LFB3807
	.uleb128 .LEHE82-.LEHB82
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB83-.LFB3807
	.uleb128 .LEHE83-.LEHB83
	.uleb128 .L539-.LFB3807
	.uleb128 0
	.uleb128 .LEHB84-.LFB3807
	.uleb128 .LEHE84-.LEHB84
	.uleb128 0
	.uleb128 0
.LLSDACSE3807:
	.text
	.size	_ZN6Matrix6toFileENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE, .-_ZN6Matrix6toFileENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
	.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	1079574528
	.align 8
.LC2:
	.long	0
	.long	1072693248
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC6:
	.long	-1
	.long	2147483647
	.long	0
	.long	0
	.section	.rodata.cst8
	.align 8
.LC13:
	.long	0
	.long	1076363264
	.align 8
.LC14:
	.long	0
	.long	-1074790400
	.align 8
.LC15:
	.long	0
	.long	1074266112
	.align 8
.LC16:
	.long	0
	.long	1073741824
	.align 8
.LC19:
	.long	-1598689907
	.long	1051772663
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.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

#define BLOCK_SIZE 256

#include "hip/hip_runtime.h"

#include <stdio.h>
#include <string>

class Matrix
{
private:
	unsigned int rows;
	unsigned int cols;
	double* mat;
public:
	//getters
	unsigned int getRows() const;
	unsigned int getCols() const;

	//constructors
	Matrix(int cols, int rows, double* mat);
	Matrix(int cols, int rows);
	Matrix(const Matrix& a);
	void operator=(const Matrix& a);
	static Matrix Stub();
	static Matrix ZeroCPU(int cols, int rows);
	static Matrix OneCPU(int cols, int rows);
	static Matrix ZeroGPU(int cols, int rows);
	static Matrix IdentityCPU(int cols, int rows);
	static Matrix FromFile(std::string path);
	static Matrix Jacobi(const Matrix& A, const Matrix& b);
	static Matrix JacobiOptimal(const Matrix& A, const Matrix& b);

	static Matrix ForwardSubstitution(const Matrix& A, const Matrix& b);
	static Matrix BackwardSubstitution(const Matrix& A, const Matrix& b);
	static Matrix GaussSeidel(const Matrix& A, const Matrix& b);
	static Matrix GaussSeidelOptimal(const Matrix& A, const Matrix& b);
	static Matrix LUMehtod(const Matrix& A, const Matrix& b);
	static Matrix LUMehtodOptimal(const Matrix& A, const Matrix& b);


	//nowy pomysl
	static void doolitle(Matrix& L, Matrix& U, const Matrix& A);
	static void doolitleGPU(Matrix& L, Matrix& U, const Matrix& A);
	static void createTest(Matrix& A, Matrix& b, Matrix& x, int size);
	static void createTask(Matrix& A, Matrix& b, const int size);
	static void createTaskC(Matrix& A, Matrix& b);

	static double maxError(Matrix& x, Matrix& r);
	static void copyGPU(Matrix& a, const Matrix& b);
	static void separateDiagonalAndInverseGPU(Matrix& d, Matrix& A);
	static void separateUpperGPU(Matrix& U, Matrix& A);
	static void additiveInverseInPlaceGPU(Matrix& A);
	static void forwardSubstitutionGPU(Matrix& result, const Matrix& A, const Matrix& b);
	static void backwardSubstitutionGPU(Matrix& result, const Matrix& A, const Matrix& b);


	void toFile(std::string path);
	Matrix separateDiagonal();
	Matrix diagonalCPU() const;
	Matrix lowerCPU() const;
	Matrix upperCPU() const;
	void inverseDiagonalInPlaceCPU();
	void transposeVectorInPlace();
	double vectorEuclideanNorm();
	Matrix lu();
	void print() const;
	~Matrix();

	friend Matrix operator*(const Matrix& a, const Matrix& b);
	friend Matrix operator+(const Matrix& a, const Matrix& b);
	friend Matrix operator-(const Matrix& a, const Matrix& b);
	friend Matrix operator-(const Matrix& a);
	static void refMul(Matrix& result, const Matrix& a, const Matrix& b);
	static void refMulDiag(Matrix& result, const Matrix& a, const Matrix& b);
	static void refAdd(Matrix& result, const Matrix& a, const Matrix& b);
	static void refSub(Matrix& result, const Matrix& a, const Matrix& b);
};

	.text
	.file	"Matrix.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
