#include <stdio.h>
#include <time.h>
#include <iostream>
#include <vector>
#include <math.h>
#include <fstream> 
void checkCUDAError(const char *msg);
#include <cuda_runtime.h>
using namespace std;


// --------------------INPUT DATA---------------------
const int Nx = 24, Ny = 120, Nz = 20; // Number of mass points
float maxtime = 60; // End time [sec]
const int Nstep = 1200; // Number of time steps
__device__ __constant__ float dt = 0.05; // maxtime / Nstep; // Time step size [sec]
float dtcpu=0.05;

const int xlength =  (4 + 2 * Nx)*(2 + Ny)*(2 + Nz); // Solution array in x-direction
const int ylength =  (2 + Nx)*(4 + 2 * Ny)*(2 + Nz); // Solution array in y-direction
const int zlength =  (2 + Nx)*(2 + Ny)*(4 + 2 * Nz); // Solution array in z-direction
const int masslength = Nx * Ny * Nz;
const int kxlength = Nz * Ny * (Nx + 1);
const int kylength = Nz * Nx * (Ny + 1);
const int kzlength = Ny * Nx * (Nz + 1);
const int bxlength = kxlength;
const int bylength = kylength;
const int bzlength = kzlength;






//------------------------DEVICE FUNCTIONS----------------------------//
//x-displacement
__device__ float fxx(int n, int i, int j, int k, float*xold)
{
	return xold[ (Ny + 2)*(4 + 2*Nx) + (k - 1)*(Ny + 2)*(4 + 2*Nx) + 4 +
		2*Nx + (i - 1)*(4 + 2*Nx) + 2 + (2*j - 1)-1];
}

//x-velocity
__device__ float fvx(int n, int i, int j, int k, float*xold)
{
	return xold[ (Ny + 2)*(4 + 2*Nx) + (k - 1)*(Ny + 2)*(4 + 2*Nx) + 4 +
		2*Nx + (i - 1)*(4 + 2*Nx) + 2 + (2*j)-1];
}

//y-displacement
__device__ float fyy(int n, int i, int j, int k, float*yold)
{
	return yold[ (Nx + 2)*(4 + 2*Ny) + (k - 1)*(Nx + 2)*(4 + 2*Ny) + 4 +
		2*Ny + (j - 1)*(4 + 2*Ny) + 2 + (2*i - 1)-1];
}

//y-velocity
__device__ float fvy(int n, int i, int j, int k, float*yold)
{
	return yold[ (Nx + 2)*(4 + 2 * Ny) + (k - 1)*(Nx + 2)*(4 + 2 * Ny) + 4 +
		2 * Ny + (j - 1)*(4 + 2 * Ny) + 2 + (2 * i)-1];
}

//z-displacement
__device__ float fzz(int n, int i, int j, int k, float*zold)
{
	return zold[ (Nx + 2)*(4 + 2*Nz) + (i - 1)*(Nx + 2)*(4 + 2*Nz) + 4 +
		2*Nz + (j - 1)*(4 + 2*Nz) + 2 + (2*k - 1)-1];
}

//z-velocity
__device__ float fvz(int n, int i, int j, int k, float*zold)
{
	return zold[ (Nx + 2)*(4 + 2 * Nz) + (i - 1)*(Nx + 2)*(4 + 2 * Nz) + 4 +
		2 * Nz + (j - 1)*(4 + 2 * Nz) + 2 + (2 * k)-1];
}

//mass
__device__ float fm(int i, int j, int k, float*m)
{
	return m[(k - 1)*Ny*Nx + (i - 1)*Nx + j-1];
}

//x-stiffness
__device__ float fkx(int i, int j, int k, float*kx)
{
	return kx[(k - 1)*Ny*(Nx + 1) + (i - 1)*(Nx + 1) + j-1];
}

//y-stiffness
__device__ float fky(int i, int j, int k, float*ky)
{
	return ky[(k - 1)*Nx*(Ny + 1) + (i - 1)*Nx + j-1];
}

//z-stiffness
__device__ float fkz(int i, int j, int k, float*kz)
{
	return kz[(k - 1)*Nx*Ny + (i - 1)*Nx + j-1];
}

//x-damping
__device__ float fbx(int i, int j, int k, float*bx)
{
	return bx[(k - 1)*Ny*(Nx + 1) + (i - 1)*(Nx + 1) + j-1];
}

//y-damping
__device__ float fby(int i, int j, int k, float*by)
{
	return by[(k - 1)*Nx*(Ny + 1) + (i - 1)*Nx + j-1];
}

//z-damping
__device__ float fbz(int i, int j, int k, float*bz)
{
	return bz[(k - 1)*Nx*Ny + (i - 1)*Nx + j-1];
}

//x-force
__device__ float fFx(int i, int j, int k, float*Fx)
{
	return Fx[(k - 1)*Ny*Nx + (i - 1)*Nx + j-1];
}

//y-force
__device__ float fFy(int i, int j, int k, float*Fy)
{
	return Fy[(k - 1)*Ny*Nx + (i - 1)*Nx + j-1];
}

//z-force
__device__ float fFz(int i, int j, int k, float*Fz)
{
	return Fz[(k - 1)*Ny*Nx + (i - 1)*Nx + j-1];
}

//x-acceleration
__device__ float ax(int i, int j, int k, float*Fx, float*xold, float*kx, float*ky, float*kz, float*bx, float*by, float*bz, float*m)
{
	return (fFx(i, j, k, Fx) - fby(i, j, k, by)*(-fvx(1, -1 + i, j, k, xold) + fvx(1, i, j, k, xold)) -
		fbx(i, j, k, bx)*(-fvx(1, i, -1 + j, k, xold) + fvx(1, i, j, k, xold)) - fbz(i, j, k, bz)*(-fvx(1, i, j, -1 + k, xold) + fvx(1, i, j, k, xold)) +
		fbz(i, j, 1 + k, bz)*(-fvx(1, i, j, k, xold) + fvx(1, i, j, 1 + k, xold)) +
		fbx(i, 1 + j, k, bx)*(-fvx(1, i, j, k, xold) + fvx(1, i, 1 + j, k, xold)) +
		fby(1 + i, j, k, by)*(-fvx(1, i, j, k, xold) + fvx(1, 1 + i, j, k, xold)) -
		fky(i, j, k, ky)*(-fxx(1, -1 + i, j, k, xold) + fxx(1, i, j, k, xold)) - fkx(i, j, k, kx)*(-fxx(1, i, -1 + j, k, xold) + fxx(1, i, j, k, xold)) -
		fkz(i, j, k, kz)*(-fxx(1, i, j, -1 + k, xold) + fxx(1, i, j, k, xold)) +
		fkz(i, j, 1 + k, kz)*(-fxx(1, i, j, k, xold) + fxx(1, i, j, 1 + k, xold)) +
		fkx(i, 1 + j, k, kx)*(-fxx(1, i, j, k, xold) + fxx(1, i, 1 + j, k, xold)) +
		fky(1 + i, j, k, ky)*(-fxx(1, i, j, k, xold) + fxx(1, 1 + i, j, k, xold))) / fm(i, j, k, m);
}

//y-acceleration
__device__ float ay(int i, int j, int k, float*Fy, float*yold, float*kx, float*ky, float*kz, float*bx, float*by, float*bz, float*m)
{
	return (fFy(i, j, k, Fy) - fby(i, j, k, by)*(-fvy(1, -1 + i, j, k, yold) + fvy(1, i, j, k, yold)) -
		fbx(i, j, k, bx)*(-fvy(1, i, -1 + j, k, yold) + fvy(1, i, j, k, yold)) - fbz(i, j, k, bz)*(-fvy(1, i, j, -1 + k, yold) + fvy(1, i, j, k, yold)) +
		fbz(i, j, 1 + k, bz)*(-fvy(1, i, j, k, yold) + fvy(1, i, j, 1 + k, yold)) +
		fbx(i, 1 + j, k, bx)*(-fvy(1, i, j, k, yold) + fvy(1, i, 1 + j, k, yold)) +
		fby(1 + i, j, k, by)*(-fvy(1, i, j, k, yold) + fvy(1, 1 + i, j, k, yold)) -
		fky(i, j, k, ky)*(-fyy(1, -1 + i, j, k, yold) + fyy(1, i, j, k, yold)) - fkx(i, j, k, kx)*(-fyy(1, i, -1 + j, k, yold) + fyy(1, i, j, k, yold)) -
		fkz(i, j, k, kz)*(-fyy(1, i, j, -1 + k, yold) + fyy(1, i, j, k, yold)) +
		fkz(i, j, 1 + k, kz)*(-fyy(1, i, j, k, yold) + fyy(1, i, j, 1 + k, yold)) +
		fkx(i, 1 + j, k, kx)*(-fyy(1, i, j, k, yold) + fyy(1, i, 1 + j, k, yold)) +
		fky(1 + i, j, k, ky)*(-fyy(1, i, j, k, yold) + fyy(1, 1 + i, j, k, yold))) / fm(i, j, k, m);
}

//z-acceleration
__device__ float az(int i, int j, int k, float*Fz, float*zold, float*kx, float*ky, float*kz, float*bx, float*by, float*bz, float*m)
{
	return (fFz(i, j, k, Fz) - fby(i, j, k, by)*(-fvz(1, -1 + i, j, k, zold) + fvz(1, i, j, k, zold)) -
		fbx(i, j, k, bx)*(-fvz(1, i, -1 + j, k, zold) + fvz(1, i, j, k, zold)) - fbz(i, j, k, bz)*(-fvz(1, i, j, -1 + k, zold) + fvz(1, i, j, k, zold)) +
		fbz(i, j, 1 + k, bz)*(-fvz(1, i, j, k, zold) + fvz(1, i, j, 1 + k, zold)) +
		fbx(i, 1 + j, k, bx)*(-fvz(1, i, j, k, zold) + fvz(1, i, 1 + j, k, zold)) +
		fby(1 + i, j, k, by)*(-fvz(1, i, j, k, zold) + fvz(1, 1 + i, j, k, zold)) -
		fky(i, j, k, ky)*(-fzz(1, -1 + i, j, k, zold) + fzz(1, i, j, k, zold)) - fkx(i, j, k, kx)*(-fzz(1, i, -1 + j, k, zold) + fzz(1, i, j, k, zold)) -
		fkz(i, j, k, kz)*(-fzz(1, i, j, -1 + k, zold) + fzz(1, i, j, k, zold)) +
		fkz(i, j, 1 + k, kz)*(-fzz(1, i, j, k, zold) + fzz(1, i, j, 1 + k, zold)) +
		fkx(i, 1 + j, k, kx)*(-fzz(1, i, j, k, zold) + fzz(1, i, 1 + j, k, zold)) +
		fky(1 + i, j, k, ky)*(-fzz(1, i, j, k, zold) + fzz(1, 1 + i, j, k, zold))) / fm(i, j, k, m);
}











__global__ void SolveKernel(int dimBlockX,int dimBlockY,int dimBlockZ,float*xoldd,float*yoldd,float*zoldd,float*xnewd,float*ynewd,float*znewd,float*md,float*kxd,float*kyd,float*kzd,float*bxd,float*byd,float*bzd,float*Fxd,float*Fyd,float*Fzd)
{
//	int tx=threadIdx.x;
//	int ty=threadIdx.y;
	int tx=blockIdx.x*dimBlockX+threadIdx.x;
	int ty=blockIdx.y*dimBlockY+threadIdx.y;
	int tz=blockIdx.z*dimBlockZ+threadIdx.z;
	
	int i=ty+1;
	int j=tx+1;
	int k=tz+1;

	xnewd[ (Ny + 2)*(4 + 2*Nx) + (k - 1)*(Ny + 2)*(4 + 2*Nx) + 4 +
		2 * Nx + (i - 1)*(4 + 2 * Nx) + 2 + (2 * j - 1) - 1] = fxx(1, i, j, k, xoldd) + fvx(1, i, j, k, xoldd)*dt;
	xnewd[ (Ny + 2)*(4 + 2 * Nx) + (k - 1)*(Ny + 2)*(4 + 2 * Nx) + 4 +
		2 * Nx + (i - 1)*(4 + 2 * Nx) + 2 + (2 * j) - 1] = fvx(1, i, j, k, xoldd) + ax(i, j, k, Fxd, xoldd, kxd, kyd, kzd, bxd, byd, bzd, md)*dt;

	ynewd[ (Nx + 2)*(4 + 2*Ny) + (k - 1)*(Nx + 2)*(4 + 2*Ny) + 4 +
		2*Ny + (j - 1)*(4 + 2*Ny) + 2 + (2*i - 1)-1] = fyy(1, i, j, k, yoldd) + fvy(1, i, j, k, yoldd)*dt;
	ynewd[ (Nx + 2)*(4 + 2*Ny) + (k - 1)*(Nx + 2)*(4 + 2*Ny) + 4 +
		2*Ny + (j - 1)*(4 + 2*Ny) + 2 + (2*i)-1] = fvy(1, i, j, k, yoldd) + ay(i, j, k, Fyd, yoldd, kxd, kyd, kzd, bxd, byd, bzd, md)*dt;

	znewd[ (Nx + 2)*(4 + 2*Nz) + (i - 1)*(Nx + 2)*(4 + 2*Nz) + 4 +
		2*Nz + (j - 1)*(4 + 2*Nz) + 2 + (2*k - 1)-1] = fzz(1, i, j, k, zoldd) + fvz(1, i, j, k, zoldd)*dt;
	znewd[ (Nx + 2)*(4 + 2*Nz) + (i - 1)*(Nx + 2)*(4 + 2*Nz) + 4 +
		2*Nz + (j - 1)*(4 + 2*Nz) + 2 + (2*k)-1] = fvz(1, i, j, k, zoldd) + az(i, j, k, Fzd, zoldd, kxd, kyd, kzd, bxd, byd, bzd, md)*dt;
}









void Solve(float*xold,float*yold,float*zold,float*xnew,float*ynew,float*znew,float*m,float*kx,float*ky,float*kz,float*bx,float*by,float*bz,float*Fx,float*Fy,float*Fz)
{
	float *xoldd,*yoldd,*zoldd,*xnewd,*ynewd,*znewd,*md,*kxd,*kyd,*kzd,*bxd,*byd,*bzd,*Fxd,*Fyd,*Fzd;
	
	int sizexoldd=xlength*sizeof(float);
	cudaMalloc((void**)&xoldd,sizexoldd);
	cudaMemcpy(xoldd,xold,sizexoldd,cudaMemcpyHostToDevice);
	int sizeyoldd=ylength*sizeof(float);
	cudaMalloc((void**)&yoldd,sizeyoldd);
	cudaMemcpy(yoldd,yold,sizeyoldd,cudaMemcpyHostToDevice);
	int sizezoldd=zlength*sizeof(float);
	cudaMalloc((void**)&zoldd,sizezoldd);
	cudaMemcpy(zoldd,zold,sizezoldd,cudaMemcpyHostToDevice);
	int sizexnewd=xlength*sizeof(float);
	cudaMalloc((void**)&xnewd,sizexnewd);
	cudaMemcpy(xnewd,xnew,sizexnewd,cudaMemcpyHostToDevice);
	int sizeynewd=ylength*sizeof(float);
	cudaMalloc((void**)&ynewd,sizeynewd);
	cudaMemcpy(ynewd,ynew,sizeynewd,cudaMemcpyHostToDevice);
	int sizeznewd=zlength*sizeof(float);
	cudaMalloc((void**)&znewd,sizeznewd);
	cudaMemcpy(znewd,znew,sizeznewd,cudaMemcpyHostToDevice);
	int sizemd=masslength*sizeof(float);
	cudaMalloc((void**)&md,sizemd);
	cudaMemcpy(md,m,sizemd,cudaMemcpyHostToDevice);
	int sizekxd=kxlength*sizeof(float);
	cudaMalloc((void**)&kxd,sizekxd);
	cudaMemcpy(kxd,kx,sizekxd,cudaMemcpyHostToDevice);
	int sizekyd=kylength*sizeof(float);
	cudaMalloc((void**)&kyd,sizekyd);
	cudaMemcpy(kyd,ky,sizekyd,cudaMemcpyHostToDevice);
	int sizekzd=kzlength*sizeof(float);
	cudaMalloc((void**)&kzd,sizekzd);
	cudaMemcpy(kzd,kz,sizekzd,cudaMemcpyHostToDevice);
	int sizebxd=bxlength*sizeof(float);
	cudaMalloc((void**)&bxd,sizebxd);
	cudaMemcpy(bxd,bx,sizebxd,cudaMemcpyHostToDevice);
	int sizebyd=bylength*sizeof(float);
	cudaMalloc((void**)&byd,sizebyd);
	cudaMemcpy(byd,by,sizebyd,cudaMemcpyHostToDevice);
	int sizebzd=bzlength*sizeof(float);
	cudaMalloc((void**)&bzd,sizebzd);
	cudaMemcpy(bzd,bz,sizebzd,cudaMemcpyHostToDevice);
	int sizeFxd=masslength*sizeof(float);
	cudaMalloc((void**)&Fxd,sizeFxd);
	cudaMemcpy(Fxd,Fx,sizeFxd,cudaMemcpyHostToDevice);
	int sizeFyd=masslength*sizeof(float);
	cudaMalloc((void**)&Fyd,sizeFyd);
	cudaMemcpy(Fyd,Fy,sizeFyd,cudaMemcpyHostToDevice);
	int sizeFzd=masslength*sizeof(float);
	cudaMalloc((void**)&Fzd,sizeFzd);
	cudaMemcpy(Fzd,Fz,sizeFzd,cudaMemcpyHostToDevice);

	//Malloc result
	//cudaMalloc((void**)&Pd,size);
	//Dimensions of the run
	//int SubMtxWidth=SubWidth;
	int NBlockX=4;
	int NBlockY=3;
	int NBlockZ=5;
	int dimBlockX=Nx/NBlockX;
	int dimBlockY=Ny/NBlockY;
	int dimBlockZ=Nz/NBlockZ;
	dim3 dimBlock(dimBlockX,dimBlockY,dimBlockZ);
	dim3 dimGrid(NBlockX,NBlockY,NBlockZ);
	//Running Kernel
	SolveKernel<<<dimGrid,dimBlock>>>(dimBlockX,dimBlockY,dimBlockZ,xoldd,yoldd,zoldd,xnewd,ynewd,znewd,md,kxd,kyd,kzd,bxd,byd,bzd,Fxd,Fyd,Fzd);
	cudaThreadSynchronize();
	//Copy data back
	cudaMemcpy(xnew,xnewd,sizexnewd,cudaMemcpyDeviceToHost);
	cudaMemcpy(ynew,ynewd,sizeynewd,cudaMemcpyDeviceToHost);
	cudaMemcpy(znew,znewd,sizeznewd,cudaMemcpyDeviceToHost);
	checkCUDAError("memcpy");
	//Free memory
	//cudaFree(Md);
	//cudaFree(Nd);
	//cudaFree(Pd);
	//NEWSHIT
	cudaFree(xoldd);
	cudaFree(yoldd);
	cudaFree(zoldd);
	cudaFree(xnewd);
	cudaFree(ynewd);
	cudaFree(znewd);
	cudaFree(md);
	cudaFree(kxd);
	cudaFree(kyd);
	cudaFree(kzd);
	cudaFree(bxd);
	cudaFree(byd);
	cudaFree(bzd);
	cudaFree(Fxd);
	cudaFree(Fyd);
	cudaFree(Fzd);
}
















int main(int argc,char* argv[])
{
float *xold,*yold,*zold,*xnew,*ynew,*znew,*m,*kx,*ky,*kz,*bx,*by,*bz,*Fx,*Fy,*Fz;

//----------------------------------INITIALIZATION START----------------------------------
// Solution vectors 
xold=(float *)malloc(xlength*sizeof(float));
yold=(float *)malloc(ylength*sizeof(float));
zold=(float *)malloc(zlength*sizeof(float));
xnew=(float *)malloc(xlength*sizeof(float));
ynew=(float *)malloc(ylength*sizeof(float));
znew=(float *)malloc(zlength*sizeof(float));

// Mass vector
m=(float *)malloc(masslength*sizeof(float));

// Stiffness vectors
kx=(float *)malloc(kxlength*sizeof(float));
ky=(float *)malloc(kylength*sizeof(float));
kz=(float *)malloc(kzlength*sizeof(float));

// Damping vectors
bx=(float *)malloc(bxlength*sizeof(float));
by=(float *)malloc(bylength*sizeof(float));
bz=(float *)malloc(bzlength*sizeof(float));

// Force vectors
Fx=(float *)malloc(masslength*sizeof(float));
Fy=(float *)malloc(masslength*sizeof(float));
Fz=(float *)malloc(masslength*sizeof(float));

// Initial conditions
for (int i = 0; i < xlength ; i++)
{
	xold[i]=0.0f;
	xnew[i]=0.0f;
}

for (int i = 0; i < ylength ; i++)
{
	yold[i]=0.0f;
	ynew[i]=0.0f;
}

for (int i = 0; i < zlength ; i++)
{
	zold[i]=0.0f;
	znew[i]=0.0f;
}

// Mass [kg] and forces
for (int i = 0; i < masslength ; i++)
{
	m[i]=1.0f;
	Fx[i]=0.0f;
	Fy[i]=0.0f;
	Fz[i]=0.0f;
}

// Stiffness [N/m] and damping [N sec/m] in x-direction
for (int i = 0; i < kxlength ; i++)
{
	kx[i]=0.2f;
	bx[i]=0.05f;
}

// Stiffness [N/m] and damping [N sec/m] in y-direction
for (int i = 0; i < kylength ; i++)
{
	ky[i]=0.2f;
	by[i]=0.05f;
}

// Stiffness [N/m] and damping [N sec/m] in z-direction
for (int i = 0; i < kzlength ; i++)
{
	kz[i]=0.2f;
	bz[i]=0.05f;
}
//----------------------------------INITIALIZATION END--------------------------------------



//-------------------------------BOUNDARY CONDITIONS START----------------------------------
// No connections with Top wall B.C.'s
for (int i = 1; i <= Nx; i++)
	{
		for (int k = 1; k <= Nz; k++)
		{
			ky[i + Nx*Ny + (-1 + k)*Nx*(1 + Ny) - 1] = 0.0f;
			by[i + Nx*Ny + (-1 + k)*Nx*(1 + Ny) - 1] = 0.0f;
		}
	}
//--------------------------------BOUNDARY CONDITIONS END-----------------------------------



//--------------------------------------SOLVER START-----------------------------------------
clock_t t;
t=clock();

for (int n = 1; n <= Nstep-1; n++)
	{
		// Excitation
		Fx[(2 - 1)*Ny*Nx + (6 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu); // omega = 3 [rad/sec]
		Fy[(2 - 1)*Ny*Nx + (6 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);
		Fz[(2 - 1)*Ny*Nx + (6 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);

		Fx[(2 - 1)*Ny*Nx + (7 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);
		Fy[(2 - 1)*Ny*Nx + (7 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);
		Fz[(2 - 1)*Ny*Nx + (7 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);

		Fx[(2 - 1)*Ny*Nx + (5 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);
		Fy[(2 - 1)*Ny*Nx + (5 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);
		Fz[(2 - 1)*Ny*Nx + (5 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);

		Solve(xold,yold,zold,xnew,ynew,znew,m,kx,ky,kz,bx,by,bz,Fx,Fy,Fz);
		cudaThreadSynchronize();
		// OLD=NEW
		for (int ix = 0; ix < xlength; ix++)
		{
			xold[ix] = xnew[ix];
		}
		for (int iy = 0; iy < ylength; iy++)
		{
			yold[iy] = ynew[iy];
		}
		for (int iz = 0; iz < zlength; iz++)
		{
			zold[iz] = znew[iz];
		}
	}
ofstream fout("test.txt");
if (fout.is_open())
	{
		//file opened successfully so we are here
		cout << "File Opened successfully!!!. Writing data from array to file" << endl;

		for (int j = 0; j < zlength; j++)
			{
				fout << znew[j] << ' '; //writing ith character of array in the file
			}
			fout << '\n';
		cout << "Array data successfully saved into the file test.txt" << endl;
	}
	else //file could not be opened
	{
		cout << "File could not be opened." << endl;
	}

t=clock()-t;
printf("%f seconds\n",((float)t)/CLOCKS_PER_SEC);
printf("%f,%f,%f\n",xold[60],yold[60],zold[60]);

free(xold);
free(yold);
free(zold);
free(xnew);
free(ynew);
free(znew);
free(m);
free(kx);
free(ky);
free(kz);
free(bx);
free(by);
free(bz);
free(Fx);
free(Fy);
free(Fz);

return 0;
}


void checkCUDAError(const char *msg)
{
cudaError_t err = cudaGetLastError();
if(cudaSuccess!= err)
{
fprintf(stderr,"Cuda error: %s: %s.\n",msg,cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
}
	.file	"tmpxft_002f113d_00000000-6_msd.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4153:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4153:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0, @function
_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0:
.LFB4892:
	.cfi_startproc
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movq	(%rdi), %rax
	movb	%sil, 12(%rsp)
	movq	-24(%rax), %rax
	cmpq	$0, 16(%rdi,%rax)
	je	.L3
	leaq	12(%rsp), %rsi
	movl	$1, %edx
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
	jmp	.L2
.L3:
	movsbl	12(%rsp), %esi
	call	_ZNSo3putEc@PLT
.L2:
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4892:
	.size	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0, .-_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	.globl	_Z3fxxiiiiPf
	.type	_Z3fxxiiiiPf, @function
_Z3fxxiiiiPf:
.LFB4129:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4129:
	.size	_Z3fxxiiiiPf, .-_Z3fxxiiiiPf
	.globl	_Z3fvxiiiiPf
	.type	_Z3fvxiiiiPf, @function
_Z3fvxiiiiPf:
.LFB4861:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4861:
	.size	_Z3fvxiiiiPf, .-_Z3fvxiiiiPf
	.globl	_Z3fyyiiiiPf
	.type	_Z3fyyiiiiPf, @function
_Z3fyyiiiiPf:
.LFB4863:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4863:
	.size	_Z3fyyiiiiPf, .-_Z3fyyiiiiPf
	.globl	_Z3fvyiiiiPf
	.type	_Z3fvyiiiiPf, @function
_Z3fvyiiiiPf:
.LFB4865:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4865:
	.size	_Z3fvyiiiiPf, .-_Z3fvyiiiiPf
	.globl	_Z3fzziiiiPf
	.type	_Z3fzziiiiPf, @function
_Z3fzziiiiPf:
.LFB4867:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4867:
	.size	_Z3fzziiiiPf, .-_Z3fzziiiiPf
	.globl	_Z3fvziiiiPf
	.type	_Z3fvziiiiPf, @function
_Z3fvziiiiPf:
.LFB4869:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4869:
	.size	_Z3fvziiiiPf, .-_Z3fvziiiiPf
	.globl	_Z2fmiiiPf
	.type	_Z2fmiiiPf, @function
_Z2fmiiiPf:
.LFB4135:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4135:
	.size	_Z2fmiiiPf, .-_Z2fmiiiPf
	.globl	_Z3fkxiiiPf
	.type	_Z3fkxiiiPf, @function
_Z3fkxiiiPf:
.LFB4871:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4871:
	.size	_Z3fkxiiiPf, .-_Z3fkxiiiPf
	.globl	_Z3fkyiiiPf
	.type	_Z3fkyiiiPf, @function
_Z3fkyiiiPf:
.LFB4873:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4873:
	.size	_Z3fkyiiiPf, .-_Z3fkyiiiPf
	.globl	_Z3fkziiiPf
	.type	_Z3fkziiiPf, @function
_Z3fkziiiPf:
.LFB4875:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4875:
	.size	_Z3fkziiiPf, .-_Z3fkziiiPf
	.globl	_Z3fbxiiiPf
	.type	_Z3fbxiiiPf, @function
_Z3fbxiiiPf:
.LFB4877:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4877:
	.size	_Z3fbxiiiPf, .-_Z3fbxiiiPf
	.globl	_Z3fbyiiiPf
	.type	_Z3fbyiiiPf, @function
_Z3fbyiiiPf:
.LFB4879:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4879:
	.size	_Z3fbyiiiPf, .-_Z3fbyiiiPf
	.globl	_Z3fbziiiPf
	.type	_Z3fbziiiPf, @function
_Z3fbziiiPf:
.LFB4881:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4881:
	.size	_Z3fbziiiPf, .-_Z3fbziiiPf
	.globl	_Z3fFxiiiPf
	.type	_Z3fFxiiiPf, @function
_Z3fFxiiiPf:
.LFB4883:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4883:
	.size	_Z3fFxiiiPf, .-_Z3fFxiiiPf
	.globl	_Z3fFyiiiPf
	.type	_Z3fFyiiiPf, @function
_Z3fFyiiiPf:
.LFB4885:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4885:
	.size	_Z3fFyiiiPf, .-_Z3fFyiiiPf
	.globl	_Z3fFziiiPf
	.type	_Z3fFziiiPf, @function
_Z3fFziiiPf:
.LFB4887:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4887:
	.size	_Z3fFziiiPf, .-_Z3fFziiiPf
	.globl	_Z2axiiiPfS_S_S_S_S_S_S_S_
	.type	_Z2axiiiPfS_S_S_S_S_S_S_S_, @function
_Z2axiiiPfS_S_S_S_S_S_S_S_:
.LFB4145:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4145:
	.size	_Z2axiiiPfS_S_S_S_S_S_S_S_, .-_Z2axiiiPfS_S_S_S_S_S_S_S_
	.globl	_Z2ayiiiPfS_S_S_S_S_S_S_S_
	.type	_Z2ayiiiPfS_S_S_S_S_S_S_S_, @function
_Z2ayiiiPfS_S_S_S_S_S_S_S_:
.LFB4889:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4889:
	.size	_Z2ayiiiPfS_S_S_S_S_S_S_S_, .-_Z2ayiiiPfS_S_S_S_S_S_S_S_
	.globl	_Z2aziiiPfS_S_S_S_S_S_S_S_
	.type	_Z2aziiiPfS_S_S_S_S_S_S_S_, @function
_Z2aziiiPfS_S_S_S_S_S_S_S_:
.LFB4891:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4891:
	.size	_Z2aziiiPfS_S_S_S_S_S_S_S_, .-_Z2aziiiPfS_S_S_S_S_S_S_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Cuda error: %s: %s.\n"
	.text
	.globl	_Z14checkCUDAErrorPKc
	.type	_Z14checkCUDAErrorPKc, @function
_Z14checkCUDAErrorPKc:
.LFB4150:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	call	cudaGetLastError@PLT
	testl	%eax, %eax
	je	.L45
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	movq	stderr(%rip), %rdi
	movq	%rbx, %rcx
	leaq	.LC0(%rip), %rdx
	movq	%rax, %r8
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L45:
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4150:
	.size	_Z14checkCUDAErrorPKc, .-_Z14checkCUDAErrorPKc
	.globl	_Z64__device_stub__Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.type	_Z64__device_stub__Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, @function
_Z64__device_stub__Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_:
.LFB4175:
	.cfi_startproc
	endbr64
	subq	$376, %rsp
	.cfi_def_cfa_offset 384
	movq	384(%rsp), %rax
	movl	%edi, 140(%rsp)
	leaq	160(%rsp), %rdi
	movl	%esi, 136(%rsp)
	leaq	172(%rsp), %rsi
	movq	%rax, 96(%rsp)
	movq	392(%rsp), %rax
	movl	%edx, 132(%rsp)
	leaq	144(%rsp), %rdx
	movq	%rax, 88(%rsp)
	movq	400(%rsp), %rax
	movq	%rcx, 120(%rsp)
	leaq	152(%rsp), %rcx
	movq	%rax, 80(%rsp)
	movq	408(%rsp), %rax
	movq	%r8, 112(%rsp)
	movq	%rax, 72(%rsp)
	movq	416(%rsp), %rax
	movq	%r9, 104(%rsp)
	movq	%rax, 64(%rsp)
	movq	424(%rsp), %rax
	movq	%rax, 56(%rsp)
	movq	432(%rsp), %rax
	movq	%rax, 48(%rsp)
	movq	440(%rsp), %rax
	movq	%rax, 40(%rsp)
	movq	448(%rsp), %rax
	movq	%rax, 32(%rsp)
	movq	456(%rsp), %rax
	movq	%rax, 24(%rsp)
	movq	464(%rsp), %rax
	movq	%rax, 16(%rsp)
	movq	472(%rsp), %rax
	movq	%rax, 8(%rsp)
	movq	480(%rsp), %rax
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 360(%rsp)
	xorl	%eax, %eax
	leaq	140(%rsp), %rax
	movl	$1, 168(%rsp)
	movq	%rax, 208(%rsp)
	leaq	136(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	132(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	120(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 336(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 344(%rsp)
	movq	%rsp, %rax
	movq	%rax, 352(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 160(%rsp)
	movq	%rax, 172(%rsp)
	movl	$1, 180(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L48
	pushq	152(%rsp)
	.cfi_def_cfa_offset 392
	leaq	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_(%rip), %rdi
	pushq	152(%rsp)
	.cfi_def_cfa_offset 400
	movq	188(%rsp), %rcx
	movl	196(%rsp), %r8d
	movq	176(%rsp), %rsi
	movl	184(%rsp), %edx
	leaq	224(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 392
	popq	%rdx
	.cfi_def_cfa_offset 384
.L48:
	movq	360(%rsp), %rax
	subq	%fs:40, %rax
	je	.L50
	call	__stack_chk_fail@PLT
.L50:
	addq	$376, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4175:
	.size	_Z64__device_stub__Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, .-_Z64__device_stub__Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.globl	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.type	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, @function
_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_:
.LFB4176:
	.cfi_startproc
	endbr64
	jmp	_Z64__device_stub__Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.cfi_endproc
.LFE4176:
	.size	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, .-_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.section	.rodata.str1.1
.LC1:
	.string	"memcpy"
	.text
	.globl	_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.type	_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, @function
_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_:
.LFB4148:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	movl	$558272, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdx, %rbx
	subq	$248, %rsp
	.cfi_def_cfa_offset 304
	movq	328(%rsp), %rax
	movq	%r8, 8(%rsp)
	leaq	80(%rsp), %rdi
	movq	%r9, 16(%rsp)
	movq	304(%rsp), %r15
	movq	%rax, 24(%rsp)
	movq	336(%rsp), %rax
	movq	%rcx, (%rsp)
	movq	312(%rsp), %r14
	movq	%rax, 32(%rsp)
	movq	344(%rsp), %rax
	movq	320(%rsp), %r13
	movq	%rax, 40(%rsp)
	movq	352(%rsp), %rax
	movq	%rax, 48(%rsp)
	movq	360(%rsp), %rax
	movq	%rax, 56(%rsp)
	movq	368(%rsp), %rax
	movq	%rax, 64(%rsp)
	movq	376(%rsp), %rax
	movq	%rax, 72(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 232(%rsp)
	xorl	%eax, %eax
	call	cudaMalloc@PLT
	movq	80(%rsp), %rdi
	movl	$1, %ecx
	movq	%r12, %rsi
	movl	$558272, %edx
	call	cudaMemcpy@PLT
	leaq	88(%rsp), %rdi
	movl	$558272, %esi
	call	cudaMalloc@PLT
	movq	88(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbp, %rsi
	movl	$558272, %edx
	call	cudaMemcpy@PLT
	leaq	96(%rsp), %rdi
	movl	$558272, %esi
	call	cudaMalloc@PLT
	movq	96(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$558272, %edx
	call	cudaMemcpy@PLT
	leaq	104(%rsp), %rdi
	movl	$558272, %esi
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	movq	104(%rsp), %rdi
	movl	$1, %ecx
	movl	$558272, %edx
	call	cudaMemcpy@PLT
	leaq	112(%rsp), %rdi
	movl	$558272, %esi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rsi
	movq	112(%rsp), %rdi
	movl	$1, %ecx
	movl	$558272, %edx
	call	cudaMemcpy@PLT
	leaq	120(%rsp), %rdi
	movl	$558272, %esi
	call	cudaMalloc@PLT
	movq	16(%rsp), %rsi
	movq	120(%rsp), %rdi
	movl	$1, %ecx
	movl	$558272, %edx
	call	cudaMemcpy@PLT
	leaq	128(%rsp), %rdi
	movl	$230400, %esi
	call	cudaMalloc@PLT
	movq	128(%rsp), %rdi
	movl	$1, %ecx
	movq	%r15, %rsi
	movl	$230400, %edx
	call	cudaMemcpy@PLT
	leaq	136(%rsp), %rdi
	movl	$240000, %esi
	call	cudaMalloc@PLT
	movq	136(%rsp), %rdi
	movl	$1, %ecx
	movq	%r14, %rsi
	movl	$240000, %edx
	call	cudaMemcpy@PLT
	leaq	144(%rsp), %rdi
	movl	$232320, %esi
	call	cudaMalloc@PLT
	movq	144(%rsp), %rdi
	movl	$1, %ecx
	movq	%r13, %rsi
	movl	$232320, %edx
	call	cudaMemcpy@PLT
	leaq	152(%rsp), %rdi
	movl	$241920, %esi
	call	cudaMalloc@PLT
	movq	24(%rsp), %rsi
	movl	$1, %ecx
	movq	152(%rsp), %rdi
	movl	$241920, %edx
	call	cudaMemcpy@PLT
	leaq	160(%rsp), %rdi
	movl	$240000, %esi
	call	cudaMalloc@PLT
	movq	32(%rsp), %rsi
	movl	$1, %ecx
	movq	160(%rsp), %rdi
	movl	$240000, %edx
	call	cudaMemcpy@PLT
	leaq	168(%rsp), %rdi
	movl	$232320, %esi
	call	cudaMalloc@PLT
	movq	40(%rsp), %rsi
	movl	$1, %ecx
	movq	168(%rsp), %rdi
	movl	$232320, %edx
	call	cudaMemcpy@PLT
	leaq	176(%rsp), %rdi
	movl	$241920, %esi
	call	cudaMalloc@PLT
	movq	48(%rsp), %rsi
	movl	$1, %ecx
	movq	176(%rsp), %rdi
	movl	$241920, %edx
	call	cudaMemcpy@PLT
	leaq	184(%rsp), %rdi
	movl	$230400, %esi
	call	cudaMalloc@PLT
	movq	56(%rsp), %rsi
	movl	$1, %ecx
	movq	184(%rsp), %rdi
	movl	$230400, %edx
	call	cudaMemcpy@PLT
	leaq	192(%rsp), %rdi
	movl	$230400, %esi
	call	cudaMalloc@PLT
	movq	64(%rsp), %rsi
	movl	$1, %ecx
	movq	192(%rsp), %rdi
	movl	$230400, %edx
	call	cudaMemcpy@PLT
	leaq	200(%rsp), %rdi
	movl	$230400, %esi
	call	cudaMalloc@PLT
	movq	72(%rsp), %rsi
	movl	$1, %ecx
	movq	200(%rsp), %rdi
	movl	$230400, %edx
	call	cudaMemcpy@PLT
	movl	$3221225473, %edi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$2, %rdi
	movl	$4, %ecx
	movl	$5, %esi
	movabsq	$171798691846, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L54
	pushq	%rax
	.cfi_def_cfa_offset 312
	movl	$4, %edx
	movl	$40, %esi
	movl	$6, %edi
	pushq	208(%rsp)
	.cfi_def_cfa_offset 320
	pushq	208(%rsp)
	.cfi_def_cfa_offset 328
	pushq	208(%rsp)
	.cfi_def_cfa_offset 336
	pushq	208(%rsp)
	.cfi_def_cfa_offset 344
	pushq	208(%rsp)
	.cfi_def_cfa_offset 352
	pushq	208(%rsp)
	.cfi_def_cfa_offset 360
	pushq	208(%rsp)
	.cfi_def_cfa_offset 368
	pushq	208(%rsp)
	.cfi_def_cfa_offset 376
	pushq	208(%rsp)
	.cfi_def_cfa_offset 384
	pushq	208(%rsp)
	.cfi_def_cfa_offset 392
	pushq	208(%rsp)
	.cfi_def_cfa_offset 400
	pushq	208(%rsp)
	.cfi_def_cfa_offset 408
	pushq	208(%rsp)
	.cfi_def_cfa_offset 416
	movq	208(%rsp), %r9
	movq	200(%rsp), %r8
	movq	192(%rsp), %rcx
	call	_Z64__device_stub__Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_iiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	addq	$112, %rsp
	.cfi_def_cfa_offset 304
.L54:
	call	cudaThreadSynchronize@PLT
	movq	104(%rsp), %rsi
	movq	(%rsp), %rdi
	movl	$2, %ecx
	movl	$558272, %edx
	call	cudaMemcpy@PLT
	movq	112(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$2, %ecx
	movl	$558272, %edx
	call	cudaMemcpy@PLT
	movq	120(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	$2, %ecx
	movl	$558272, %edx
	call	cudaMemcpy@PLT
	leaq	.LC1(%rip), %rdi
	call	_Z14checkCUDAErrorPKc
	movq	80(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	movq	112(%rsp), %rdi
	call	cudaFree@PLT
	movq	120(%rsp), %rdi
	call	cudaFree@PLT
	movq	128(%rsp), %rdi
	call	cudaFree@PLT
	movq	136(%rsp), %rdi
	call	cudaFree@PLT
	movq	144(%rsp), %rdi
	call	cudaFree@PLT
	movq	152(%rsp), %rdi
	call	cudaFree@PLT
	movq	160(%rsp), %rdi
	call	cudaFree@PLT
	movq	168(%rsp), %rdi
	call	cudaFree@PLT
	movq	176(%rsp), %rdi
	call	cudaFree@PLT
	movq	184(%rsp), %rdi
	call	cudaFree@PLT
	movq	192(%rsp), %rdi
	call	cudaFree@PLT
	movq	200(%rsp), %rdi
	call	cudaFree@PLT
	movq	232(%rsp), %rax
	subq	%fs:40, %rax
	je	.L55
	call	__stack_chk_fail@PLT
.L55:
	addq	$248, %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
.LFE4148:
	.size	_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, .-_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.section	.rodata.str1.1
.LC2:
	.string	"_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_"
.LC3:
	.string	"dt"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4178:
	.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	.LC2(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	movl	$4, %r9d
	leaq	_ZL2dt(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.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
.LFE4178:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata.str1.1
.LC8:
	.string	"test.txt"
.LC9:
	.string	"File Opened successfully!!!. Writing data from array to file"
.LC10:
	.string	"Array data successfully saved into the file test.txt"
.LC12:
	.string	"%f seconds\n"
.LC13:
	.string	"File could not be opened."
.LC14:
	.string	"%f,%f,%f\n"
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB4149:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4149
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$558272, %edi
	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
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$648, %rsp
	.cfi_def_cfa_offset 704
	movq	%fs:40, %rax
	movq	%rax, 632(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	movl	$558272, %edi
	movq	%rax, 8(%rsp)
	call	malloc@PLT
	movl	$558272, %edi
	movq	%rax, 16(%rsp)
	call	malloc@PLT
	movl	$558272, %edi
	movq	%rax, 24(%rsp)
	call	malloc@PLT
	movl	$558272, %edi
	movq	%rax, 40(%rsp)
	call	malloc@PLT
	movl	$558272, %edi
	movq	%rax, 48(%rsp)
	call	malloc@PLT
	movl	$230400, %edi
	movq	%rax, 32(%rsp)
	call	malloc@PLT
	movl	$240000, %edi
	movq	%rax, 56(%rsp)
	call	malloc@PLT
	movl	$232320, %edi
	movq	%rax, 64(%rsp)
	call	malloc@PLT
	movl	$241920, %edi
	movq	%rax, %r13
	call	malloc@PLT
	movl	$240000, %edi
	movq	%rax, 72(%rsp)
	call	malloc@PLT
	movl	$232320, %edi
	movq	%rax, 80(%rsp)
	call	malloc@PLT
	movl	$241920, %edi
	movq	%rax, %r15
	call	malloc@PLT
	movl	$230400, %edi
	movq	%rax, 88(%rsp)
	call	malloc@PLT
	movl	$230400, %edi
	movq	%rax, %r12
	call	malloc@PLT
	movl	$230400, %edi
	movq	%rax, %rbp
	call	malloc@PLT
	movq	8(%rsp), %rdi
	movl	$139568, %ecx
	xorl	%edx, %edx
	movq	%rax, %rbx
	xorl	%eax, %eax
	movss	.LC4(%rip), %xmm0
	rep stosl
	movq	40(%rsp), %rdi
	movl	$139568, %ecx
	rep stosl
	movq	16(%rsp), %rdi
	movl	$139568, %ecx
	rep stosl
	movq	48(%rsp), %rdi
	movl	$139568, %ecx
	rep stosl
	movq	24(%rsp), %rdi
	movl	$139568, %ecx
	rep stosl
	movq	32(%rsp), %rdi
	movl	$139568, %ecx
	rep stosl
.L60:
	movq	56(%rsp), %rax
	movss	%xmm0, (%rax,%rdx,4)
	incq	%rdx
	cmpq	$57600, %rdx
	jne	.L60
	xorl	%eax, %eax
	movq	%r12, %rdi
	movq	%rdx, %rcx
	movss	.LC5(%rip), %xmm1
	rep stosl
	movq	%rbp, %rdi
	movq	%rdx, %rcx
	movss	.LC6(%rip), %xmm0
	rep stosl
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	rep stosl
	xorl	%eax, %eax
.L61:
	movq	64(%rsp), %rdx
	movss	%xmm1, (%rdx,%rax)
	movq	80(%rsp), %rdx
	movss	%xmm0, (%rdx,%rax)
	addq	$4, %rax
	cmpq	$240000, %rax
	jne	.L61
	movss	.LC5(%rip), %xmm1
	movss	.LC6(%rip), %xmm0
	xorl	%eax, %eax
.L62:
	movss	%xmm1, 0(%r13,%rax)
	movss	%xmm0, (%r15,%rax)
	addq	$4, %rax
	cmpq	$232320, %rax
	jne	.L62
	movss	.LC5(%rip), %xmm1
	movss	.LC6(%rip), %xmm0
	xorl	%eax, %eax
.L63:
	movq	72(%rsp), %rdx
	movss	%xmm1, (%rdx,%rax)
	movq	88(%rsp), %rdx
	movss	%xmm0, (%rdx,%rax)
	addq	$4, %rax
	cmpq	$241920, %rax
	jne	.L63
	movl	$4, %edx
.L64:
	movq	%rdx, %rax
	movl	$20, %ecx
.L65:
	movl	$0x00000000, 11516(%r13,%rax)
	movl	$0x00000000, 11516(%r15,%rax)
	addq	$11616, %rax
	decl	%ecx
	jne	.L65
	addq	$4, %rdx
	cmpq	$100, %rdx
	jne	.L64
	call	clock@PLT
	movl	$3, %r14d
	movq	%rax, 96(%rsp)
.L67:
	cvtsi2ssl	%r14d, %xmm0
	mulss	dtcpu(%rip), %xmm0
	addl	$3, %r14d
	call	sinf@PLT
	movss	%xmm0, 12028(%r12)
	movss	%xmm0, 12028(%rbp)
	movss	%xmm0, 12028(%rbx)
	movss	%xmm0, 12124(%r12)
	movss	%xmm0, 12124(%rbp)
	movss	%xmm0, 12124(%rbx)
	movss	%xmm0, 11932(%r12)
	movss	%xmm0, 11932(%rbp)
	movss	%xmm0, 11932(%rbx)
	pushq	%rbx
	.cfi_def_cfa_offset 712
	pushq	%rbp
	.cfi_def_cfa_offset 720
	pushq	%r12
	.cfi_def_cfa_offset 728
	pushq	112(%rsp)
	.cfi_def_cfa_offset 736
	pushq	%r15
	.cfi_def_cfa_offset 744
	pushq	120(%rsp)
	.cfi_def_cfa_offset 752
	pushq	120(%rsp)
	.cfi_def_cfa_offset 760
	pushq	%r13
	.cfi_def_cfa_offset 768
	pushq	128(%rsp)
	.cfi_def_cfa_offset 776
	pushq	128(%rsp)
	.cfi_def_cfa_offset 784
	movq	120(%rsp), %rcx
	movq	112(%rsp), %r9
	movq	128(%rsp), %r8
	movq	104(%rsp), %rdx
	movq	96(%rsp), %rsi
	movq	88(%rsp), %rdi
.LEHB0:
	call	_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	addq	$80, %rsp
	.cfi_def_cfa_offset 704
	call	cudaThreadSynchronize@PLT
	movq	8(%rsp), %rdi
	movq	40(%rsp), %rsi
	movl	$139568, %ecx
	rep movsl
	movq	16(%rsp), %rdi
	movq	48(%rsp), %rsi
	movl	$139568, %ecx
	rep movsl
	movq	24(%rsp), %rdi
	movq	32(%rsp), %rsi
	movl	$139568, %ecx
	rep movsl
	cmpl	$3600, %r14d
	jne	.L67
	leaq	120(%rsp), %rdi
	movl	$16, %edx
	leaq	.LC8(%rip), %rsi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode@PLT
.LEHE0:
	leaq	232(%rsp), %rdi
	call	_ZNKSt12__basic_fileIcE7is_openEv@PLT
	testb	%al, %al
	je	.L68
	leaq	.LC9(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
.LEHB1:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	leaq	120(%rsp), %rax
	xorl	%r14d, %r14d
	movq	%rax, 104(%rsp)
.L69:
	movq	32(%rsp), %rax
	movq	104(%rsp), %rdi
	cvtss2sd	(%rax,%r14,4), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movl	$32, %esi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	incq	%r14
	cmpq	$139568, %r14
	jne	.L69
	leaq	120(%rsp), %rdi
	movl	$10, %esi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	leaq	.LC10(%rip), %rsi
.L89:
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	call	clock@PLT
	movq	96(%rsp), %rdx
	movl	$2, %edi
	leaq	.LC12(%rip), %rsi
	subq	%rdx, %rax
	cvtsi2ssq	%rax, %xmm0
	movb	$1, %al
	divss	.LC11(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	jmp	.L90
.L68:
	leaq	.LC13(%rip), %rsi
	jmp	.L89
.L90:
	movq	8(%rsp), %rax
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	cvtss2sd	240(%rax), %xmm0
	movq	24(%rsp), %rax
	cvtss2sd	240(%rax), %xmm2
	movq	16(%rsp), %rax
	cvtss2sd	240(%rax), %xmm1
	movb	$3, %al
	call	__printf_chk@PLT
.LEHE1:
	movq	8(%rsp), %rdi
	call	free@PLT
	movq	16(%rsp), %rdi
	call	free@PLT
	movq	24(%rsp), %rdi
	call	free@PLT
	movq	40(%rsp), %rdi
	call	free@PLT
	movq	48(%rsp), %rdi
	call	free@PLT
	movq	32(%rsp), %rdi
	call	free@PLT
	movq	56(%rsp), %rdi
	call	free@PLT
	movq	64(%rsp), %rdi
	call	free@PLT
	movq	%r13, %rdi
	call	free@PLT
	movq	72(%rsp), %rdi
	call	free@PLT
	movq	80(%rsp), %rdi
	call	free@PLT
	movq	%r15, %rdi
	call	free@PLT
	movq	88(%rsp), %rdi
	call	free@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	leaq	120(%rsp), %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	632(%rsp), %rax
	subq	%fs:40, %rax
	je	.L74
	jmp	.L84
.L75:
	endbr64
	movq	%rax, %rbx
.L72:
	leaq	120(%rsp), %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	632(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L84
	movq	%rbx, %rdi
.LEHB2:
	call	_Unwind_Resume@PLT
.LEHE2:
.L84:
	call	__stack_chk_fail@PLT
.L74:
	addq	$648, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	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
.LFE4149:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4149:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4149-.LLSDACSB4149
.LLSDACSB4149:
	.uleb128 .LEHB0-.LFB4149
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB4149
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L75-.LFB4149
	.uleb128 0
	.uleb128 .LEHB2-.LFB4149
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
.LLSDACSE4149:
	.section	.text.startup
	.size	main, .-main
	.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
	.globl	dtcpu
	.data
	.align 4
	.type	dtcpu, @object
	.size	dtcpu, 4
dtcpu:
	.long	1028443341
	.local	_ZL2dt
	.comm	_ZL2dt,4,4
	.globl	maxtime
	.align 4
	.type	maxtime, @object
	.size	maxtime, 4
maxtime:
	.long	1114636288
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC4:
	.long	1065353216
	.align 4
.LC5:
	.long	1045220557
	.align 4
.LC6:
	.long	1028443341
	.align 4
.LC11:
	.long	1232348160
	.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:

#include <stdio.h>
#include <time.h>
#include <iostream>
#include <vector>
#include <math.h>
#include <fstream> 
void checkCUDAError(const char *msg);
#include <hip/hip_runtime.h>
using namespace std;


// --------------------INPUT DATA---------------------
const int Nx = 24, Ny = 120, Nz = 20; // Number of mass points
float maxtime = 60; // End time [sec]
const int Nstep = 1200; // Number of time steps
__device__ __constant__ float dt = 0.05; // maxtime / Nstep; // Time step size [sec]
float dtcpu=0.05;

const int xlength =  (4 + 2 * Nx)*(2 + Ny)*(2 + Nz); // Solution array in x-direction
const int ylength =  (2 + Nx)*(4 + 2 * Ny)*(2 + Nz); // Solution array in y-direction
const int zlength =  (2 + Nx)*(2 + Ny)*(4 + 2 * Nz); // Solution array in z-direction
const int masslength = Nx * Ny * Nz;
const int kxlength = Nz * Ny * (Nx + 1);
const int kylength = Nz * Nx * (Ny + 1);
const int kzlength = Ny * Nx * (Nz + 1);
const int bxlength = kxlength;
const int bylength = kylength;
const int bzlength = kzlength;






//------------------------DEVICE FUNCTIONS----------------------------//
//x-displacement
__device__ float fxx(int n, int i, int j, int k, float*xold)
{
	return xold[ (Ny + 2)*(4 + 2*Nx) + (k - 1)*(Ny + 2)*(4 + 2*Nx) + 4 +
		2*Nx + (i - 1)*(4 + 2*Nx) + 2 + (2*j - 1)-1];
}

//x-velocity
__device__ float fvx(int n, int i, int j, int k, float*xold)
{
	return xold[ (Ny + 2)*(4 + 2*Nx) + (k - 1)*(Ny + 2)*(4 + 2*Nx) + 4 +
		2*Nx + (i - 1)*(4 + 2*Nx) + 2 + (2*j)-1];
}

//y-displacement
__device__ float fyy(int n, int i, int j, int k, float*yold)
{
	return yold[ (Nx + 2)*(4 + 2*Ny) + (k - 1)*(Nx + 2)*(4 + 2*Ny) + 4 +
		2*Ny + (j - 1)*(4 + 2*Ny) + 2 + (2*i - 1)-1];
}

//y-velocity
__device__ float fvy(int n, int i, int j, int k, float*yold)
{
	return yold[ (Nx + 2)*(4 + 2 * Ny) + (k - 1)*(Nx + 2)*(4 + 2 * Ny) + 4 +
		2 * Ny + (j - 1)*(4 + 2 * Ny) + 2 + (2 * i)-1];
}

//z-displacement
__device__ float fzz(int n, int i, int j, int k, float*zold)
{
	return zold[ (Nx + 2)*(4 + 2*Nz) + (i - 1)*(Nx + 2)*(4 + 2*Nz) + 4 +
		2*Nz + (j - 1)*(4 + 2*Nz) + 2 + (2*k - 1)-1];
}

//z-velocity
__device__ float fvz(int n, int i, int j, int k, float*zold)
{
	return zold[ (Nx + 2)*(4 + 2 * Nz) + (i - 1)*(Nx + 2)*(4 + 2 * Nz) + 4 +
		2 * Nz + (j - 1)*(4 + 2 * Nz) + 2 + (2 * k)-1];
}

//mass
__device__ float fm(int i, int j, int k, float*m)
{
	return m[(k - 1)*Ny*Nx + (i - 1)*Nx + j-1];
}

//x-stiffness
__device__ float fkx(int i, int j, int k, float*kx)
{
	return kx[(k - 1)*Ny*(Nx + 1) + (i - 1)*(Nx + 1) + j-1];
}

//y-stiffness
__device__ float fky(int i, int j, int k, float*ky)
{
	return ky[(k - 1)*Nx*(Ny + 1) + (i - 1)*Nx + j-1];
}

//z-stiffness
__device__ float fkz(int i, int j, int k, float*kz)
{
	return kz[(k - 1)*Nx*Ny + (i - 1)*Nx + j-1];
}

//x-damping
__device__ float fbx(int i, int j, int k, float*bx)
{
	return bx[(k - 1)*Ny*(Nx + 1) + (i - 1)*(Nx + 1) + j-1];
}

//y-damping
__device__ float fby(int i, int j, int k, float*by)
{
	return by[(k - 1)*Nx*(Ny + 1) + (i - 1)*Nx + j-1];
}

//z-damping
__device__ float fbz(int i, int j, int k, float*bz)
{
	return bz[(k - 1)*Nx*Ny + (i - 1)*Nx + j-1];
}

//x-force
__device__ float fFx(int i, int j, int k, float*Fx)
{
	return Fx[(k - 1)*Ny*Nx + (i - 1)*Nx + j-1];
}

//y-force
__device__ float fFy(int i, int j, int k, float*Fy)
{
	return Fy[(k - 1)*Ny*Nx + (i - 1)*Nx + j-1];
}

//z-force
__device__ float fFz(int i, int j, int k, float*Fz)
{
	return Fz[(k - 1)*Ny*Nx + (i - 1)*Nx + j-1];
}

//x-acceleration
__device__ float ax(int i, int j, int k, float*Fx, float*xold, float*kx, float*ky, float*kz, float*bx, float*by, float*bz, float*m)
{
	return (fFx(i, j, k, Fx) - fby(i, j, k, by)*(-fvx(1, -1 + i, j, k, xold) + fvx(1, i, j, k, xold)) -
		fbx(i, j, k, bx)*(-fvx(1, i, -1 + j, k, xold) + fvx(1, i, j, k, xold)) - fbz(i, j, k, bz)*(-fvx(1, i, j, -1 + k, xold) + fvx(1, i, j, k, xold)) +
		fbz(i, j, 1 + k, bz)*(-fvx(1, i, j, k, xold) + fvx(1, i, j, 1 + k, xold)) +
		fbx(i, 1 + j, k, bx)*(-fvx(1, i, j, k, xold) + fvx(1, i, 1 + j, k, xold)) +
		fby(1 + i, j, k, by)*(-fvx(1, i, j, k, xold) + fvx(1, 1 + i, j, k, xold)) -
		fky(i, j, k, ky)*(-fxx(1, -1 + i, j, k, xold) + fxx(1, i, j, k, xold)) - fkx(i, j, k, kx)*(-fxx(1, i, -1 + j, k, xold) + fxx(1, i, j, k, xold)) -
		fkz(i, j, k, kz)*(-fxx(1, i, j, -1 + k, xold) + fxx(1, i, j, k, xold)) +
		fkz(i, j, 1 + k, kz)*(-fxx(1, i, j, k, xold) + fxx(1, i, j, 1 + k, xold)) +
		fkx(i, 1 + j, k, kx)*(-fxx(1, i, j, k, xold) + fxx(1, i, 1 + j, k, xold)) +
		fky(1 + i, j, k, ky)*(-fxx(1, i, j, k, xold) + fxx(1, 1 + i, j, k, xold))) / fm(i, j, k, m);
}

//y-acceleration
__device__ float ay(int i, int j, int k, float*Fy, float*yold, float*kx, float*ky, float*kz, float*bx, float*by, float*bz, float*m)
{
	return (fFy(i, j, k, Fy) - fby(i, j, k, by)*(-fvy(1, -1 + i, j, k, yold) + fvy(1, i, j, k, yold)) -
		fbx(i, j, k, bx)*(-fvy(1, i, -1 + j, k, yold) + fvy(1, i, j, k, yold)) - fbz(i, j, k, bz)*(-fvy(1, i, j, -1 + k, yold) + fvy(1, i, j, k, yold)) +
		fbz(i, j, 1 + k, bz)*(-fvy(1, i, j, k, yold) + fvy(1, i, j, 1 + k, yold)) +
		fbx(i, 1 + j, k, bx)*(-fvy(1, i, j, k, yold) + fvy(1, i, 1 + j, k, yold)) +
		fby(1 + i, j, k, by)*(-fvy(1, i, j, k, yold) + fvy(1, 1 + i, j, k, yold)) -
		fky(i, j, k, ky)*(-fyy(1, -1 + i, j, k, yold) + fyy(1, i, j, k, yold)) - fkx(i, j, k, kx)*(-fyy(1, i, -1 + j, k, yold) + fyy(1, i, j, k, yold)) -
		fkz(i, j, k, kz)*(-fyy(1, i, j, -1 + k, yold) + fyy(1, i, j, k, yold)) +
		fkz(i, j, 1 + k, kz)*(-fyy(1, i, j, k, yold) + fyy(1, i, j, 1 + k, yold)) +
		fkx(i, 1 + j, k, kx)*(-fyy(1, i, j, k, yold) + fyy(1, i, 1 + j, k, yold)) +
		fky(1 + i, j, k, ky)*(-fyy(1, i, j, k, yold) + fyy(1, 1 + i, j, k, yold))) / fm(i, j, k, m);
}

//z-acceleration
__device__ float az(int i, int j, int k, float*Fz, float*zold, float*kx, float*ky, float*kz, float*bx, float*by, float*bz, float*m)
{
	return (fFz(i, j, k, Fz) - fby(i, j, k, by)*(-fvz(1, -1 + i, j, k, zold) + fvz(1, i, j, k, zold)) -
		fbx(i, j, k, bx)*(-fvz(1, i, -1 + j, k, zold) + fvz(1, i, j, k, zold)) - fbz(i, j, k, bz)*(-fvz(1, i, j, -1 + k, zold) + fvz(1, i, j, k, zold)) +
		fbz(i, j, 1 + k, bz)*(-fvz(1, i, j, k, zold) + fvz(1, i, j, 1 + k, zold)) +
		fbx(i, 1 + j, k, bx)*(-fvz(1, i, j, k, zold) + fvz(1, i, 1 + j, k, zold)) +
		fby(1 + i, j, k, by)*(-fvz(1, i, j, k, zold) + fvz(1, 1 + i, j, k, zold)) -
		fky(i, j, k, ky)*(-fzz(1, -1 + i, j, k, zold) + fzz(1, i, j, k, zold)) - fkx(i, j, k, kx)*(-fzz(1, i, -1 + j, k, zold) + fzz(1, i, j, k, zold)) -
		fkz(i, j, k, kz)*(-fzz(1, i, j, -1 + k, zold) + fzz(1, i, j, k, zold)) +
		fkz(i, j, 1 + k, kz)*(-fzz(1, i, j, k, zold) + fzz(1, i, j, 1 + k, zold)) +
		fkx(i, 1 + j, k, kx)*(-fzz(1, i, j, k, zold) + fzz(1, i, 1 + j, k, zold)) +
		fky(1 + i, j, k, ky)*(-fzz(1, i, j, k, zold) + fzz(1, 1 + i, j, k, zold))) / fm(i, j, k, m);
}











__global__ void SolveKernel(int dimBlockX,int dimBlockY,int dimBlockZ,float*xoldd,float*yoldd,float*zoldd,float*xnewd,float*ynewd,float*znewd,float*md,float*kxd,float*kyd,float*kzd,float*bxd,float*byd,float*bzd,float*Fxd,float*Fyd,float*Fzd)
{
//	int tx=threadIdx.x;
//	int ty=threadIdx.y;
	int tx=blockIdx.x*dimBlockX+threadIdx.x;
	int ty=blockIdx.y*dimBlockY+threadIdx.y;
	int tz=blockIdx.z*dimBlockZ+threadIdx.z;
	
	int i=ty+1;
	int j=tx+1;
	int k=tz+1;

	xnewd[ (Ny + 2)*(4 + 2*Nx) + (k - 1)*(Ny + 2)*(4 + 2*Nx) + 4 +
		2 * Nx + (i - 1)*(4 + 2 * Nx) + 2 + (2 * j - 1) - 1] = fxx(1, i, j, k, xoldd) + fvx(1, i, j, k, xoldd)*dt;
	xnewd[ (Ny + 2)*(4 + 2 * Nx) + (k - 1)*(Ny + 2)*(4 + 2 * Nx) + 4 +
		2 * Nx + (i - 1)*(4 + 2 * Nx) + 2 + (2 * j) - 1] = fvx(1, i, j, k, xoldd) + ax(i, j, k, Fxd, xoldd, kxd, kyd, kzd, bxd, byd, bzd, md)*dt;

	ynewd[ (Nx + 2)*(4 + 2*Ny) + (k - 1)*(Nx + 2)*(4 + 2*Ny) + 4 +
		2*Ny + (j - 1)*(4 + 2*Ny) + 2 + (2*i - 1)-1] = fyy(1, i, j, k, yoldd) + fvy(1, i, j, k, yoldd)*dt;
	ynewd[ (Nx + 2)*(4 + 2*Ny) + (k - 1)*(Nx + 2)*(4 + 2*Ny) + 4 +
		2*Ny + (j - 1)*(4 + 2*Ny) + 2 + (2*i)-1] = fvy(1, i, j, k, yoldd) + ay(i, j, k, Fyd, yoldd, kxd, kyd, kzd, bxd, byd, bzd, md)*dt;

	znewd[ (Nx + 2)*(4 + 2*Nz) + (i - 1)*(Nx + 2)*(4 + 2*Nz) + 4 +
		2*Nz + (j - 1)*(4 + 2*Nz) + 2 + (2*k - 1)-1] = fzz(1, i, j, k, zoldd) + fvz(1, i, j, k, zoldd)*dt;
	znewd[ (Nx + 2)*(4 + 2*Nz) + (i - 1)*(Nx + 2)*(4 + 2*Nz) + 4 +
		2*Nz + (j - 1)*(4 + 2*Nz) + 2 + (2*k)-1] = fvz(1, i, j, k, zoldd) + az(i, j, k, Fzd, zoldd, kxd, kyd, kzd, bxd, byd, bzd, md)*dt;
}









void Solve(float*xold,float*yold,float*zold,float*xnew,float*ynew,float*znew,float*m,float*kx,float*ky,float*kz,float*bx,float*by,float*bz,float*Fx,float*Fy,float*Fz)
{
	float *xoldd,*yoldd,*zoldd,*xnewd,*ynewd,*znewd,*md,*kxd,*kyd,*kzd,*bxd,*byd,*bzd,*Fxd,*Fyd,*Fzd;
	
	int sizexoldd=xlength*sizeof(float);
	hipMalloc((void**)&xoldd,sizexoldd);
	hipMemcpy(xoldd,xold,sizexoldd,hipMemcpyHostToDevice);
	int sizeyoldd=ylength*sizeof(float);
	hipMalloc((void**)&yoldd,sizeyoldd);
	hipMemcpy(yoldd,yold,sizeyoldd,hipMemcpyHostToDevice);
	int sizezoldd=zlength*sizeof(float);
	hipMalloc((void**)&zoldd,sizezoldd);
	hipMemcpy(zoldd,zold,sizezoldd,hipMemcpyHostToDevice);
	int sizexnewd=xlength*sizeof(float);
	hipMalloc((void**)&xnewd,sizexnewd);
	hipMemcpy(xnewd,xnew,sizexnewd,hipMemcpyHostToDevice);
	int sizeynewd=ylength*sizeof(float);
	hipMalloc((void**)&ynewd,sizeynewd);
	hipMemcpy(ynewd,ynew,sizeynewd,hipMemcpyHostToDevice);
	int sizeznewd=zlength*sizeof(float);
	hipMalloc((void**)&znewd,sizeznewd);
	hipMemcpy(znewd,znew,sizeznewd,hipMemcpyHostToDevice);
	int sizemd=masslength*sizeof(float);
	hipMalloc((void**)&md,sizemd);
	hipMemcpy(md,m,sizemd,hipMemcpyHostToDevice);
	int sizekxd=kxlength*sizeof(float);
	hipMalloc((void**)&kxd,sizekxd);
	hipMemcpy(kxd,kx,sizekxd,hipMemcpyHostToDevice);
	int sizekyd=kylength*sizeof(float);
	hipMalloc((void**)&kyd,sizekyd);
	hipMemcpy(kyd,ky,sizekyd,hipMemcpyHostToDevice);
	int sizekzd=kzlength*sizeof(float);
	hipMalloc((void**)&kzd,sizekzd);
	hipMemcpy(kzd,kz,sizekzd,hipMemcpyHostToDevice);
	int sizebxd=bxlength*sizeof(float);
	hipMalloc((void**)&bxd,sizebxd);
	hipMemcpy(bxd,bx,sizebxd,hipMemcpyHostToDevice);
	int sizebyd=bylength*sizeof(float);
	hipMalloc((void**)&byd,sizebyd);
	hipMemcpy(byd,by,sizebyd,hipMemcpyHostToDevice);
	int sizebzd=bzlength*sizeof(float);
	hipMalloc((void**)&bzd,sizebzd);
	hipMemcpy(bzd,bz,sizebzd,hipMemcpyHostToDevice);
	int sizeFxd=masslength*sizeof(float);
	hipMalloc((void**)&Fxd,sizeFxd);
	hipMemcpy(Fxd,Fx,sizeFxd,hipMemcpyHostToDevice);
	int sizeFyd=masslength*sizeof(float);
	hipMalloc((void**)&Fyd,sizeFyd);
	hipMemcpy(Fyd,Fy,sizeFyd,hipMemcpyHostToDevice);
	int sizeFzd=masslength*sizeof(float);
	hipMalloc((void**)&Fzd,sizeFzd);
	hipMemcpy(Fzd,Fz,sizeFzd,hipMemcpyHostToDevice);

	//Malloc result
	//cudaMalloc((void**)&Pd,size);
	//Dimensions of the run
	//int SubMtxWidth=SubWidth;
	int NBlockX=4;
	int NBlockY=3;
	int NBlockZ=5;
	int dimBlockX=Nx/NBlockX;
	int dimBlockY=Ny/NBlockY;
	int dimBlockZ=Nz/NBlockZ;
	dim3 dimBlock(dimBlockX,dimBlockY,dimBlockZ);
	dim3 dimGrid(NBlockX,NBlockY,NBlockZ);
	//Running Kernel
	SolveKernel<<<dimGrid,dimBlock>>>(dimBlockX,dimBlockY,dimBlockZ,xoldd,yoldd,zoldd,xnewd,ynewd,znewd,md,kxd,kyd,kzd,bxd,byd,bzd,Fxd,Fyd,Fzd);
	hipDeviceSynchronize();
	//Copy data back
	hipMemcpy(xnew,xnewd,sizexnewd,hipMemcpyDeviceToHost);
	hipMemcpy(ynew,ynewd,sizeynewd,hipMemcpyDeviceToHost);
	hipMemcpy(znew,znewd,sizeznewd,hipMemcpyDeviceToHost);
	checkCUDAError("memcpy");
	//Free memory
	//cudaFree(Md);
	//cudaFree(Nd);
	//cudaFree(Pd);
	//NEWSHIT
	hipFree(xoldd);
	hipFree(yoldd);
	hipFree(zoldd);
	hipFree(xnewd);
	hipFree(ynewd);
	hipFree(znewd);
	hipFree(md);
	hipFree(kxd);
	hipFree(kyd);
	hipFree(kzd);
	hipFree(bxd);
	hipFree(byd);
	hipFree(bzd);
	hipFree(Fxd);
	hipFree(Fyd);
	hipFree(Fzd);
}
















int main(int argc,char* argv[])
{
float *xold,*yold,*zold,*xnew,*ynew,*znew,*m,*kx,*ky,*kz,*bx,*by,*bz,*Fx,*Fy,*Fz;

//----------------------------------INITIALIZATION START----------------------------------
// Solution vectors 
xold=(float *)malloc(xlength*sizeof(float));
yold=(float *)malloc(ylength*sizeof(float));
zold=(float *)malloc(zlength*sizeof(float));
xnew=(float *)malloc(xlength*sizeof(float));
ynew=(float *)malloc(ylength*sizeof(float));
znew=(float *)malloc(zlength*sizeof(float));

// Mass vector
m=(float *)malloc(masslength*sizeof(float));

// Stiffness vectors
kx=(float *)malloc(kxlength*sizeof(float));
ky=(float *)malloc(kylength*sizeof(float));
kz=(float *)malloc(kzlength*sizeof(float));

// Damping vectors
bx=(float *)malloc(bxlength*sizeof(float));
by=(float *)malloc(bylength*sizeof(float));
bz=(float *)malloc(bzlength*sizeof(float));

// Force vectors
Fx=(float *)malloc(masslength*sizeof(float));
Fy=(float *)malloc(masslength*sizeof(float));
Fz=(float *)malloc(masslength*sizeof(float));

// Initial conditions
for (int i = 0; i < xlength ; i++)
{
	xold[i]=0.0f;
	xnew[i]=0.0f;
}

for (int i = 0; i < ylength ; i++)
{
	yold[i]=0.0f;
	ynew[i]=0.0f;
}

for (int i = 0; i < zlength ; i++)
{
	zold[i]=0.0f;
	znew[i]=0.0f;
}

// Mass [kg] and forces
for (int i = 0; i < masslength ; i++)
{
	m[i]=1.0f;
	Fx[i]=0.0f;
	Fy[i]=0.0f;
	Fz[i]=0.0f;
}

// Stiffness [N/m] and damping [N sec/m] in x-direction
for (int i = 0; i < kxlength ; i++)
{
	kx[i]=0.2f;
	bx[i]=0.05f;
}

// Stiffness [N/m] and damping [N sec/m] in y-direction
for (int i = 0; i < kylength ; i++)
{
	ky[i]=0.2f;
	by[i]=0.05f;
}

// Stiffness [N/m] and damping [N sec/m] in z-direction
for (int i = 0; i < kzlength ; i++)
{
	kz[i]=0.2f;
	bz[i]=0.05f;
}
//----------------------------------INITIALIZATION END--------------------------------------



//-------------------------------BOUNDARY CONDITIONS START----------------------------------
// No connections with Top wall B.C.'s
for (int i = 1; i <= Nx; i++)
	{
		for (int k = 1; k <= Nz; k++)
		{
			ky[i + Nx*Ny + (-1 + k)*Nx*(1 + Ny) - 1] = 0.0f;
			by[i + Nx*Ny + (-1 + k)*Nx*(1 + Ny) - 1] = 0.0f;
		}
	}
//--------------------------------BOUNDARY CONDITIONS END-----------------------------------



//--------------------------------------SOLVER START-----------------------------------------
clock_t t;
t=clock();

for (int n = 1; n <= Nstep-1; n++)
	{
		// Excitation
		Fx[(2 - 1)*Ny*Nx + (6 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu); // omega = 3 [rad/sec]
		Fy[(2 - 1)*Ny*Nx + (6 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);
		Fz[(2 - 1)*Ny*Nx + (6 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);

		Fx[(2 - 1)*Ny*Nx + (7 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);
		Fy[(2 - 1)*Ny*Nx + (7 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);
		Fz[(2 - 1)*Ny*Nx + (7 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);

		Fx[(2 - 1)*Ny*Nx + (5 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);
		Fy[(2 - 1)*Ny*Nx + (5 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);
		Fz[(2 - 1)*Ny*Nx + (5 - 1)*Nx + 8 - 1] = sin(3 * n*dtcpu);

		Solve(xold,yold,zold,xnew,ynew,znew,m,kx,ky,kz,bx,by,bz,Fx,Fy,Fz);
		hipDeviceSynchronize();
		// OLD=NEW
		for (int ix = 0; ix < xlength; ix++)
		{
			xold[ix] = xnew[ix];
		}
		for (int iy = 0; iy < ylength; iy++)
		{
			yold[iy] = ynew[iy];
		}
		for (int iz = 0; iz < zlength; iz++)
		{
			zold[iz] = znew[iz];
		}
	}
ofstream fout("test.txt");
if (fout.is_open())
	{
		//file opened successfully so we are here
		cout << "File Opened successfully!!!. Writing data from array to file" << endl;

		for (int j = 0; j < zlength; j++)
			{
				fout << znew[j] << ' '; //writing ith character of array in the file
			}
			fout << '\n';
		cout << "Array data successfully saved into the file test.txt" << endl;
	}
	else //file could not be opened
	{
		cout << "File could not be opened." << endl;
	}

t=clock()-t;
printf("%f seconds\n",((float)t)/CLOCKS_PER_SEC);
printf("%f,%f,%f\n",xold[60],yold[60],zold[60]);

free(xold);
free(yold);
free(zold);
free(xnew);
free(ynew);
free(znew);
free(m);
free(kx);
free(ky);
free(kz);
free(bx);
free(by);
free(bz);
free(Fx);
free(Fy);
free(Fz);

return 0;
}


void checkCUDAError(const char *msg)
{
hipError_t err = hipGetLastError();
if(hipSuccess!= err)
{
fprintf(stderr,"Cuda error: %s: %s.\n",msg,hipGetErrorString(err));
exit(EXIT_FAILURE);
}
}
	.text
	.file	"msd.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z26__device_stub__SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_ # -- Begin function _Z26__device_stub__SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.type	_Z26__device_stub__SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_,@function
_Z26__device_stub__SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_: # @_Z26__device_stub__SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$256, %rsp                      # imm = 0x100
	.cfi_def_cfa_offset 304
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	20(%rsp), %rax
	movl	%edi, (%rax)
	leaq	16(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	56(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	48(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	40(%rsp), %r8
	movq	%r9, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	304(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	312(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	320(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	328(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	336(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	344(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	352(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	360(%rsp), %rax
	movq	%rax, 104(%rbx)
	leaq	368(%rsp), %rax
	movq	%rax, 112(%rbx)
	leaq	376(%rsp), %rax
	movq	%rax, 120(%rbx)
	leaq	384(%rsp), %rax
	movq	%rax, 128(%rbx)
	leaq	392(%rsp), %rax
	movq	%rax, 136(%rbx)
	leaq	400(%rsp), %rax
	movq	%rax, 144(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$272, %rsp                      # imm = 0x110
	.cfi_adjust_cfa_offset -272
	popq	%rbx
	.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
	retq
.Lfunc_end0:
	.size	_Z26__device_stub__SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, .Lfunc_end0-_Z26__device_stub__SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_ # -- Begin function _Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.type	_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_,@function
_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_: # @_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$136, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %r13
	movq	%r8, %r14
	movq	%rcx, %r15
	movq	%rdx, 128(%rsp)                 # 8-byte Spill
	movq	%rsi, %rbp
	movq	%rdi, %r12
	leaq	120(%rsp), %rbx
	movl	$558272, %esi                   # imm = 0x884C0
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$558272, %edx                   # imm = 0x884C0
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	112(%rsp), %rbx
	movl	$558272, %esi                   # imm = 0x884C0
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$558272, %edx                   # imm = 0x884C0
	movq	%rbp, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	104(%rsp), %rbx
	movl	$558272, %esi                   # imm = 0x884C0
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$558272, %edx                   # imm = 0x884C0
	movq	128(%rsp), %rsi                 # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	16(%rsp), %rbx
	movl	$558272, %esi                   # imm = 0x884C0
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$558272, %edx                   # imm = 0x884C0
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	8(%rsp), %rbx
	movl	$558272, %esi                   # imm = 0x884C0
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$558272, %edx                   # imm = 0x884C0
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	%rsp, %rbx
	movl	$558272, %esi                   # imm = 0x884C0
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$558272, %edx                   # imm = 0x884C0
	movq	%r13, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	96(%rsp), %rbx
	movl	$230400, %esi                   # imm = 0x38400
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$230400, %edx                   # imm = 0x38400
	movq	192(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	88(%rsp), %rbx
	movl	$240000, %esi                   # imm = 0x3A980
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$240000, %edx                   # imm = 0x3A980
	movq	200(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	80(%rsp), %rbx
	movl	$232320, %esi                   # imm = 0x38B80
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$232320, %edx                   # imm = 0x38B80
	movq	208(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	72(%rsp), %rbx
	movl	$241920, %esi                   # imm = 0x3B100
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$241920, %edx                   # imm = 0x3B100
	movq	216(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	64(%rsp), %rbx
	movl	$240000, %esi                   # imm = 0x3A980
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$240000, %edx                   # imm = 0x3A980
	movq	224(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	56(%rsp), %rbx
	movl	$232320, %esi                   # imm = 0x38B80
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$232320, %edx                   # imm = 0x38B80
	movq	232(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	48(%rsp), %rbx
	movl	$241920, %esi                   # imm = 0x3B100
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$241920, %edx                   # imm = 0x3B100
	movq	240(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	40(%rsp), %rbx
	movl	$230400, %esi                   # imm = 0x38400
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$230400, %edx                   # imm = 0x38400
	movq	248(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	32(%rsp), %rbx
	movl	$230400, %esi                   # imm = 0x38400
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$230400, %edx                   # imm = 0x38400
	movq	256(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	24(%rsp), %rbx
	movl	$230400, %esi                   # imm = 0x38400
	movq	%rbx, %rdi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movl	$230400, %edx                   # imm = 0x38400
	movq	264(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$12884901892, %rdi              # imm = 0x300000004
	movabsq	$171798691846, %rdx             # imm = 0x2800000006
	movl	$5, %esi
	movl	$4, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	120(%rsp), %rcx
	movq	112(%rsp), %r8
	movq	104(%rsp), %r9
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	$6, %edi
	movl	$40, %esi
	movl	$4, %edx
	pushq	32(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	48(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	64(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	80(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	96(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	112(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	144(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	160(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	176(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	88(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	104(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	120(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z26__device_stub__SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	addq	$112, %rsp
	.cfi_adjust_cfa_offset -112
.LBB1_2:
	callq	hipDeviceSynchronize
	movq	16(%rsp), %rsi
	movl	$558272, %edx                   # imm = 0x884C0
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	8(%rsp), %rsi
	movl	$558272, %edx                   # imm = 0x884C0
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	(%rsp), %rsi
	movl	$558272, %edx                   # imm = 0x884C0
	movq	%r13, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.L.str, %edi
	callq	_Z14checkCUDAErrorPKc
	movq	120(%rsp), %rdi
	callq	hipFree
	movq	112(%rsp), %rdi
	callq	hipFree
	movq	104(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	movq	96(%rsp), %rdi
	callq	hipFree
	movq	88(%rsp), %rdi
	callq	hipFree
	movq	80(%rsp), %rdi
	callq	hipFree
	movq	72(%rsp), %rdi
	callq	hipFree
	movq	64(%rsp), %rdi
	callq	hipFree
	movq	56(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	addq	$136, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, .Lfunc_end1-_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z14checkCUDAErrorPKc           # -- Begin function _Z14checkCUDAErrorPKc
	.type	_Z14checkCUDAErrorPKc,@function
_Z14checkCUDAErrorPKc:                  # @_Z14checkCUDAErrorPKc
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rdi, %rbx
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB2_2
# %bb.1:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.LBB2_2:
	.cfi_def_cfa_offset 32
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movq	%r14, %rdi
	movq	%rbx, %rdx
	movq	%rax, %rcx
	xorl	%eax, %eax
	callq	fprintf
	movl	$1, %edi
	callq	exit
.Lfunc_end2:
	.size	_Z14checkCUDAErrorPKc, .Lfunc_end2-_Z14checkCUDAErrorPKc
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI3_0:
	.long	0x49742400                      # float 1.0E+6
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
.Lfunc_begin0:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception0
# %bb.0:                                # %.preheader187.preheader
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$664, %rsp                      # imm = 0x298
	.cfi_def_cfa_offset 720
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	.cfi_escape 0x2e, 0x00
	movl	$1, %edi
	movl	$558272, %esi                   # imm = 0x884C0
	callq	calloc@PLT
	movq	%rax, %rbx
	.cfi_escape 0x2e, 0x00
	movl	$1, %edi
	movl	$558272, %esi                   # imm = 0x884C0
	callq	calloc@PLT
	movq	%rax, 56(%rsp)                  # 8-byte Spill
	.cfi_escape 0x2e, 0x00
	movl	$1, %edi
	movl	$558272, %esi                   # imm = 0x884C0
	callq	calloc@PLT
	movq	%rax, %r12
	.cfi_escape 0x2e, 0x00
	movl	$1, %edi
	movl	$558272, %esi                   # imm = 0x884C0
	callq	calloc@PLT
	movq	%rax, %r14
	.cfi_escape 0x2e, 0x00
	movl	$1, %edi
	movl	$558272, %esi                   # imm = 0x884C0
	callq	calloc@PLT
	movq	%rax, %rbp
	.cfi_escape 0x2e, 0x00
	movl	$1, %edi
	movl	$558272, %esi                   # imm = 0x884C0
	callq	calloc@PLT
	movq	%rax, %r15
	.cfi_escape 0x2e, 0x00
	movl	$230400, %edi                   # imm = 0x38400
	callq	malloc
	movq	%rax, %r13
	.cfi_escape 0x2e, 0x00
	movl	$240000, %edi                   # imm = 0x3A980
	callq	malloc
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	.cfi_escape 0x2e, 0x00
	movl	$232320, %edi                   # imm = 0x38B80
	callq	malloc
	movq	%rax, 40(%rsp)                  # 8-byte Spill
	.cfi_escape 0x2e, 0x00
	movl	$241920, %edi                   # imm = 0x3B100
	callq	malloc
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	.cfi_escape 0x2e, 0x00
	movl	$240000, %edi                   # imm = 0x3A980
	callq	malloc
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	.cfi_escape 0x2e, 0x00
	movl	$232320, %edi                   # imm = 0x38B80
	callq	malloc
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	.cfi_escape 0x2e, 0x00
	movl	$241920, %edi                   # imm = 0x3B100
	callq	malloc
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	.cfi_escape 0x2e, 0x00
	movl	$1, %edi
	movl	$230400, %esi                   # imm = 0x38400
	callq	calloc@PLT
	movq	%rax, 104(%rsp)                 # 8-byte Spill
	.cfi_escape 0x2e, 0x00
	movl	$1, %edi
	movl	$230400, %esi                   # imm = 0x38400
	callq	calloc@PLT
	movq	%rax, 96(%rsp)                  # 8-byte Spill
	.cfi_escape 0x2e, 0x00
	movl	$1, %edi
	movl	$230400, %esi                   # imm = 0x38400
	callq	calloc@PLT
	xorl	%ecx, %ecx
.LBB3_1:                                # %.preheader185
                                        # =>This Inner Loop Header: Depth=1
	movl	$1065353216, (%r13,%rcx,4)      # imm = 0x3F800000
	incq	%rcx
	cmpq	$57600, %rcx                    # imm = 0xE100
	jne	.LBB3_1
# %bb.2:                                # %.preheader184.preheader
	movq	%rax, 88(%rsp)                  # 8-byte Spill
	xorl	%eax, %eax
	movq	48(%rsp), %rcx                  # 8-byte Reload
	movq	24(%rsp), %rdx                  # 8-byte Reload
.LBB3_3:                                # %.preheader184
                                        # =>This Inner Loop Header: Depth=1
	movl	$1045220557, (%rcx,%rax,4)      # imm = 0x3E4CCCCD
	movl	$1028443341, (%rdx,%rax,4)      # imm = 0x3D4CCCCD
	incq	%rax
	cmpq	$60000, %rax                    # imm = 0xEA60
	jne	.LBB3_3
# %bb.4:                                # %.preheader183.preheader
	xorl	%eax, %eax
	movq	40(%rsp), %rcx                  # 8-byte Reload
	movq	32(%rsp), %rdx                  # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	8(%rsp), %rdi                   # 8-byte Reload
.LBB3_5:                                # %.preheader183
                                        # =>This Inner Loop Header: Depth=1
	movl	$1045220557, (%rcx,%rax,4)      # imm = 0x3E4CCCCD
	movl	$1028443341, (%rsi,%rax,4)      # imm = 0x3D4CCCCD
	incq	%rax
	cmpq	$58080, %rax                    # imm = 0xE2E0
	jne	.LBB3_5
# %bb.6:                                # %.preheader182.preheader
	xorl	%eax, %eax
.LBB3_7:                                # %.preheader182
                                        # =>This Inner Loop Header: Depth=1
	movl	$1045220557, (%rdx,%rax,4)      # imm = 0x3E4CCCCD
	movl	$1028443341, (%rdi,%rax,4)      # imm = 0x3D4CCCCD
	incq	%rax
	cmpq	$60480, %rax                    # imm = 0xEC40
	jne	.LBB3_7
# %bb.8:                                # %.preheader180.preheader
	movl	$11520, %edx                    # imm = 0x2D00
	movq	%rcx, %rax
	addq	%rdx, %rax
	movq	%rsi, %rcx
	addq	%rdx, %rcx
	movl	$1, %edx
	xorl	%esi, %esi
.LBB3_9:                                # %.preheader180
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_10 Depth 2
	xorl	%edi, %edi
.LBB3_10:                               #   Parent Loop BB3_9 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%esi, (%rax,%rdi)
	movl	%esi, (%rcx,%rdi)
	addq	$11616, %rdi                    # imm = 0x2D60
	cmpq	$232320, %rdi                   # imm = 0x38B80
	jne	.LBB3_10
# %bb.11:                               #   in Loop: Header=BB3_9 Depth=1
	incq	%rdx
	addq	$4, %rax
	addq	$4, %rcx
	cmpq	$25, %rdx
	jne	.LBB3_9
# %bb.12:
	.cfi_escape 0x2e, 0x00
	movq	%r14, 136(%rsp)                 # 8-byte Spill
	movq	%r15, 72(%rsp)                  # 8-byte Spill
	movq	%r12, 64(%rsp)                  # 8-byte Spill
	movq	%rbx, 128(%rsp)                 # 8-byte Spill
	movl	$3, %ebx
	callq	clock
	movq	%rax, 144(%rsp)                 # 8-byte Spill
	movq	%rbp, 112(%rsp)                 # 8-byte Spill
	movq	%r13, 120(%rsp)                 # 8-byte Spill
	movq	104(%rsp), %r14                 # 8-byte Reload
	movq	96(%rsp), %r12                  # 8-byte Reload
	movq	128(%rsp), %rbp                 # 8-byte Reload
	movq	136(%rsp), %r13                 # 8-byte Reload
	movq	112(%rsp), %r15                 # 8-byte Reload
.LBB3_13:                               # %.preheader179.preheader
                                        # =>This Inner Loop Header: Depth=1
	movl	%ebx, 84(%rsp)                  # 4-byte Spill
	cvtsi2ss	%ebx, %xmm1
	movss	%xmm1, 4(%rsp)                  # 4-byte Spill
	movss	dtcpu(%rip), %xmm0              # xmm0 = mem[0],zero,zero,zero
	mulss	%xmm1, %xmm0
	.cfi_escape 0x2e, 0x00
	callq	sinf
	movss	%xmm0, 12028(%r14)
	movss	dtcpu(%rip), %xmm0              # xmm0 = mem[0],zero,zero,zero
	mulss	4(%rsp), %xmm0                  # 4-byte Folded Reload
	.cfi_escape 0x2e, 0x00
	callq	sinf
	movss	%xmm0, 12028(%r12)
	movss	dtcpu(%rip), %xmm0              # xmm0 = mem[0],zero,zero,zero
	mulss	4(%rsp), %xmm0                  # 4-byte Folded Reload
	.cfi_escape 0x2e, 0x00
	callq	sinf
	movq	88(%rsp), %rbx                  # 8-byte Reload
	movss	%xmm0, 12028(%rbx)
	movss	dtcpu(%rip), %xmm0              # xmm0 = mem[0],zero,zero,zero
	mulss	4(%rsp), %xmm0                  # 4-byte Folded Reload
	.cfi_escape 0x2e, 0x00
	callq	sinf
	movss	%xmm0, 12124(%r14)
	movss	dtcpu(%rip), %xmm0              # xmm0 = mem[0],zero,zero,zero
	mulss	4(%rsp), %xmm0                  # 4-byte Folded Reload
	.cfi_escape 0x2e, 0x00
	callq	sinf
	movss	%xmm0, 12124(%r12)
	movss	dtcpu(%rip), %xmm0              # xmm0 = mem[0],zero,zero,zero
	mulss	4(%rsp), %xmm0                  # 4-byte Folded Reload
	.cfi_escape 0x2e, 0x00
	callq	sinf
	movss	%xmm0, 12124(%rbx)
	movss	dtcpu(%rip), %xmm0              # xmm0 = mem[0],zero,zero,zero
	mulss	4(%rsp), %xmm0                  # 4-byte Folded Reload
	.cfi_escape 0x2e, 0x00
	callq	sinf
	movss	%xmm0, 11932(%r14)
	movss	dtcpu(%rip), %xmm0              # xmm0 = mem[0],zero,zero,zero
	mulss	4(%rsp), %xmm0                  # 4-byte Folded Reload
	.cfi_escape 0x2e, 0x00
	callq	sinf
	movss	%xmm0, 11932(%r12)
	movss	4(%rsp), %xmm0                  # 4-byte Reload
                                        # xmm0 = mem[0],zero,zero,zero
	mulss	dtcpu(%rip), %xmm0
	.cfi_escape 0x2e, 0x00
	callq	sinf
	movss	%xmm0, 11932(%rbx)
	.cfi_escape 0x2e, 0x50
	movq	%rbp, %rdi
	movq	56(%rsp), %rsi                  # 8-byte Reload
	movq	64(%rsp), %rdx                  # 8-byte Reload
	movq	%r13, %rcx
	movq	%r15, %r8
	movq	72(%rsp), %r9                   # 8-byte Reload
	pushq	%rbx
	movl	92(%rsp), %ebx                  # 4-byte Reload
	.cfi_adjust_cfa_offset 8
	pushq	%r12
	.cfi_adjust_cfa_offset 8
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	pushq	32(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	48(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	64(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	80(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	96(%rsp)                        # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	112(%rsp)                       # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	pushq	192(%rsp)                       # 8-byte Folded Reload
	.cfi_adjust_cfa_offset 8
	callq	_Z5SolvePfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
	.cfi_escape 0x2e, 0x00
	callq	hipDeviceSynchronize
	.cfi_escape 0x2e, 0x00
	movl	$558272, %edx                   # imm = 0x884C0
	movq	%rbp, %rdi
	movq	%r13, %rsi
	callq	memcpy@PLT
	.cfi_escape 0x2e, 0x00
	movl	$558272, %edx                   # imm = 0x884C0
	movq	56(%rsp), %rdi                  # 8-byte Reload
	movq	%r15, %rsi
	callq	memcpy@PLT
	.cfi_escape 0x2e, 0x00
	movl	$558272, %edx                   # imm = 0x884C0
	movq	64(%rsp), %rdi                  # 8-byte Reload
	movq	72(%rsp), %rsi                  # 8-byte Reload
	callq	memcpy@PLT
	addl	$3, %ebx
	cmpl	$3600, %ebx                     # imm = 0xE10
	jne	.LBB3_13
# %bb.14:
	movq	120(%rsp), %r13                 # 8-byte Reload
	movq	72(%rsp), %r15                  # 8-byte Reload
	movq	136(%rsp), %rbp                 # 8-byte Reload
	movq	64(%rsp), %r14                  # 8-byte Reload
	movq	128(%rsp), %r12                 # 8-byte Reload
	.cfi_escape 0x2e, 0x00
	leaq	152(%rsp), %rdi
	movl	$.L.str.1, %esi
	movl	$16, %edx
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode
	leaq	264(%rsp), %rdi
	.cfi_escape 0x2e, 0x00
	callq	_ZNKSt12__basic_fileIcE7is_openEv
	testb	%al, %al
	je	.LBB3_15
# %bb.16:
.Ltmp0:
	.cfi_escape 0x2e, 0x00
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.2, %esi
	movl	$60, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp1:
# %bb.17:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
	movq	_ZSt4cout(%rip), %rax
	addq	-24(%rax), %rbx
.Ltmp2:
	.cfi_escape 0x2e, 0x00
	movq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp3:
# %bb.18:                               # %.noexc168
.Ltmp4:
	.cfi_escape 0x2e, 0x00
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
.Ltmp5:
# %bb.19:                               # %.noexc169
.Ltmp6:
	.cfi_escape 0x2e, 0x00
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp7:
# %bb.20:                               # %_ZNSolsEPFRSoS_E.exit.preheader.preheader
	xorl	%ebx, %ebx
.LBB3_21:                               # %_ZNSolsEPFRSoS_E.exit.preheader
                                        # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r15,%rbx,4), %xmm0
.Ltmp8:
	.cfi_escape 0x2e, 0x00
	leaq	152(%rsp), %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
.Ltmp9:
# %bb.22:                               # %_ZNSolsEf.exit
                                        #   in Loop: Header=BB3_21 Depth=1
	movb	$32, 3(%rsp)
	movq	(%rax), %rcx
	movq	-24(%rcx), %rcx
	cmpq	$0, 16(%rax,%rcx)
	je	.LBB3_26
# %bb.23:                               #   in Loop: Header=BB3_21 Depth=1
.Ltmp10:
	.cfi_escape 0x2e, 0x00
	movl	$1, %edx
	movq	%rax, %rdi
	leaq	3(%rsp), %rsi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp11:
	jmp	.LBB3_27
.LBB3_26:                               #   in Loop: Header=BB3_21 Depth=1
.Ltmp12:
	.cfi_escape 0x2e, 0x00
	movq	%rax, %rdi
	movl	$32, %esi
	callq	_ZNSo3putEc
.Ltmp13:
.LBB3_27:                               # %_ZNSolsEPFRSoS_E.exit
                                        #   in Loop: Header=BB3_21 Depth=1
	incq	%rbx
	cmpq	$139568, %rbx                   # imm = 0x22130
	jne	.LBB3_21
# %bb.24:
	movb	$10, 3(%rsp)
	movq	152(%rsp), %rax
	movq	-24(%rax), %rax
	cmpq	$0, 168(%rsp,%rax)
	je	.LBB3_29
# %bb.25:
.Ltmp15:
	.cfi_escape 0x2e, 0x00
	leaq	152(%rsp), %rdi
	leaq	3(%rsp), %rsi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp16:
	movq	%r12, %rbx
	jmp	.LBB3_30
.LBB3_15:
	movl	$25, %edx
	movl	$.L.str.4, %esi
	movq	%r12, %rbx
	jmp	.LBB3_31
.LBB3_29:
.Ltmp17:
	.cfi_escape 0x2e, 0x00
	leaq	152(%rsp), %rdi
	movl	$10, %esi
	callq	_ZNSo3putEc
.Ltmp18:
	movq	%r12, %rbx
.LBB3_30:
	movl	$52, %edx
	movl	$.L.str.3, %esi
.LBB3_31:                               # %.invoke
.Ltmp19:
	.cfi_escape 0x2e, 0x00
	movl	$_ZSt4cout, %r12d
	movl	$_ZSt4cout, %edi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp20:
# %bb.32:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit165.invoke
	movq	_ZSt4cout(%rip), %rax
	addq	-24(%rax), %r12
.Ltmp21:
	.cfi_escape 0x2e, 0x00
	movq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp22:
# %bb.33:                               # %.noexc175.invoke
.Ltmp23:
	.cfi_escape 0x2e, 0x00
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
.Ltmp24:
# %bb.34:                               # %.noexc176.invoke
.Ltmp25:
	.cfi_escape 0x2e, 0x00
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp26:
# %bb.35:                               # %_ZNSolsEPFRSoS_E.exit163
	.cfi_escape 0x2e, 0x00
	callq	clock
	subq	144(%rsp), %rax                 # 8-byte Folded Reload
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rax, %xmm0
	divss	.LCPI3_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	.cfi_escape 0x2e, 0x00
	movl	$.L.str.5, %edi
	movb	$1, %al
	callq	printf
	xorps	%xmm0, %xmm0
	cvtss2sd	240(%rbx), %xmm0
	movq	56(%rsp), %r12                  # 8-byte Reload
	cvtss2sd	240(%r12), %xmm1
	cvtss2sd	240(%r14), %xmm2
	.cfi_escape 0x2e, 0x00
	movl	$.L.str.6, %edi
	movb	$3, %al
	callq	printf
	.cfi_escape 0x2e, 0x00
	movq	%rbx, %rdi
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	%r12, %rdi
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	%r14, %rdi
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	%rbp, %rdi
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	112(%rsp), %rdi                 # 8-byte Reload
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	%r15, %rdi
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	%r13, %rdi
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	48(%rsp), %rdi                  # 8-byte Reload
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	40(%rsp), %rdi                  # 8-byte Reload
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	32(%rsp), %rdi                  # 8-byte Reload
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	24(%rsp), %rdi                  # 8-byte Reload
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	16(%rsp), %rdi                  # 8-byte Reload
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	8(%rsp), %rdi                   # 8-byte Reload
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	104(%rsp), %rdi                 # 8-byte Reload
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	96(%rsp), %rdi                  # 8-byte Reload
	callq	free
	.cfi_escape 0x2e, 0x00
	movq	88(%rsp), %rdi                  # 8-byte Reload
	callq	free
	movq	_ZTTSt14basic_ofstreamIcSt11char_traitsIcEE(%rip), %rax
	leaq	160(%rsp), %rdi
	movq	%rax, -8(%rdi)
	movq	_ZTTSt14basic_ofstreamIcSt11char_traitsIcEE+24(%rip), %rcx
	movq	-24(%rax), %rax
	movq	%rcx, 152(%rsp,%rax)
	.cfi_escape 0x2e, 0x00
	callq	_ZNSt13basic_filebufIcSt11char_traitsIcEED2Ev
	leaq	400(%rsp), %rdi
	.cfi_escape 0x2e, 0x00
	callq	_ZNSt8ios_baseD2Ev
	xorl	%eax, %eax
	addq	$664, %rsp                      # imm = 0x298
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB3_36:
	.cfi_def_cfa_offset 720
.Ltmp27:
	jmp	.LBB3_37
.LBB3_28:
.Ltmp14:
.LBB3_37:
	movq	%rax, %rbx
	.cfi_escape 0x2e, 0x00
	leaq	152(%rsp), %rdi
	callq	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev
	.cfi_escape 0x2e, 0x00
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end3:
	.size	main, .Lfunc_end3-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table3:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end0-.Lcst_begin0
.Lcst_begin0:
	.uleb128 .Lfunc_begin0-.Lfunc_begin0    # >> Call Site 1 <<
	.uleb128 .Ltmp0-.Lfunc_begin0           #   Call between .Lfunc_begin0 and .Ltmp0
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp0-.Lfunc_begin0           # >> Call Site 2 <<
	.uleb128 .Ltmp7-.Ltmp0                  #   Call between .Ltmp0 and .Ltmp7
	.uleb128 .Ltmp27-.Lfunc_begin0          #     jumps to .Ltmp27
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp8-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp13-.Ltmp8                 #   Call between .Ltmp8 and .Ltmp13
	.uleb128 .Ltmp14-.Lfunc_begin0          #     jumps to .Ltmp14
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp15-.Lfunc_begin0          # >> Call Site 4 <<
	.uleb128 .Ltmp26-.Ltmp15                #   Call between .Ltmp15 and .Ltmp26
	.uleb128 .Ltmp27-.Lfunc_begin0          #     jumps to .Ltmp27
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp26-.Lfunc_begin0          # >> Call Site 5 <<
	.uleb128 .Lfunc_end3-.Ltmp26            #   Call between .Ltmp26 and .Lfunc_end3
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.p2align	2, 0x0
                                        # -- End function
	.text
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$dt, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$4, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB5_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	maxtime,@object                 # @maxtime
	.data
	.globl	maxtime
	.p2align	2, 0x0
maxtime:
	.long	0x42700000                      # float 60
	.size	maxtime, 4

	.type	dt,@object                      # @dt
	.local	dt
	.comm	dt,4,4
	.type	dtcpu,@object                   # @dtcpu
	.globl	dtcpu
	.p2align	2, 0x0
dtcpu:
	.long	0x3d4ccccd                      # float 0.0500000007
	.size	dtcpu, 4

	.type	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_,@object # @_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.section	.rodata,"a",@progbits
	.globl	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.p2align	3, 0x0
_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_:
	.quad	_Z26__device_stub__SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.size	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"memcpy"
	.size	.L.str, 7

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"test.txt"
	.size	.L.str.1, 9

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"File Opened successfully!!!. Writing data from array to file"
	.size	.L.str.2, 61

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Array data successfully saved into the file test.txt"
	.size	.L.str.3, 53

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"File could not be opened."
	.size	.L.str.4, 26

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"%f seconds\n"
	.size	.L.str.5, 12

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"%f,%f,%f\n"
	.size	.L.str.6, 10

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Cuda error: %s: %s.\n"
	.size	.L.str.7, 21

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_"
	.size	.L__unnamed_1, 51

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"dt"
	.size	.L__unnamed_2, 3

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.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 _Z26__device_stub__SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym dt
	.addrsig_sym _Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_ ; -- Begin function _Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.globl	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
	.p2align	8
	.type	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_,@function
_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_: ; @_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
; %bb.0:
	s_load_b128 s[4:7], s[0:1], 0x0
	v_bfe_u32 v1, v0, 10, 10
	v_and_b32_e32 v5, 0x3ff, v0
	v_bfe_u32 v6, v0, 20, 10
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, dt@rel32@lo+4
	s_addc_u32 s3, s3, dt@rel32@hi+12
	s_waitcnt lgkmcnt(0)
	v_mad_u64_u32 v[2:3], null, s14, s5, v[1:2]
	v_mad_u64_u32 v[3:4], null, s13, s4, v[5:6]
	v_mad_u64_u32 v[0:1], null, s15, s6, v[6:7]
	s_load_b512 s[4:19], s[0:1], 0x10
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_mul_lo_u32 v31, v2, 52
	v_add_nc_u32_e32 v36, 1, v2
	v_add_nc_u32_e32 v9, 1, v3
	v_add_nc_u32_e32 v17, 2, v3
	v_mul_lo_u32 v32, 0x18c8, v0
	v_add_nc_u32_e32 v1, 1, v0
	v_mul_lo_u32 v12, v36, 25
	v_lshlrev_b32_e32 v35, 1, v9
	v_add_nc_u32_e32 v8, 53, v31
	v_add_nc_u32_e32 v10, 54, v31
	v_add_nc_u32_e32 v40, 0x68, v31
	v_add_nc_u32_e32 v33, 0x18c8, v32
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_nc_u32_e32 v11, v8, v35
	v_add_nc_u32_e32 v13, v10, v35
	v_add3_u32 v63, v31, v32, 0x18fe
	v_add_nc_u32_e32 v19, -2, v35
	v_ashrrev_i32_e32 v18, 31, v35
	v_add_nc_u32_e32 v4, v11, v33
	v_add_nc_u32_e32 v6, v13, v33
	v_ashrrev_i32_e32 v20, 31, v63
	v_add_nc_u32_e32 v25, v33, v35
	v_add_nc_u32_e32 v34, 0x3190, v32
	v_ashrrev_i32_e32 v5, 31, v4
	v_ashrrev_i32_e32 v7, 31, v6
	v_add_nc_u32_e32 v26, v33, v10
	v_add_nc_u32_e32 v38, 2, v25
	v_or_b32_e32 v42, 1, v25
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	v_lshlrev_b64 v[6:7], 2, v[6:7]
	v_add_nc_u32_e32 v30, v32, v11
	v_add_nc_u32_e32 v22, v13, v34
	v_add_nc_u32_e32 v37, v11, v34
	v_add_nc_u32_e32 v39, v38, v31
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v4, vcc_lo, s4, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
	v_add_co_u32 v14, vcc_lo, s4, v6
	v_add_co_ci_u32_e32 v15, vcc_lo, s5, v7, vcc_lo
	v_mul_lo_u32 v7, 0xbb8, v1
	s_clause 0x1
	global_load_b32 v61, v[4:5], off offset:-4
	global_load_b32 v62, v[14:15], off offset:-4
	s_load_b512 s[36:51], s[0:1], 0x50
	s_load_b32 s0, s[2:3], 0x0
	v_mul_lo_u32 v5, v2, 24
	v_mul_lo_u32 v4, 0xb40, v0
	v_mul_lo_u32 v6, 0xb58, v1
	v_add_nc_u32_e32 v43, v42, v31
	v_add3_u32 v24, v12, v7, 0xfffff42f
	v_ashrrev_i32_e32 v31, 31, v30
	v_add_nc_u32_e32 v41, v38, v40
	v_ashrrev_i32_e32 v38, 31, v37
	v_add_nc_u32_e32 v16, 24, v5
	v_add_nc_u32_e32 v12, v24, v9
	v_add3_u32 v23, v3, v6, 0xfffff491
	v_add_co_u32 v6, vcc_lo, v35, v63
	s_delay_alu instid0(VALU_DEP_4)
	v_add3_u32 v21, v3, v16, v4
	v_add_co_ci_u32_e32 v7, vcc_lo, v18, v20, vcc_lo
	v_add_nc_u32_e32 v24, v24, v17
	v_add3_u32 v28, v5, v23, 48
	v_add_nc_u32_e32 v10, v23, v16
	v_add_nc_u32_e32 v16, v26, v19
	v_add_nc_u32_e32 v20, 0xb29, v21
	v_lshl_add_u32 v26, v17, 1, v26
	v_lshlrev_b64 v[6:7], 2, v[6:7]
	v_ashrrev_i32_e32 v25, 31, v24
	v_ashrrev_i32_e32 v11, 31, v10
	v_ashrrev_i32_e32 v23, 31, v22
	v_ashrrev_i32_e32 v27, 31, v26
	v_add_nc_u32_e32 v18, v32, v13
	v_lshlrev_b64 v[53:54], 2, v[24:25]
	v_lshlrev_b64 v[24:25], 2, v[30:31]
	v_add_co_u32 v30, vcc_lo, s10, v6
	v_ashrrev_i32_e32 v13, 31, v12
	v_add_co_ci_u32_e32 v31, vcc_lo, s11, v7, vcc_lo
	v_ashrrev_i32_e32 v17, 31, v16
	v_ashrrev_i32_e32 v29, 31, v28
	s_delay_alu instid0(VALU_DEP_4)
	v_lshlrev_b64 v[12:13], 2, v[12:13]
	v_add_nc_u32_e32 v45, v42, v40
	v_ashrrev_i32_e32 v40, 31, v39
	v_lshlrev_b64 v[49:50], 2, v[16:17]
	v_lshlrev_b64 v[28:29], 2, v[28:29]
	v_ashrrev_i32_e32 v42, 31, v41
	v_ashrrev_i32_e32 v44, 31, v43
	v_lshlrev_b64 v[39:40], 2, v[39:40]
	v_ashrrev_i32_e32 v46, 31, v45
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[41:42], 2, v[41:42]
	v_lshlrev_b64 v[43:44], 2, v[43:44]
	s_delay_alu instid0(VALU_DEP_3)
	v_lshlrev_b64 v[45:46], 2, v[45:46]
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_fmac_f32_e32 v61, s0, v62
	v_add3_u32 v4, v33, v8, v19
	v_subrev_nc_u32_e32 v8, 23, v21
	v_ashrrev_i32_e32 v21, 31, v20
	v_ashrrev_i32_e32 v19, 31, v18
	global_store_b32 v[30:31], v61, off offset:-8
	v_ashrrev_i32_e32 v5, 31, v4
	v_ashrrev_i32_e32 v9, 31, v8
	v_lshlrev_b64 v[51:52], 2, v[20:21]
	v_lshlrev_b64 v[20:21], 2, v[22:23]
	v_lshlrev_b64 v[22:23], 2, v[26:27]
	v_lshlrev_b64 v[47:48], 2, v[4:5]
	v_lshlrev_b64 v[4:5], 2, v[8:9]
	v_lshlrev_b64 v[8:9], 2, v[10:11]
	v_lshlrev_b64 v[26:27], 2, v[37:38]
	v_lshlrev_b64 v[18:19], 2, v[18:19]
	v_add_co_u32 v37, vcc_lo, s4, v47
	v_add_co_ci_u32_e32 v38, vcc_lo, s5, v48, vcc_lo
	v_add_co_u32 v47, vcc_lo, s46, v4
	v_add_co_ci_u32_e32 v48, vcc_lo, s47, v5, vcc_lo
	v_add_co_u32 v16, vcc_lo, s42, v8
	v_add_co_ci_u32_e32 v17, vcc_lo, s43, v9, vcc_lo
	v_add_co_u32 v10, vcc_lo, s40, v12
	v_add_co_ci_u32_e32 v11, vcc_lo, s41, v13, vcc_lo
	v_add_co_u32 v49, vcc_lo, s4, v49
	v_add_co_ci_u32_e32 v50, vcc_lo, s5, v50, vcc_lo
	v_add_co_u32 v6, vcc_lo, s44, v4
	v_add_co_ci_u32_e32 v7, vcc_lo, s45, v5, vcc_lo
	v_add_co_u32 v55, vcc_lo, s4, v18
	v_add_co_ci_u32_e32 v56, vcc_lo, s5, v19, vcc_lo
	v_add_co_u32 v18, vcc_lo, s44, v51
	v_add_co_ci_u32_e32 v19, vcc_lo, s45, v52, vcc_lo
	v_add_co_u32 v57, vcc_lo, s4, v20
	v_add_co_ci_u32_e32 v58, vcc_lo, s5, v21, vcc_lo
	v_add_co_u32 v20, vcc_lo, s40, v53
	v_add_co_ci_u32_e32 v21, vcc_lo, s41, v54, vcc_lo
	v_add_co_u32 v59, vcc_lo, s4, v22
	v_add_co_ci_u32_e32 v60, vcc_lo, s5, v23, vcc_lo
	v_add_co_u32 v22, vcc_lo, s42, v28
	v_add_co_ci_u32_e32 v23, vcc_lo, s43, v29, vcc_lo
	v_add_co_u32 v8, vcc_lo, s36, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s37, v9, vcc_lo
	v_add_co_u32 v39, vcc_lo, s4, v39
	v_add_co_ci_u32_e32 v40, vcc_lo, s5, v40, vcc_lo
	s_clause 0x1
	global_load_b64 v[61:62], v[14:15], off offset:-4
	global_load_b32 v64, v[39:40], off offset:-4
	global_load_b32 v47, v[47:48], off offset:-4
	global_load_b32 v48, v[16:17], off offset:-4
	global_load_b64 v[39:40], v[49:50], off offset:-4
	v_add_co_u32 v12, vcc_lo, s18, v12
	v_add_co_ci_u32_e32 v13, vcc_lo, s19, v13, vcc_lo
	global_load_b32 v49, v[55:56], off offset:-4
	global_load_b32 v50, v[10:11], off offset:-4
	global_load_b32 v55, v[6:7], off offset:-4
	global_load_b32 v56, v[57:58], off offset:-4
	v_add_co_u32 v14, vcc_lo, s38, v4
	v_add_co_ci_u32_e32 v15, vcc_lo, s39, v5, vcc_lo
	v_add_co_u32 v30, vcc_lo, s4, v24
	v_add_co_ci_u32_e32 v31, vcc_lo, s5, v25, vcc_lo
	v_add_co_u32 v24, vcc_lo, s4, v41
	global_load_b32 v57, v[18:19], off offset:-4
	global_load_b32 v58, v[59:60], off offset:-4
	v_add_co_ci_u32_e32 v25, vcc_lo, s5, v42, vcc_lo
	v_add_co_u32 v41, vcc_lo, s4, v43
	v_add_co_ci_u32_e32 v42, vcc_lo, s5, v44, vcc_lo
	global_load_b32 v43, v[20:21], off offset:-4
	global_load_b32 v44, v[24:25], off offset:-4
	v_add_co_u32 v24, vcc_lo, s38, v51
	v_add_co_ci_u32_e32 v25, vcc_lo, s39, v52, vcc_lo
	global_load_b32 v41, v[41:42], off offset:-4
	global_load_b32 v42, v[22:23], off offset:-4
	global_load_b32 v51, v[8:9], off offset:-4
	global_load_b32 v52, v[37:38], off offset:-4
	v_add_co_u32 v37, vcc_lo, s4, v26
	global_load_b32 v59, v[12:13], off offset:-4
	global_load_b32 v60, v[30:31], off offset:-4
	v_add_co_ci_u32_e32 v38, vcc_lo, s5, v27, vcc_lo
	v_add_co_u32 v26, vcc_lo, s18, v53
	global_load_b32 v53, v[14:15], off offset:-4
	global_load_b32 v37, v[37:38], off offset:-4
	v_add_co_ci_u32_e32 v27, vcc_lo, s19, v54, vcc_lo
	v_add_co_u32 v30, vcc_lo, s4, v45
	v_add_co_ci_u32_e32 v31, vcc_lo, s5, v46, vcc_lo
	global_load_b32 v38, v[24:25], off offset:-4
	v_add_co_u32 v28, vcc_lo, s36, v28
	v_add_co_ci_u32_e32 v29, vcc_lo, s37, v29, vcc_lo
	global_load_b32 v45, v[30:31], off offset:-4
	global_load_b32 v46, v[26:27], off offset:-4
	v_add_co_u32 v30, vcc_lo, s16, v4
	v_add_co_ci_u32_e32 v31, vcc_lo, s17, v5, vcc_lo
	global_load_b32 v54, v[28:29], off offset:-4
	global_load_b32 v65, v[30:31], off offset:-4
	s_waitcnt vmcnt(24)
	v_sub_f32_e32 v64, v61, v64
	s_waitcnt vmcnt(21)
	v_sub_f32_e32 v39, v61, v39
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_fma_f32 v47, -v48, v64, v47
	s_waitcnt vmcnt(20)
	v_sub_f32_e32 v48, v61, v49
	s_waitcnt vmcnt(19)
	v_fma_f32 v39, -v50, v39, v47
	s_waitcnt vmcnt(17)
	v_sub_f32_e32 v47, v56, v61
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f32 v39, -v55, v48, v39
	s_waitcnt vmcnt(15)
	v_dual_sub_f32 v48, v58, v61 :: v_dual_fmac_f32 v39, v57, v47
	s_waitcnt vmcnt(13)
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_sub_f32 v44, v44, v61 :: v_dual_fmac_f32 v39, v43, v48
	s_waitcnt vmcnt(12)
	v_sub_f32_e32 v41, v40, v41
	v_mul_lo_u32 v43, 0xf4, v3
	v_mul_lo_u32 v3, v3, 44
	s_waitcnt vmcnt(11)
	v_fmac_f32_e32 v39, v42, v44
	s_waitcnt vmcnt(9)
	v_dual_sub_f32 v42, v40, v52 :: v_dual_add_nc_u32 v35, v63, v35
	s_delay_alu instid0(VALU_DEP_2)
	v_fma_f32 v39, -v51, v41, v39
	s_waitcnt vmcnt(7)
	v_sub_f32_e32 v41, v40, v60
	s_waitcnt vmcnt(5)
	v_sub_f32_e32 v37, v37, v40
	v_fma_f32 v39, -v59, v42, v39
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f32 v42, -v53, v41, v39
	v_sub_f32_e32 v39, v62, v40
	v_lshlrev_b32_e32 v62, 1, v36
	v_ashrrev_i32_e32 v36, 31, v35
	s_waitcnt vmcnt(3)
	v_dual_fmac_f32 v42, v38, v37 :: v_dual_sub_f32 v37, v45, v40
	v_add_nc_u32_e32 v45, 0xf4, v43
	v_add_nc_u32_e32 v48, 2, v62
	v_lshlrev_b64 v[35:36], 2, v[35:36]
	s_waitcnt vmcnt(2)
	v_fmac_f32_e32 v42, v46, v39
	v_or_b32_e32 v46, 1, v62
	v_add_nc_u32_e32 v50, v48, v45
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v42, v54, v37
	v_add_nc_u32_e32 v49, v46, v45
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v39, v50, v33
	s_waitcnt vmcnt(0)
	v_div_scale_f32 v40, null, v65, v65, v42
	v_div_scale_f32 v44, vcc_lo, v42, v65, v42
	v_add_nc_u32_e32 v54, v49, v34
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_rcp_f32_e32 v41, v40
	v_add_nc_u32_e32 v52, v32, v49
	v_ashrrev_i32_e32 v55, 31, v54
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v53, 31, v52
	v_lshlrev_b64 v[54:55], 2, v[54:55]
	s_waitcnt_depctr 0xfff
	v_fma_f32 v37, -v40, v41, 1.0
	v_lshlrev_b64 v[52:53], 2, v[52:53]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v41, v37, v41
	v_add_nc_u32_e32 v37, v49, v33
	v_mul_f32_e32 v47, v44, v41
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v38, -v40, v47, v44
	v_fmac_f32_e32 v47, v38, v41
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v38, 31, v37
	v_fma_f32 v44, -v40, v47, v44
	v_ashrrev_i32_e32 v40, 31, v39
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[37:38], 2, v[37:38]
	v_div_fmas_f32 v44, v44, v41, v47
	v_add_co_u32 v35, vcc_lo, s10, v35
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[40:41], 2, v[39:40]
	v_add_co_ci_u32_e32 v36, vcc_lo, s11, v36, vcc_lo
	v_div_fixup_f32 v42, v44, v65, v42
	v_add_co_u32 v37, vcc_lo, s6, v37
	v_add_co_ci_u32_e32 v38, vcc_lo, s7, v38, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_fmac_f32_e32 v61, s0, v42
	v_add_co_u32 v40, vcc_lo, s6, v40
	v_add_co_ci_u32_e32 v41, vcc_lo, s7, v41, vcc_lo
	global_store_b32 v[35:36], v61, off offset:-4
	s_clause 0x1
	global_load_b32 v63, v[37:38], off offset:-4
	global_load_b32 v64, v[40:41], off offset:-4
	v_add3_u32 v65, v43, v32, 0x19be
	v_add_nc_u32_e32 v36, -1, v62
	v_add_nc_u32_e32 v37, v33, v45
	v_add_nc_u32_e32 v45, v33, v48
	v_add_nc_u32_e32 v47, v33, v46
	v_add_nc_u32_e32 v38, 0x1e8, v43
	v_add_nc_u32_e32 v35, v36, v65
	v_add_nc_u32_e32 v33, v37, v62
	v_add_nc_u32_e32 v37, v36, v37
	v_add_nc_u32_e32 v46, v50, v34
	v_add_nc_u32_e32 v42, v45, v43
	v_ashrrev_i32_e32 v36, 31, v35
	v_ashrrev_i32_e32 v34, 31, v33
	v_add_nc_u32_e32 v48, v45, v38
	v_add_nc_u32_e32 v56, v47, v38
	v_ashrrev_i32_e32 v38, 31, v37
	v_lshlrev_b64 v[35:36], 2, v[35:36]
	v_add_nc_u32_e32 v44, v32, v50
	v_lshlrev_b64 v[58:59], 2, v[33:34]
	v_add_nc_u32_e32 v50, v47, v43
	v_ashrrev_i32_e32 v43, 31, v42
	v_lshlrev_b64 v[37:38], 2, v[37:38]
	v_add_co_u32 v34, vcc_lo, s12, v35
	v_add_co_ci_u32_e32 v35, vcc_lo, s13, v36, vcc_lo
	v_add_co_u32 v60, vcc_lo, s48, v4
	v_ashrrev_i32_e32 v45, 31, v44
	v_add_co_ci_u32_e32 v61, vcc_lo, s49, v5, vcc_lo
	v_lshlrev_b64 v[42:43], 2, v[42:43]
	v_add_co_u32 v58, vcc_lo, s6, v58
	v_ashrrev_i32_e32 v47, 31, v46
	v_add_co_ci_u32_e32 v59, vcc_lo, s7, v59, vcc_lo
	v_lshlrev_b64 v[44:45], 2, v[44:45]
	v_add_co_u32 v36, vcc_lo, s6, v37
	v_add_co_ci_u32_e32 v37, vcc_lo, s7, v38, vcc_lo
	v_lshlrev_b64 v[46:47], 2, v[46:47]
	v_add_co_u32 v38, vcc_lo, s6, v42
	v_add_nc_u32_e32 v32, 2, v39
	v_add_co_ci_u32_e32 v39, vcc_lo, s7, v43, vcc_lo
	v_add_co_u32 v42, vcc_lo, s6, v44
	v_ashrrev_i32_e32 v49, 31, v48
	v_add_co_ci_u32_e32 v43, vcc_lo, s7, v45, vcc_lo
	v_add_co_u32 v44, vcc_lo, s6, v46
	v_add_co_ci_u32_e32 v45, vcc_lo, s7, v47, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[48:49], 2, v[48:49]
	v_ashrrev_i32_e32 v33, 31, v32
	v_ashrrev_i32_e32 v51, 31, v50
	v_ashrrev_i32_e32 v57, 31, v56
	v_lshlrev_b64 v[32:33], 2, v[32:33]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[50:51], 2, v[50:51]
	v_lshlrev_b64 v[56:57], 2, v[56:57]
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v63, s0, v64
	global_store_b32 v[34:35], v63, off offset:-4
	s_clause 0x1
	global_load_b64 v[34:35], v[40:41], off offset:-4
	global_load_b64 v[40:41], v[58:59], off offset:-4
	global_load_b32 v46, v[60:61], off offset:-4
	global_load_b32 v47, v[38:39], off offset:-4
	global_load_b32 v58, v[16:17], off offset:-4
	global_load_b32 v59, v[10:11], off offset:-4
	global_load_b32 v42, v[42:43], off offset:-4
	v_add_co_u32 v38, vcc_lo, s6, v48
	global_load_b32 v43, v[6:7], off offset:-4
	global_load_b32 v44, v[44:45], off offset:-4
	v_add_co_ci_u32_e32 v39, vcc_lo, s7, v49, vcc_lo
	v_add_co_u32 v32, vcc_lo, s6, v32
	global_load_b32 v45, v[18:19], off offset:-4
	global_load_b32 v48, v[38:39], off offset:-4
	v_add_co_ci_u32_e32 v33, vcc_lo, s7, v33, vcc_lo
	v_add_co_u32 v38, vcc_lo, s6, v50
	v_add_co_ci_u32_e32 v39, vcc_lo, s7, v51, vcc_lo
	global_load_b32 v49, v[32:33], off offset:-4
	global_load_b32 v50, v[20:21], off offset:-4
	global_load_b32 v51, v[22:23], off offset:-4
	global_load_b32 v60, v[36:37], off offset:-4
	v_add_co_u32 v32, vcc_lo, s6, v52
	global_load_b32 v52, v[8:9], off offset:-4
	global_load_b32 v38, v[38:39], off offset:-4
	v_add_co_ci_u32_e32 v33, vcc_lo, s7, v53, vcc_lo
	global_load_b32 v39, v[12:13], off offset:-4
	global_load_b32 v53, v[32:33], off offset:-4
	v_add_co_u32 v36, vcc_lo, s6, v54
	v_add_co_ci_u32_e32 v37, vcc_lo, s7, v55, vcc_lo
	v_add_co_u32 v32, vcc_lo, s6, v56
	v_add_co_ci_u32_e32 v33, vcc_lo, s7, v57, vcc_lo
	global_load_b32 v36, v[36:37], off offset:-4
	global_load_b32 v37, v[14:15], off offset:-4
	global_load_b32 v32, v[32:33], off offset:-4
	global_load_b32 v33, v[24:25], off offset:-4
	global_load_b32 v54, v[26:27], off offset:-4
	global_load_b32 v55, v[28:29], off offset:-4
	global_load_b32 v56, v[30:31], off offset:-4
	s_waitcnt vmcnt(22)
	v_sub_f32_e32 v47, v34, v47
	v_sub_f32_e32 v40, v34, v40
	s_waitcnt vmcnt(19)
	v_sub_f32_e32 v42, v34, v42
	s_waitcnt vmcnt(17)
	v_sub_f32_e32 v44, v44, v34
	v_fma_f32 v40, -v58, v40, v46
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v40, -v59, v47, v40
	v_fma_f32 v40, -v43, v42, v40
	s_waitcnt vmcnt(15)
	v_sub_f32_e32 v42, v48, v34
	s_waitcnt vmcnt(14)
	v_sub_f32_e32 v43, v49, v34
	s_waitcnt vmcnt(9)
	v_sub_f32_e32 v38, v41, v38
	v_fmac_f32_e32 v40, v45, v44
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v40, v50, v42
	v_sub_f32_e32 v42, v41, v60
	s_waitcnt vmcnt(6)
	v_sub_f32_e32 v36, v36, v41
	v_fmac_f32_e32 v40, v51, v43
	s_waitcnt vmcnt(4)
	v_sub_f32_e32 v32, v32, v41
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f32 v40, -v52, v42, v40
	v_sub_f32_e32 v42, v41, v53
	v_fma_f32 v38, -v39, v38, v40
	v_add_nc_u32_e32 v40, 44, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f32 v37, -v37, v42, v38
	v_mul_lo_u32 v42, 0x478, v2
	s_waitcnt vmcnt(3)
	v_fmac_f32_e32 v37, v33, v36
	v_sub_f32_e32 v33, v35, v41
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v43, 0x479, v42
	s_waitcnt vmcnt(2)
	v_dual_fmac_f32 v37, v54, v32 :: v_dual_add_nc_u32 v44, 0x47a, v42
	v_add3_u32 v57, v3, v42, 0x4a6
	v_add_nc_u32_e32 v48, 0x8f0, v42
	v_add_nc_u32_e32 v46, v43, v40
	s_waitcnt vmcnt(1)
	v_fmac_f32_e32 v37, v55, v33
	v_lshlrev_b32_e32 v55, 1, v1
	v_add_nc_u32_e32 v1, v65, v62
	v_add_nc_u32_e32 v47, v44, v40
	s_waitcnt vmcnt(0)
	v_div_scale_f32 v38, null, v56, v56, v37
	v_div_scale_f32 v41, vcc_lo, v37, v56, v37
	v_add_nc_u32_e32 v32, v46, v55
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_rcp_f32_e32 v39, v38
	v_add_nc_u32_e32 v35, v47, v55
	v_add_nc_u32_e32 v40, v55, v40
	v_ashrrev_i32_e32 v36, 31, v35
	s_delay_alu instid0(VALU_DEP_2)
	v_add_nc_u32_e32 v49, 2, v40
	v_or_b32_e32 v50, 1, v40
	s_waitcnt_depctr 0xfff
	v_fma_f32 v2, -v38, v39, 1.0
	v_lshlrev_b64 v[35:36], 2, v[35:36]
	v_add_nc_u32_e32 v53, v50, v48
	v_add_nc_u32_e32 v51, v50, v42
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v39, v2, v39
	v_ashrrev_i32_e32 v2, 31, v1
	v_ashrrev_i32_e32 v54, 31, v53
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_ashrrev_i32_e32 v52, 31, v51
	v_mul_f32_e32 v45, v41, v39
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_lshlrev_b64 v[51:52], 2, v[51:52]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v33, -v38, v45, v41
	v_fmac_f32_e32 v45, v33, v39
	v_ashrrev_i32_e32 v33, 31, v32
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v38, -v38, v45, v41
	v_lshlrev_b64 v[32:33], 2, v[32:33]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_div_fmas_f32 v38, v38, v39, v45
	v_add_co_u32 v1, vcc_lo, s12, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s13, v2, vcc_lo
	v_div_fixup_f32 v37, v38, v56, v37
	v_add_co_u32 v32, vcc_lo, s8, v32
	v_add_co_ci_u32_e32 v33, vcc_lo, s9, v33, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_fmac_f32_e32 v34, s0, v37
	v_add_co_u32 v35, vcc_lo, s8, v35
	v_add_co_ci_u32_e32 v36, vcc_lo, s9, v36, vcc_lo
	global_store_b32 v[1:2], v34, off offset:-4
	s_clause 0x1
	global_load_b32 v34, v[32:33], off offset:-4
	global_load_b32 v56, v[35:36], off offset:-4
	v_ashrrev_i32_e32 v1, 31, v55
	v_ashrrev_i32_e32 v2, 31, v57
	v_add_nc_u32_e32 v39, v55, v44
	v_lshlrev_b32_e32 v33, 1, v0
	v_add_co_u32 v0, vcc_lo, v55, v57
	v_add_nc_u32_e32 v45, -2, v55
	v_add_co_ci_u32_e32 v1, vcc_lo, v1, v2, vcc_lo
	v_add_nc_u32_e32 v44, v55, v43
	v_add_nc_u32_e32 v2, v39, v3
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_nc_u32_e32 v32, v45, v47
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	v_add_nc_u32_e32 v38, 0x58, v3
	v_add_nc_u32_e32 v41, v44, v3
	v_ashrrev_i32_e32 v3, 31, v2
	v_add3_u32 v37, v33, v47, 4
	v_add_nc_u32_e32 v47, v49, v42
	v_ashrrev_i32_e32 v33, 31, v32
	v_add_co_u32 v0, vcc_lo, s14, v0
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	v_add_nc_u32_e32 v49, v49, v48
	v_ashrrev_i32_e32 v48, 31, v47
	v_add_co_ci_u32_e32 v1, vcc_lo, s15, v1, vcc_lo
	v_lshlrev_b64 v[32:33], 2, v[32:33]
	v_add_co_u32 v4, vcc_lo, s50, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s51, v5, vcc_lo
	v_lshlrev_b64 v[47:48], 2, v[47:48]
	v_add_co_u32 v2, vcc_lo, s8, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v3, vcc_lo
	v_add_co_u32 v32, vcc_lo, s8, v32
	v_add_co_ci_u32_e32 v33, vcc_lo, s9, v33, vcc_lo
	v_add_co_u32 v47, vcc_lo, s8, v47
	v_add_co_ci_u32_e32 v48, vcc_lo, s9, v48, vcc_lo
	v_add_nc_u32_e32 v43, v45, v46
	v_add_nc_u32_e32 v45, v44, v38
	v_ashrrev_i32_e32 v50, 31, v49
	v_ashrrev_i32_e32 v42, 31, v41
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_ashrrev_i32_e32 v44, 31, v43
	v_ashrrev_i32_e32 v46, 31, v45
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[49:50], 2, v[49:50]
	v_lshlrev_b64 v[41:42], 2, v[41:42]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[43:44], 2, v[43:44]
	v_lshlrev_b64 v[45:46], 2, v[45:46]
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v34, s0, v56
	global_store_b32 v[0:1], v34, off offset:-8
	s_clause 0x1
	global_load_b64 v[0:1], v[35:36], off offset:-4
	global_load_b32 v34, v[47:48], off offset:-4
	global_load_b32 v35, v[4:5], off offset:-4
	global_load_b32 v16, v[16:17], off offset:-4
	global_load_b32 v17, v[2:3], off offset:-4
	v_add_nc_u32_e32 v39, v39, v38
	v_ashrrev_i32_e32 v38, 31, v37
	global_load_b32 v36, v[10:11], off offset:-4
	global_load_b64 v[4:5], v[32:33], off offset:-4
	v_ashrrev_i32_e32 v40, 31, v39
	v_lshlrev_b64 v[37:38], 2, v[37:38]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[39:40], 2, v[39:40]
	v_add_co_u32 v37, vcc_lo, s8, v37
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v38, vcc_lo, s9, v38, vcc_lo
	v_add_co_u32 v2, vcc_lo, s8, v39
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v40, vcc_lo
	global_load_b32 v32, v[37:38], off offset:-4
	s_clause 0x1
	global_load_b32 v33, v[6:7], off offset:-4
	global_load_b32 v18, v[18:19], off offset:-4
	global_load_b32 v19, v[2:3], off offset:-4
	v_add_co_u32 v10, vcc_lo, s8, v49
	v_add_co_ci_u32_e32 v11, vcc_lo, s9, v50, vcc_lo
	v_add_co_u32 v2, vcc_lo, s8, v51
	global_load_b32 v20, v[20:21], off offset:-4
	global_load_b32 v10, v[10:11], off offset:-4
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v52, vcc_lo
	v_add_co_u32 v6, vcc_lo, s8, v41
	v_add_co_ci_u32_e32 v7, vcc_lo, s9, v42, vcc_lo
	global_load_b32 v11, v[22:23], off offset:-4
	global_load_b32 v21, v[2:3], off offset:-4
	v_add_co_u32 v2, vcc_lo, s8, v43
	global_load_b32 v22, v[6:7], off offset:-4
	global_load_b32 v8, v[8:9], off offset:-4
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v44, vcc_lo
	global_load_b32 v9, v[12:13], off offset:-4
	global_load_b32 v12, v[2:3], off offset:-4
	v_lshlrev_b64 v[6:7], 2, v[53:54]
	global_load_b32 v13, v[14:15], off offset:-4
	v_add_co_u32 v2, vcc_lo, s8, v45
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v46, vcc_lo
	v_add_co_u32 v6, vcc_lo, s8, v6
	v_add_co_ci_u32_e32 v7, vcc_lo, s9, v7, vcc_lo
	global_load_b32 v2, v[2:3], off offset:-4
	global_load_b32 v3, v[24:25], off offset:-4
	global_load_b32 v6, v[6:7], off offset:-4
	global_load_b32 v7, v[26:27], off offset:-4
	global_load_b32 v14, v[28:29], off offset:-4
	global_load_b32 v15, v[30:31], off offset:-4
	s_waitcnt vmcnt(24)
	v_sub_f32_e32 v23, v0, v34
	s_waitcnt vmcnt(21)
	v_sub_f32_e32 v17, v0, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f32 v16, -v16, v23, v35
	s_waitcnt vmcnt(19)
	v_dual_sub_f32 v4, v0, v4 :: v_dual_sub_f32 v1, v1, v5
	v_fma_f32 v16, -v36, v17, v16
	s_waitcnt vmcnt(18)
	v_sub_f32_e32 v17, v32, v0
	s_waitcnt vmcnt(17)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f32 v4, -v33, v4, v16
	s_waitcnt vmcnt(15)
	v_sub_f32_e32 v16, v19, v0
	v_fmac_f32_e32 v4, v18, v17
	s_waitcnt vmcnt(13)
	v_sub_f32_e32 v10, v10, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v4, v20, v16
	s_waitcnt vmcnt(12)
	v_fmac_f32_e32 v4, v11, v10
	s_waitcnt vmcnt(11)
	v_sub_f32_e32 v16, v5, v21
	s_waitcnt vmcnt(10)
	v_sub_f32_e32 v10, v5, v22
	s_waitcnt vmcnt(9)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f32 v4, -v8, v16, v4
	s_waitcnt vmcnt(7)
	v_sub_f32_e32 v8, v5, v12
	v_fma_f32 v4, -v9, v10, v4
	s_waitcnt vmcnt(6)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_fma_f32 v4, -v13, v8, v4
	s_waitcnt vmcnt(5)
	v_sub_f32_e32 v2, v2, v5
	s_waitcnt vmcnt(4)
	v_fmac_f32_e32 v4, v3, v1
	s_waitcnt vmcnt(2)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v1, v6, v5 :: v_dual_fmac_f32 v4, v7, v2
	s_waitcnt vmcnt(1)
	v_fmac_f32_e32 v4, v14, v1
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v1, null, v15, v15, v4
	v_div_scale_f32 v5, vcc_lo, v4, v15, v4
	v_rcp_f32_e32 v2, v1
	s_waitcnt_depctr 0xfff
	v_fma_f32 v3, -v1, v2, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v2, v3, v2
	v_mul_f32_e32 v3, v5, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v6, -v1, v3, v5
	v_fmac_f32_e32 v3, v6, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f32 v5, -v1, v3, v5
	v_add_nc_u32_e32 v1, v57, v55
	v_div_fmas_f32 v3, v5, v2, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v2, 31, v1
	v_div_fixup_f32 v3, v3, v15, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_fmac_f32_e32 v0, s0, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v1, vcc_lo, s14, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s15, v2, vcc_lo
	global_store_b32 v[1:2], v0, off offset:-4
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 144
		.amdhsa_user_sgpr_count 13
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 66
		.amdhsa_next_free_sgpr 52
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_, .Lfunc_end0-_Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 3880
; NumSgprs: 54
; NumVgprs: 66
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 6
; VGPRBlocks: 8
; NumSGPRsForWavesPerEU: 54
; NumVGPRsForWavesPerEU: 66
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	dt                      ; @dt
	.type	dt,@object
	.data
	.globl	dt
	.p2align	2, 0x0
dt:
	.long	0x3d4ccccd                      ; float 0.0500000007
	.size	dt, 4

	.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 dt
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         80
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         88
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         96
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         104
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         112
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         120
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         128
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         136
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 144
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     54
    .sgpr_spill_count: 0
    .symbol:         _Z11SolveKerneliiiPfS_S_S_S_S_S_S_S_S_S_S_S_S_S_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     66
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
