/*

	Implement your CUDA kernel in this file

*/



/* 
 * Solves the Panfilov model using an explicit numerical scheme.
 * Based on code orginally provided by Xing Cai, Simula Research Laboratory 
 * and reimplementation by Scott B. Baden, UCSD
 * 
 * Modified and  restructured by Didem Unat, Koc University
 *
 * Refer to "Detailed Numerical Analyses of the Aliev-Panfilov Model on GPGPU"
 * https://www.simula.no/publications/detailed-numerical-analyses-aliev-panfilov-model-gpgpu
 * by Xing Cai, Didem Unat and Scott Baden
 *
*/

#include <stdio.h>
#include <assert.h>
#include <stdlib.h>
#include <iostream>
#include <iomanip>
#include <string.h>
#include <math.h>
#include <sys/time.h>
#include <getopt.h>

using namespace std;

#define TILE_DIM 16


// External functions
//extern "C" void splot(double **E, double T, int niter, int m, int n);

void cmdLine(int argc, char *argv[], double &T, int &n, int &px, int &py, int &plot_freq, int &no_comm, int &num_threads){
	/// Command line arguments
	// Default value of the domain sizes
	static struct option long_options[] = {
		{"n", required_argument, 0, 'n'},
		{"px", required_argument, 0, 'x'},
		{"py", required_argument, 0, 'y'},
		{"tfinal", required_argument, 0, 't'},
		{"plot", required_argument, 0, 'p'},
		{"nocomm", no_argument, 0, 'k'},
		{"numthreads", required_argument, 0, 'o'},
	};
	// Process command line arguments
	int ac;
	for (ac = 1; ac < argc; ac++)
	{
		int c;
		while ((c = getopt_long(argc, argv, "n:x:y:t:kp:o:", long_options, NULL)) != -1)
		{
			switch (c)
			{

				// Size of the computational box
			case 'n':
				n = atoi(optarg);
				break;

				// X processor geometry
			case 'x':
				px = atoi(optarg);

				// Y processor geometry
			case 'y':
				py = atoi(optarg);

				// Length of simulation, in simulated time units
			case 't':
				T = atof(optarg);
				break;
				// Turn off communication
			case 'k':
				no_comm = 1;
				break;

				// Plot the excitation variable
			case 'p':
				plot_freq = atoi(optarg);
				break;

				// Plot the excitation variable
			case 'o':
				num_threads = atoi(optarg);
				break;

				// Error
			default:
				printf("Usage: a.out [-n <domain size>] [-t <final time >]\n\t [-p <plot frequency>]\n\t[-px <x processor geometry> [-py <y proc. geometry] [-k turn off communication] [-o <Number of OpenMP threads>]\n");
				exit(-1);
			}
		}
	}
}

// Utilities
//

// Timer
// Make successive calls and take a difference to get the elapsed time.
static const double kMicro = 1.0e-6;

double getTime(){
	struct timeval TV;
	struct timezone TZ;

	const int RC = gettimeofday(&TV, &TZ);
	if (RC == -1)
	{
		cerr << "ERROR: Bad call to gettimeofday" << endl;
		return (-1);
	}

	return (((double)TV.tv_sec) + kMicro * ((double)TV.tv_usec));

} 

// end getTime()

// Allocate a 2D array
double **alloc2D(int m, int n){
	double **E;
	int nx = n, ny = m;
	E = (double **)malloc(sizeof(double *) * ny + sizeof(double) * nx * ny);
	assert(E);
	int row;
	for (row = 0; row < ny; row++)
		E[row] = (double *)(E + ny) + row * nx;
	return (E);
}

// Reports statistics about the computation
// These values should not vary (except to within roundoff)
// when we use different numbers of  processes to solve the problem
double stats(double **E, int m, int n, double *_mx){
	double mx = -1;
	double l2norm = 0;
	int col, row;
	for (row = 1; row <= m; row++)
		for (col = 1; col <= n; col++)
		{
			l2norm += E[row][col] * E[row][col];
			if (E[row][col] > mx)
				mx = E[row][col];
		}
	*_mx = mx;
	l2norm /= (double)((m) * (n));
	l2norm = sqrt(l2norm);
	return l2norm;
}

double stats1D(double *E, int m, int n, double *_mx, int WIDTH){
	double mx = -1;
	double l2norm = 0;
	int col, row;
	int index = -1;
	for (row = 1; row <= m; row++) {
		for (col = 1; col <= n; col++){
			index = row * WIDTH + col;
			l2norm += E[index] * E[index];
			if (E[index] > mx) {
				mx = E[index];
			}
		}
	}
		
	*_mx = mx;
	l2norm /= (double)((m) * (n));
	l2norm = sqrt(l2norm);
	return l2norm;
}

__global__ void simulate_version1(const double alpha, const int n, const int m, const double kk, const double dt, const double a, const double epsilon, const double M1, const double M2, const double b, double* E_1D, double* E_prev_1D, double* R_1D, int WIDTH){

   
  int RADIUS = 1;
  int row = blockIdx.y * blockDim.y + threadIdx.y + RADIUS;
  int col = blockIdx.x * blockDim.x + threadIdx.x + RADIUS;
  
  if (row < (WIDTH -1) && col < (WIDTH -1)) {
      E_1D[row*WIDTH+col] = E_prev_1D[row*WIDTH+col] + alpha * (E_prev_1D[row*WIDTH+(col+1)] + E_prev_1D[row*WIDTH + (col-1)] - 4 * E_prev_1D[row*WIDTH+ col] + E_prev_1D[(row + 1)*WIDTH + col] + E_prev_1D[(row - 1) * WIDTH + col]);
  }   

}

void simulate(double **E, double **E_prev, double **R, const double alpha, const int n, const int m, const double kk, const double dt, const double a, const double epsilon, const double M1, const double M2, const double b, double* E_1D, double* E_prev_1D, double* R_1D, int WIDTH){
	
	int col, row;
	/* 
	* Copy data from boundary of the computational box 
	* to the padding region, set up for differencing
	* on the boundary of the computational box
	* Using mirror boundaries
	*/

	for (row = 1; row <= m; row++) {
		E_prev[row][0] = E_prev[row][2];
		E_prev_1D[row* WIDTH + 0] = E_prev_1D[row*WIDTH + 2];
	}
	
	for (row = 1; row <= m; row++) {
		E_prev[row][n + 1] = E_prev[row][n - 1];
		E_prev_1D[row*WIDTH+ (n+1)] = E_prev_1D[row*WIDTH + (n-1)];
	}

	for (col = 1; col <= n; col++) {
		E_prev[0][col] = E_prev[2][col];
		E_prev_1D[0*WIDTH+col] = E_prev_1D[2*WIDTH+ col];
	}

	for (col = 1; col <= n; col++) {
		E_prev[m + 1][col] = E_prev[m - 1][col];
		E_prev_1D[(m+1)*WIDTH+ col] = E_prev_1D[(m-1)*WIDTH+ col];
    }
    
    // ================ Kernel Setup ==============
    // double *d_E_1D, *d_E_prev_1D, *d_R_1D;
    // const dim3 block_size(TILE_DIM, TILE_DIM);
    // const dim3 num_blocks(WIDTH / block_size.x, WIDTH / block_size.y);
    // int Total_Bytes = WIDTH * WIDTH * sizeof(double);

    // // Allocate space on the device
    // cudaMalloc( &d_E_1D, Total_Bytes) ;
    // cudaMalloc( &d_E_prev_1D,Total_Bytes) ;
    // cudaMalloc(&d_R_1D, Total_Bytes);
    
    // // Copy input data to device
    // cudaMemcpy(d_E_1D, E_1D, Total_Bytes, cudaMemcpyHostToDevice) ;
    // cudaMemcpy(d_E_prev_1D, E_prev_1D, Total_Bytes, cudaMemcpyHostToDevice) ;
    // cudaMemcpy(d_R_1D, R_1D, Total_Bytes, cudaMemcpyHostToDevice) ;
    
    // // simulate_version1<<<num_blocks, block_size>>>(alpha, n, m, kk, dt, a, epsilon, M1, M2, b, d_E_1D, d_E_prev_1D, d_R_1D, WIDTH);

    // cudaMemcpy(E_1D ,d_E_1D, Total_Bytes, cudaMemcpyDeviceToHost);

    // cudaFree(d_E_1D);
    // cudaFree(d_E_prev_1D);
    // cudaFree(d_R_1D);
    // =================================================

	// Solve for the excitation, the PDE
	for (row = 1; row <= m; row++){
		for (col = 1; col <= n; col++){
            E[row][col] = E_prev[row][col] + alpha * (E_prev[row][col + 1] + E_prev[row][col - 1] - 4 * E_prev[row][col] + E_prev[row + 1][col] + E_prev[row - 1][col]);
            
			E_1D[row*WIDTH+col] = E_prev_1D[row*WIDTH+col] + alpha * (E_prev_1D[row*WIDTH+(col+1)] + E_prev_1D[row*WIDTH + (col-1)] - 4 * E_prev_1D[row*WIDTH+ col] + E_prev_1D[(row + 1)*WIDTH + col] + E_prev_1D[(row - 1) * WIDTH + col]);
		}
	}

	/* 
	* Solve the ODE, advancing excitation and recovery to the
	*     next timtestep
	*/
	int index = -1;
	for (row = 1; row <= m; row++){
		for (col = 1; col <= n; col++) {
			E[row][col] = E[row][col] - dt * (kk * E[row][col] * (E[row][col] - a) * (E[row][col] - 1) + E[row][col] * R[row][col]);
			index = row * WIDTH + col;	
			E_1D[index] = E_1D[index] - dt * (kk * E_1D[index] * (E_1D[index] - a) * (E_1D[index] - 1) + E_1D[index] * R_1D[index]);

		}
	}

	for (row = 1; row <= m; row++){
		for (col = 1; col <= n; col++) {
			
			R[row][col] = R[row][col] + dt * (epsilon + M1 * R[row][col] / (E[row][col] + M2)) * (-R[row][col] - kk * E[row][col] * (E[row][col] - b - 1));

			index =  row * WIDTH + col;
			R_1D[index] = R_1D[index] + dt * (epsilon + M1 * R_1D[index] / (E_1D[index] + M2)) * (-R_1D[index] - kk * E_1D[index] * (E_1D[index] - b - 1));
		}
	}

}

// Main program
int main(int argc, char **argv){
	/*
	*  Solution arrays
	*   E is the "Excitation" variable, a voltage
	*   R is the "Recovery" variable
	*   E_prev is the Excitation variable for the previous timestep,
	*      and is used in time integration
	*/
    cout<< "\n Hello I am waris"<<endl;
	double **E, **R, **E_prev;

	double * E_1D, *R_1D, *E_prev_1D; 

	// Various constants - these definitions shouldn't change
	const double a = 0.1, b = 0.1, kk = 8.0, M1 = 0.07, M2 = 0.3, epsilon = 0.01, d = 5e-5;

	double T = 1000.0;
	int m = 200, n = 200;
	int plot_freq = 0;
	int px = 1, py = 1;
	int no_comm = 0;
	int num_threads = 1;
	int WIDTH;

	//int size;

	cmdLine(argc, argv, T, n, px, py, plot_freq, no_comm, num_threads);
	m = n;
	// Allocate contiguous memory for solution arrays
	// The computational box is defined on [1:m+1,1:n+1]
	// We pad the arrays in order to facilitate differencing on the
	// boundaries of the computation box
	int Total_Bytes = (m+2) * (n+2) *sizeof(double);
	WIDTH =  m+2;


	E = alloc2D(m + 2, n + 2);
	E_prev = alloc2D(m + 2, n + 2);
	R = alloc2D(m + 2, n + 2);
	
	E_1D = (double *) malloc (Total_Bytes);
	E_prev_1D = (double *) malloc (Total_Bytes);
    R_1D = (double *) malloc (Total_Bytes);
    
    


	int col, row;
	// Initialization
	
	for (row = 1; row <= m; row++) {
		for (col = 1; col <= n; col++) {
			E_prev[row][col] = 0;
			R[row][col] = 0;
			E_prev_1D[row*WIDTH + col] = 0;
			R_1D[row*WIDTH + col] = 0;
		}
	}
		
	for (row = 1; row <= m; row++) {
		for (col = n / 2 + 1; col <= n; col++) {
			E_prev[row][col] = 1.0;
			E_prev_1D[row*WIDTH + col] = 1.0;
		}
	}
		
	for (row = m / 2 + 1; row <= m; row++) {
		for (col = 1; col <= n; col++) {
			R[row][col] = 1.0;
			R_1D[row*WIDTH + col] = 1.0;
		}
	}

    //cout<< "check 1"<<endl;
		

	double dx = 1.0 / n;

	// For time integration, these values shouldn't change
	double rp = kk * (b + 1) * (b + 1) / 4;
	double dte = (dx * dx) / (d * 4 + ((dx * dx)) * (rp + kk));
	double dtr = 1 / (epsilon + ((M1 / M2) * rp));
	double dt = (dte < dtr) ? 0.95 * dte : 0.95 * dtr;
	double alpha = d * dt / (dx * dx);

	cout << "Grid Size       : " << n << endl;
	cout << "Duration of Sim : " << T << endl;
	cout << "Time step dt    : " << dt << endl;
	cout << "Process geometry: " << px << " x " << py << endl;
	
	if (no_comm) {
		cout << "Communication   : DISABLED" << endl;
	}

	cout << endl;

	// Start the timer
	double t0 = getTime();

	// Simulated time is different from the integer timestep number
	// Simulated time
	double t = 0.0;
	// Integer timestep number
	int niter = 0;

	while (t < T){

		t += dt;
		niter++;

		simulate(E, E_prev, R, alpha, n, m, kk, dt, a, epsilon, M1, M2, b, E_1D, E_prev_1D, R_1D, WIDTH);

		//swap current E with previous E
		double **tmp = E;
		E = E_prev;
		E_prev = tmp;

		double *tmp2 = E_1D;
		E_1D = E_prev_1D;
		E_prev_1D = tmp2;


		// if (plot_freq){
		// 	int k = (int)(t / plot_freq);
		// 	if ((t - k * plot_freq) < dt){
		// 		splot(E, t, niter, m + 2, n + 2);
		// 	}
		// }

	} //end of while loop

	double time_elapsed = getTime() - t0;

	double Gflops = (double)(niter * (1E-9 * n * n) * 28.0) / time_elapsed;
	double BW = (double)(niter * 1E-9 * (n * n * sizeof(double) * 4.0)) / time_elapsed;

	cout << "Number of Iterations        : " << niter << endl;
	cout << "Elapsed Time (sec)          : " << time_elapsed << endl;
	cout << "Sustained Gflops Rate       : " << Gflops << endl;
	cout << "Sustained Bandwidth (GB/sec): " << BW << endl << endl;

	double mx;
	double l2norm = stats(E_prev, m, n, &mx);
	cout << "Max: " << mx << " L2norm: " << l2norm << endl;

	l2norm = stats1D(E_prev_1D, m, n, &mx, WIDTH);
	cout << "1D Max : " << mx << " 1D L2norm: " << l2norm << endl;

	// if (plot_freq){
	// 	cout << "\n\nEnter any input to close the program and the plot..." << endl;
	// 	getchar();
	// }

	free(E);
	free(E_prev);
	free(R);

	free(E_1D);
	free(E_prev_1D);
	free(R_1D);

	return 0;
}


/* **********************************************************
*  Author : Urvashi R.V. [04/06/2004]
*      Modified by Didem Unat [03/23/18]
*************************************************************/

//#include <stdio.h>

/* Function to plot the 2D array
* 'gnuplot' is instantiated via a pipe and 
* the values to be plotted are passed through, along 
* with gnuplot commands */

// FILE *gnu = NULL;

// void splot(double **U, double T, int niter, int m, int n){
// 	int col, row;
// 	if (gnu == NULL)
// 		gnu = popen("gnuplot", "w");

// 	double mx = -1, mn = 32768;
// 	for (row = 0; row < m; row++)
// 		for (col = 0; col < n; col++)
// 		{
// 			if (U[row][col] > mx)
// 				mx = U[row][col];
// 			if (U[row][col] < mn)
// 				mn = U[row][col];
// 		}

// 	fprintf(gnu, "set title \"T = %f [niter = %d]\"\n", T, niter);
// 	fprintf(gnu, "set size square\n");
// 	fprintf(gnu, "set key off\n");
// 	fprintf(gnu, "set pm3d map\n");
// 	// Various color schemes
// 	fprintf(gnu, "set palette defined (-3 \"blue\", 0 \"white\", 1 \"red\")\n");

// 	//    fprintf(gnu,"set palette rgbformulae 22, 13, 31\n");
// 	//    fprintf(gnu,"set palette rgbformulae 30, 31, 32\n");

// 	fprintf(gnu, "splot [0:%d] [0:%d][%f:%f] \"-\"\n", m - 1, n - 1, mn, mx);
// 	for (row = 0; row < m; row++)
// 	{
// 		for (col = 0; col < n; col++)
// 		{
// 			fprintf(gnu, "%d %d %f\n", col, row, U[col][row]);
// 		}
// 		fprintf(gnu, "\n");
// 	}
// 	fprintf(gnu, "e\n");
// 	fflush(gnu);
// 	return;
// }
	.file	"tmpxft_00261fd1_00000000-6_vecAdd_UVM.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3924:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3924:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Usage: a.out [-n <domain size>] [-t <final time >]\n\t [-p <plot frequency>]\n\t[-px <x processor geometry> [-py <y proc. geometry] [-k turn off communication] [-o <Number of OpenMP threads>]\n"
.LC1:
	.string	"n:x:y:t:kp:o:"
	.text
	.globl	_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_
	.type	_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_, @function
_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_:
.LFB3915:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%r9, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%r8, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rcx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	$1, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%edi, %ebx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movq	%rsi, 8(%rsp)
.L3:
	cmpl	%ebx, %ebp
	jge	.L20
.L15:
	movq	8(%rsp), %rsi
	xorl	%r8d, %r8d
	leaq	_ZZ7cmdLineiPPcRdRiS2_S2_S2_S2_S2_E12long_options(%rip), %rcx
	movl	%ebx, %edi
	leaq	.LC1(%rip), %rdx
	call	getopt_long@PLT
	cmpl	$-1, %eax
	je	.L21
	subl	$107, %eax
	cmpl	$14, %eax
	ja	.L4
	leaq	.L6(%rip), %rdx
	movslq	(%rdx,%rax,4), %rax
	addq	%rdx, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L6:
	.long	.L12-.L6
	.long	.L4-.L6
	.long	.L4-.L6
	.long	.L11-.L6
	.long	.L10-.L6
	.long	.L9-.L6
	.long	.L4-.L6
	.long	.L4-.L6
	.long	.L4-.L6
	.long	.L8-.L6
	.long	.L4-.L6
	.long	.L4-.L6
	.long	.L4-.L6
	.long	.L7-.L6
	.long	.L5-.L6
	.text
.L11:
	movq	optarg(%rip), %rdi
	call	atoi@PLT
	movl	%eax, 0(%r13)
	jmp	.L15
.L7:
	movq	optarg(%rip), %rdi
	call	atoi@PLT
	movl	%eax, (%r14)
.L5:
	movq	optarg(%rip), %rdi
	call	atoi@PLT
	movl	%eax, (%r15)
.L8:
	movq	optarg(%rip), %rdi
	call	atof@PLT
	movsd	%xmm0, (%r12)
	jmp	.L15
.L12:
	movq	88(%rsp), %rax
	movl	$1, (%rax)
	jmp	.L15
.L9:
	movq	optarg(%rip), %rdi
	call	atoi@PLT
	movq	80(%rsp), %rcx
	jmp	.L18
.L10:
	movq	optarg(%rip), %rdi
	call	atoi@PLT
	movq	96(%rsp), %rcx
.L18:
	movl	%eax, (%rcx)
	jmp	.L15
.L4:
	movl	$2, %edi
	leaq	.LC0(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	orl	$-1, %edi
	call	exit@PLT
.L21:
	incl	%ebp
	jmp	.L3
.L20:
	addq	$24, %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
.LFE3915:
	.size	_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_, .-_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_
	.section	.rodata.str1.1
.LC3:
	.string	"ERROR: Bad call to gettimeofday"
	.text
	.globl	_Z7getTimev
	.type	_Z7getTimev, @function
_Z7getTimev:
.LFB3916:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rsi
	leaq	8(%rsp), %rdi
	call	gettimeofday@PLT
	incl	%eax
	jne	.L23
	leaq	.LC3(%rip), %rsi
	leaq	_ZSt4cerr(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movsd	.LC2(%rip), %xmm0
	jmp	.L22
.L23:
	cvtsi2sdq	16(%rsp), %xmm0
	mulsd	.LC4(%rip), %xmm0
	cvtsi2sdq	8(%rsp), %xmm1
	addsd	%xmm1, %xmm0
.L22:
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3916:
	.size	_Z7getTimev, .-_Z7getTimev
	.globl	_Z7alloc2Dii
	.type	_Z7alloc2Dii, @function
_Z7alloc2Dii:
.LFB3917:
	.cfi_startproc
	endbr64
	movslq	%edi, %rdx
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movslq	%esi, %rbp
	movq	%rdx, %r12
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	leaq	1(%rbp), %rdi
	leaq	0(,%rdx,8), %rbx
	salq	$3, %rbp
	imulq	%rbx, %rdi
	call	malloc@PLT
	xorl	%ecx, %ecx
	leaq	(%rax,%rbx), %rdx
.L28:
	cmpl	%ecx, %r12d
	jle	.L31
	movq	%rdx, (%rax,%rcx,8)
	incq	%rcx
	addq	%rbp, %rdx
	jmp	.L28
.L31:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3917:
	.size	_Z7alloc2Dii, .-_Z7alloc2Dii
	.globl	_Z5statsPPdiiS_
	.type	_Z5statsPPdiiS_, @function
_Z5statsPPdiiS_:
.LFB3918:
	.cfi_startproc
	endbr64
	movsd	.LC2(%rip), %xmm1
	movl	$1, %eax
	xorps	%xmm0, %xmm0
.L33:
	cmpl	%eax, %esi
	jl	.L38
	movl	$8, %r9d
	movl	$1, %r8d
.L36:
	cmpl	%edx, %r8d
	jg	.L39
	movq	(%rdi,%rax,8), %r10
	incl	%r8d
	movsd	(%r10,%r9), %xmm2
	addq	$8, %r9
	movaps	%xmm2, %xmm3
	mulsd	%xmm2, %xmm3
	maxsd	%xmm1, %xmm2
	movaps	%xmm2, %xmm1
	addsd	%xmm3, %xmm0
	jmp	.L36
.L39:
	incq	%rax
	jmp	.L33
.L38:
	imull	%edx, %esi
	movsd	%xmm1, (%rcx)
	cvtsi2sdl	%esi, %xmm1
	divsd	%xmm1, %xmm0
	jmp	sqrt@PLT
	.cfi_endproc
.LFE3918:
	.size	_Z5statsPPdiiS_, .-_Z5statsPPdiiS_
	.globl	_Z7stats1DPdiiS_i
	.type	_Z7stats1DPdiiS_i, @function
_Z7stats1DPdiiS_i:
.LFB3919:
	.cfi_startproc
	endbr64
	movsd	.LC2(%rip), %xmm1
	movl	%esi, %eax
	movl	%r8d, %r9d
	xorps	%xmm0, %xmm0
	movl	$1, %r10d
.L41:
	cmpl	%eax, %r10d
	jg	.L44
	movslq	%r9d, %rsi
	leaq	(%rdi,%rsi,8), %r11
	movl	$1, %esi
.L45:
	cmpl	%esi, %edx
	jl	.L46
	movsd	(%r11,%rsi,8), %xmm2
	incq	%rsi
	movaps	%xmm2, %xmm3
	mulsd	%xmm2, %xmm3
	maxsd	%xmm1, %xmm2
	movaps	%xmm2, %xmm1
	addsd	%xmm3, %xmm0
	jmp	.L45
.L46:
	incl	%r10d
	addl	%r8d, %r9d
	jmp	.L41
.L44:
	imull	%edx, %eax
	movsd	%xmm1, (%rcx)
	cvtsi2sdl	%eax, %xmm1
	divsd	%xmm1, %xmm0
	jmp	sqrt@PLT
	.cfi_endproc
.LFE3919:
	.size	_Z7stats1DPdiiS_i, .-_Z7stats1DPdiiS_i
	.globl	_Z8simulatePPdS0_S0_diidddddddS_S_S_i
	.type	_Z8simulatePPdS0_S0_diidddddddS_S_S_i, @function
_Z8simulatePPdS0_S0_diidddddddS_S_S_i:
.LFB3920:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdi, %r11
	movl	%ecx, %edi
	movaps	%xmm0, %xmm9
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	xorl	%eax, %eax
	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
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movslq	72(%rsp), %r14
	movq	56(%rsp), %rcx
	movq	64(%rsp), %r13
	leaq	0(,%r14,8), %r10
	movq	%r14, %rdx
	leaq	(%rcx,%r10), %rbx
.L48:
	incq	%rax
	cmpl	%eax, %r8d
	jl	.L69
	movq	(%rsi,%rax,8), %r12
	movsd	16(%r12), %xmm0
	movsd	%xmm0, (%r12)
	movsd	16(%rbx), %xmm0
	movsd	%xmm0, (%rbx)
	addq	%r10, %rbx
	jmp	.L48
.L69:
	movslq	%edi, %rax
	leaq	-8(,%rax,8), %rbx
	leaq	8(,%rax,8), %r12
	addq	%r14, %rax
	leaq	(%rcx,%rax,8), %r14
	xorl	%eax, %eax
.L50:
	incq	%rax
	cmpl	%eax, %r8d
	jl	.L70
	movq	(%rsi,%rax,8), %r15
	movsd	(%r15,%rbx), %xmm0
	movsd	%xmm0, (%r15,%r12)
	movsd	-8(%r14), %xmm0
	movsd	%xmm0, 8(%r14)
	addq	%r10, %r14
	jmp	.L50
.L70:
	leal	(%rdx,%rdx), %r14d
	movslq	%r14d, %rax
	leaq	(%rcx,%rax,8), %r10
	xorl	%eax, %eax
.L52:
	incq	%rax
	cmpl	%eax, %edi
	jl	.L71
	movq	16(%rsi), %rbx
	movsd	(%rbx,%rax,8), %xmm0
	movq	(%rsi), %rbx
	movsd	%xmm0, (%rbx,%rax,8)
	movsd	(%r10,%rax,8), %xmm0
	movsd	%xmm0, (%rcx,%rax,8)
	jmp	.L52
.L71:
	movslq	%r8d, %rax
	leaq	-8(%rsi,%rax,8), %r10
	leaq	8(%rsi,%rax,8), %rbx
	leal	-1(%r8), %eax
	imull	%edx, %eax
	movslq	%eax, %r12
	addl	%r14d, %eax
	cltq
	leaq	(%rcx,%r12,8), %r12
	leaq	(%rcx,%rax,8), %r14
	xorl	%eax, %eax
.L54:
	incq	%rax
	cmpl	%eax, %edi
	jl	.L72
	movq	(%r10), %r15
	movsd	(%r15,%rax,8), %xmm0
	movq	(%rbx), %r15
	movsd	%xmm0, (%r15,%rax,8)
	movsd	(%r12,%rax,8), %xmm0
	movsd	%xmm0, (%r14,%rax,8)
	jmp	.L54
.L72:
	movsd	.LC6(%rip), %xmm10
	movl	%edx, %r12d
	xorl	%r14d, %r14d
	movl	$1, %r10d
.L56:
	cmpl	%r10d, %r8d
	jl	.L66
	movslq	%r12d, %rax
	addl	%edx, %r12d
	movl	$1, -36(%rsp)
	movslq	%r12d, %r15
	salq	$3, %rax
	leaq	(%rcx,%r15,8), %r15
	leaq	(%rcx,%rax), %rbx
	addq	%r9, %rax
	movq	%r15, -24(%rsp)
	movslq	%r14d, %r15
	leaq	(%rcx,%r15,8), %r15
	movq	%rax, -8(%rsp)
	xorl	%eax, %eax
	movq	%r15, -16(%rsp)
.L59:
	leaq	8(%rax), %r15
	addq	$8, %rbx
	movq	%r15, -32(%rsp)
	cmpl	%edi, -36(%rsp)
	jg	.L73
	movq	(%rsi,%r10,8), %r15
	incl	-36(%rsp)
	movsd	8(%r15,%rax), %xmm11
	movsd	16(%r15,%rax), %xmm8
	movsd	(%r15,%rax), %xmm0
	movq	8(%rsi,%r10,8), %r15
	addsd	%xmm8, %xmm0
	movaps	%xmm11, %xmm8
	mulsd	%xmm10, %xmm8
	subsd	%xmm8, %xmm0
	addsd	8(%r15,%rax), %xmm0
	movq	-8(%rsi,%r10,8), %r15
	addsd	8(%r15,%rax), %xmm0
	movq	(%r11,%r10,8), %r15
	mulsd	%xmm9, %xmm0
	addsd	%xmm11, %xmm0
	movsd	%xmm0, 8(%r15,%rax)
	movsd	(%rbx), %xmm11
	movsd	8(%rbx), %xmm8
	movsd	-8(%rbx), %xmm0
	movq	-24(%rsp), %r15
	addsd	%xmm8, %xmm0
	movaps	%xmm11, %xmm8
	mulsd	%xmm10, %xmm8
	subsd	%xmm8, %xmm0
	addsd	8(%r15,%rax), %xmm0
	movq	-16(%rsp), %r15
	addsd	8(%r15,%rax), %xmm0
	movq	-8(%rsp), %r15
	mulsd	%xmm9, %xmm0
	addsd	%xmm11, %xmm0
	movsd	%xmm0, 8(%r15,%rax)
	movq	-32(%rsp), %rax
	jmp	.L59
.L73:
	incq	%r10
	addl	%edx, %r14d
	jmp	.L56
.L66:
	movsd	.LC7(%rip), %xmm8
	movl	%edx, %r10d
	movl	$1, %esi
.L58:
	cmpl	%esi, %r8d
	jl	.L67
	movslq	%r10d, %rax
	xorl	%ecx, %ecx
	salq	$3, %rax
.L62:
	incq	%rcx
	addq	$8, %rax
	cmpl	%ecx, %edi
	jl	.L74
	movq	(%r11,%rsi,8), %rbx
	movq	0(%rbp,%rsi,8), %r12
	leaq	(%rbx,%rcx,8), %rbx
	movsd	(%rbx), %xmm9
	movaps	%xmm9, %xmm0
	movaps	%xmm9, %xmm10
	mulsd	%xmm1, %xmm0
	subsd	%xmm3, %xmm10
	mulsd	%xmm10, %xmm0
	movaps	%xmm9, %xmm10
	subsd	%xmm8, %xmm10
	mulsd	%xmm10, %xmm0
	movsd	(%r12,%rcx,8), %xmm10
	mulsd	%xmm9, %xmm10
	addsd	%xmm10, %xmm0
	mulsd	%xmm2, %xmm0
	subsd	%xmm0, %xmm9
	movsd	%xmm9, (%rbx)
	movsd	(%r9,%rax), %xmm9
	movaps	%xmm9, %xmm0
	movaps	%xmm9, %xmm10
	mulsd	%xmm1, %xmm0
	subsd	%xmm3, %xmm10
	mulsd	%xmm10, %xmm0
	movaps	%xmm9, %xmm10
	subsd	%xmm8, %xmm10
	mulsd	%xmm10, %xmm0
	movsd	0(%r13,%rax), %xmm10
	mulsd	%xmm9, %xmm10
	addsd	%xmm10, %xmm0
	mulsd	%xmm2, %xmm0
	subsd	%xmm0, %xmm9
	movsd	%xmm9, (%r9,%rax)
	jmp	.L62
.L74:
	incq	%rsi
	addl	%edx, %r10d
	jmp	.L58
.L67:
	movq	.LC8(%rip), %xmm3
	movl	%edx, %r10d
	movl	$1, %esi
	movsd	.LC7(%rip), %xmm8
.L61:
	cmpl	%esi, %r8d
	jl	.L47
	movslq	%r10d, %rax
	xorl	%ecx, %ecx
	salq	$3, %rax
.L65:
	incq	%rcx
	addq	$8, %rax
	cmpl	%ecx, %edi
	jl	.L75
	movq	(%r11,%rsi,8), %r12
	movq	0(%rbp,%rsi,8), %rbx
	movsd	(%r12,%rcx,8), %xmm10
	leaq	(%rbx,%rcx,8), %rbx
	movsd	(%rbx), %xmm11
	movaps	%xmm10, %xmm12
	movaps	%xmm10, %xmm9
	addsd	%xmm6, %xmm10
	subsd	%xmm7, %xmm12
	mulsd	%xmm1, %xmm9
	movaps	%xmm11, %xmm0
	xorps	%xmm3, %xmm0
	subsd	%xmm8, %xmm12
	mulsd	%xmm9, %xmm12
	subsd	%xmm12, %xmm0
	movaps	%xmm0, %xmm9
	movaps	%xmm11, %xmm0
	mulsd	%xmm5, %xmm0
	divsd	%xmm10, %xmm0
	addsd	%xmm4, %xmm0
	mulsd	%xmm2, %xmm0
	mulsd	%xmm9, %xmm0
	addsd	%xmm11, %xmm0
	movsd	%xmm0, (%rbx)
	movsd	(%r9,%rax), %xmm10
	movsd	0(%r13,%rax), %xmm11
	movaps	%xmm10, %xmm12
	movaps	%xmm10, %xmm9
	movaps	%xmm11, %xmm0
	subsd	%xmm7, %xmm12
	mulsd	%xmm1, %xmm9
	xorps	%xmm3, %xmm0
	addsd	%xmm6, %xmm10
	subsd	%xmm8, %xmm12
	mulsd	%xmm9, %xmm12
	subsd	%xmm12, %xmm0
	movaps	%xmm0, %xmm9
	movaps	%xmm11, %xmm0
	mulsd	%xmm5, %xmm0
	divsd	%xmm10, %xmm0
	addsd	%xmm4, %xmm0
	mulsd	%xmm2, %xmm0
	mulsd	%xmm9, %xmm0
	addsd	%xmm11, %xmm0
	movsd	%xmm0, 0(%r13,%rax)
	jmp	.L65
.L75:
	incq	%rsi
	addl	%edx, %r10d
	jmp	.L61
.L47:
	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
.LFE3920:
	.size	_Z8simulatePPdS0_S0_diidddddddS_S_S_i, .-_Z8simulatePPdS0_S0_diidddddddS_S_S_i
	.section	.rodata.str1.1
.LC10:
	.string	"\n Hello I am waris"
.LC17:
	.string	"Grid Size       : "
.LC18:
	.string	"Duration of Sim : "
.LC19:
	.string	"Time step dt    : "
.LC20:
	.string	"Process geometry: "
.LC21:
	.string	" x "
.LC22:
	.string	"Communication   : DISABLED"
.LC30:
	.string	"Number of Iterations        : "
.LC31:
	.string	"Elapsed Time (sec)          : "
.LC32:
	.string	"Sustained Gflops Rate       : "
.LC33:
	.string	"Sustained Bandwidth (GB/sec): "
.LC34:
	.string	"Max: "
.LC35:
	.string	" L2norm: "
.LC36:
	.string	"1D Max : "
.LC37:
	.string	" 1D L2norm: "
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3921:
	.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
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	leaq	.LC10(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%edi, %ebx
	leaq	_ZSt4cout(%rip), %rdi
	subq	$120, %rsp
	.cfi_def_cfa_offset 176
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	.LC11(%rip), %rax
	xorl	%edx, %edx
	xorl	%ecx, %ecx
	movl	%edx, 68(%rsp)
	leaq	88(%rsp), %rdx
	movl	%ebx, %edi
	movq	%rax, 88(%rsp)
	movl	%ecx, 80(%rsp)
	leaq	64(%rsp), %rcx
	movl	$200, 64(%rsp)
	movl	$1, 72(%rsp)
	movl	$1, 76(%rsp)
	movl	$1, 84(%rsp)
	pushq	%rsi
	.cfi_def_cfa_offset 184
	movq	%rbp, %rsi
	leaq	92(%rsp), %rax
	pushq	%rax
	.cfi_def_cfa_offset 192
	leaq	96(%rsp), %rax
	pushq	%rax
	.cfi_def_cfa_offset 200
	leaq	92(%rsp), %rax
	pushq	%rax
	.cfi_def_cfa_offset 208
	leaq	108(%rsp), %r9
	leaq	104(%rsp), %r8
	call	_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_
	movl	96(%rsp), %ebx
	addq	$32, %rsp
	.cfi_def_cfa_offset 176
	leal	2(%rbx), %ebp
	movl	%ebp, %r13d
	movl	%ebp, %esi
	movl	%ebp, %edi
	call	_Z7alloc2Dii
	imull	%ebp, %r13d
	movl	%ebp, %esi
	movl	%ebp, %edi
	movq	%rax, 8(%rsp)
	call	_Z7alloc2Dii
	movl	%ebp, %esi
	movl	%ebp, %edi
	sall	$3, %r13d
	movq	%rax, %r12
	call	_Z7alloc2Dii
	movslq	%r13d, %r13
	movq	%r13, %rdi
	movq	%rax, 24(%rsp)
	call	malloc@PLT
	movq	%r13, %rdi
	movq	%rax, 16(%rsp)
	call	malloc@PLT
	movq	%r13, %rdi
	movq	%rax, %r15
	call	malloc@PLT
	movslq	%ebp, %rsi
	movl	$1, %r8d
	salq	$3, %rsi
	movq	%rax, %r13
	leaq	(%rax,%rsi), %r10
	leal	-1(%rbx), %eax
	leaq	16(,%rax,8), %rdx
	leaq	8(,%rax,8), %rax
	movq	%rax, (%rsp)
	leaq	(%r15,%rsi), %r9
	xorl	%eax, %eax
.L77:
	cmpl	%r8d, %ebx
	jl	.L103
	movq	24(%rsp), %rcx
	movq	(%r12,%r8,8), %rdi
	movq	(%rcx,%r8,8), %rcx
	leaq	8(%rdi), %r14
	movq	%r14, 32(%rsp)
	leaq	(%rdi,%rdx), %r14
	leaq	8(%rcx), %r11
	cmpq	%r14, %r11
	jnb	.L98
	leaq	(%rcx,%rdx), %r14
	cmpq	%r14, 32(%rsp)
	jb	.L96
.L98:
	movq	32(%rsp), %rdi
	movq	(%rsp), %rcx
	rep stosb
	movq	(%rsp), %rcx
	movq	%r11, %rdi
	leaq	8(%r9), %r11
	rep stosb
	movq	(%rsp), %rcx
	movq	%r11, %rdi
	leaq	8(%r10), %r11
	rep stosb
	movq	(%rsp), %rcx
	movq	%r11, %rdi
	rep stosb
	jmp	.L80
.L96:
	movl	$1, %r11d
.L78:
	movq	$0x000000000, (%rdi,%r11,8)
	movq	$0x000000000, (%r9,%r11,8)
	movq	$0x000000000, (%rcx,%r11,8)
	movq	$0x000000000, (%r10,%r11,8)
	incq	%r11
	cmpl	%r11d, %ebx
	jge	.L78
.L80:
	incq	%r8
	addq	%rsi, %r9
	addq	%rsi, %r10
	jmp	.L77
.L103:
	movl	%ebx, %ecx
	movsd	.LC7(%rip), %xmm0
	movl	%ebp, %esi
	movl	$1, %edx
	sarl	%ecx
	incl	%ecx
.L82:
	cmpl	%edx, %ebx
	jl	.L104
	movslq	%esi, %rdi
	movl	%ecx, %eax
	leaq	(%r15,%rdi,8), %rdi
.L83:
	cmpl	%eax, %ebx
	jl	.L105
	movq	(%r12,%rdx,8), %r8
	movsd	%xmm0, (%rdi,%rax,8)
	movsd	%xmm0, (%r8,%rax,8)
	incq	%rax
	jmp	.L83
.L105:
	incq	%rdx
	addl	%ebp, %esi
	jmp	.L82
.L104:
	movl	%ebx, %eax
	movl	$2, %ecx
	movsd	.LC7(%rip), %xmm0
	cltd
	idivl	%ecx
	incl	%eax
	movslq	%eax, %rcx
	imull	%ebp, %eax
.L86:
	cmpl	%ecx, %ebx
	jl	.L88
	movslq	%eax, %rdx
	leaq	0(%r13,%rdx,8), %rsi
	xorl	%edx, %edx
.L89:
	incq	%rdx
	cmpl	%edx, %ebx
	jl	.L106
	movq	24(%rsp), %rdi
	movsd	%xmm0, (%rsi,%rdx,8)
	movq	(%rdi,%rcx,8), %rdi
	movsd	%xmm0, (%rdi,%rdx,8)
	jmp	.L89
.L106:
	incq	%rcx
	addl	%ebp, %eax
	jmp	.L86
.L88:
	cvtsi2sdl	%ebx, %xmm4
	movsd	.LC7(%rip), %xmm0
	movsd	.LC12(%rip), %xmm1
	divsd	%xmm4, %xmm0
	movsd	%xmm4, 32(%rsp)
	mulsd	%xmm0, %xmm0
	mulsd	%xmm0, %xmm1
	movaps	%xmm0, %xmm2
	addsd	.LC13(%rip), %xmm1
	divsd	%xmm1, %xmm2
	movsd	.LC14(%rip), %xmm1
	comisd	%xmm2, %xmm1
	jbe	.L101
	mulsd	.LC15(%rip), %xmm2
	movsd	%xmm2, (%rsp)
	jmp	.L90
.L101:
	movsd	.LC9(%rip), %xmm6
	movsd	%xmm6, (%rsp)
.L90:
	movsd	(%rsp), %xmm1
	mulsd	.LC16(%rip), %xmm1
	leaq	.LC17(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	divsd	%xmm0, %xmm1
	movsd	%xmm1, 48(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC18(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	88(%rsp), %xmm7
	movq	%rax, %rdi
	movaps	%xmm7, %xmm0
	movsd	%xmm7, 40(%rsp)
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC19(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC20(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	72(%rsp), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC21(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	76(%rsp), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	cmpl	$0, 80(%rsp)
	je	.L92
	leaq	.LC22(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
.L92:
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	xorl	%r14d, %r14d
	call	_ZNSolsEPFRSoS_E@PLT
	call	_Z7getTimev
	xorps	%xmm13, %xmm13
	movsd	%xmm0, 56(%rsp)
.L93:
	movsd	40(%rsp), %xmm5
	comisd	%xmm13, %xmm5
	jbe	.L107
	addsd	(%rsp), %xmm13
	movsd	.LC23(%rip), %xmm7
	pushq	%rax
	.cfi_def_cfa_offset 184
	movq	%r12, %rsi
	pushq	%rbp
	.cfi_def_cfa_offset 192
	movsd	.LC24(%rip), %xmm6
	movl	%ebx, %r8d
	movl	%ebx, %ecx
	pushq	%r13
	.cfi_def_cfa_offset 200
	movsd	.LC25(%rip), %xmm5
	movaps	%xmm7, %xmm3
	incl	%r14d
	pushq	%r15
	.cfi_def_cfa_offset 208
	movsd	.LC26(%rip), %xmm4
	movq	48(%rsp), %r9
	movsd	32(%rsp), %xmm2
	movsd	80(%rsp), %xmm0
	movq	56(%rsp), %rdx
	movq	40(%rsp), %rdi
	movsd	.LC27(%rip), %xmm1
	call	_Z8simulatePPdS0_S0_diidddddddS_S_S_i
	movq	48(%rsp), %rax
	addq	$32, %rsp
	.cfi_def_cfa_offset 176
	movq	%r15, 16(%rsp)
	movq	%rax, %r15
	movq	8(%rsp), %rax
	movq	%r12, 8(%rsp)
	movq	%rax, %r12
	jmp	.L93
.L107:
	call	_Z7getTimev
	movsd	.LC28(%rip), %xmm3
	cvtsi2sdl	%r14d, %xmm2
	movl	%ebx, %eax
	movsd	32(%rsp), %xmm1
	imull	%ebx, %eax
	subsd	56(%rsp), %xmm0
	leaq	.LC30(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	mulsd	%xmm3, %xmm1
	mulsd	32(%rsp), %xmm1
	cltq
	movsd	%xmm0, 40(%rsp)
	salq	$3, %rax
	mulsd	%xmm2, %xmm1
	mulsd	.LC29(%rip), %xmm1
	mulsd	%xmm3, %xmm2
	divsd	%xmm0, %xmm1
	movsd	%xmm1, (%rsp)
	cvtsi2sdq	%rax, %xmm1
	mulsd	.LC6(%rip), %xmm1
	mulsd	%xmm2, %xmm1
	divsd	%xmm0, %xmm1
	movsd	%xmm1, 32(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r14d, %esi
	leaq	96(%rsp), %r14
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC31(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	40(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC32(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	leaq	.LC33(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	32(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movl	%ebx, %edx
	movq	%r14, %rcx
	movl	%ebx, %esi
	movq	%r12, %rdi
	call	_Z5statsPPdiiS_
	leaq	.LC34(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movsd	%xmm0, (%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	96(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC35(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movl	%ebx, %edx
	movl	%ebp, %r8d
	movq	%r14, %rcx
	movl	%ebx, %esi
	movq	%r15, %rdi
	call	_Z7stats1DPdiiS_i
	leaq	.LC36(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movq	%xmm0, %rbx
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movsd	96(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC37(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rbx, %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	8(%rsp), %rdi
	call	free@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	24(%rsp), %rdi
	call	free@PLT
	movq	16(%rsp), %rdi
	call	free@PLT
	movq	%r15, %rdi
	call	free@PLT
	movq	%r13, %rdi
	call	free@PLT
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L95
	call	__stack_chk_fail@PLT
.L95:
	addq	$120, %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
.LFE3921:
	.size	main, .-main
	.text
	.globl	_Z52__device_stub__Z17simulate_version1diidddddddPdS_S_idiidddddddPdS_S_i
	.type	_Z52__device_stub__Z17simulate_version1diidddddddPdS_S_idiidddddddPdS_S_i, @function
_Z52__device_stub__Z17simulate_version1diidddddddPdS_S_idiidddddddPdS_S_i:
.LFB3946:
	.cfi_startproc
	endbr64
	subq	$312, %rsp
	.cfi_def_cfa_offset 320
	movl	%edi, 100(%rsp)
	leaq	136(%rsp), %rdi
	movl	%esi, 96(%rsp)
	leaq	148(%rsp), %rsi
	movq	%rdx, 32(%rsp)
	leaq	120(%rsp), %rdx
	movq	%rcx, 24(%rsp)
	leaq	128(%rsp), %rcx
	movq	%r8, 16(%rsp)
	movl	%r9d, 12(%rsp)
	movsd	%xmm0, 104(%rsp)
	movsd	%xmm1, 88(%rsp)
	movsd	%xmm2, 80(%rsp)
	movsd	%xmm3, 72(%rsp)
	movsd	%xmm4, 64(%rsp)
	movsd	%xmm5, 56(%rsp)
	movsd	%xmm6, 48(%rsp)
	movsd	%xmm7, 40(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 296(%rsp)
	xorl	%eax, %eax
	leaq	104(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	100(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 288(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 136(%rsp)
	movl	$1, 144(%rsp)
	movq	%rax, 148(%rsp)
	movl	$1, 156(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L108
	pushq	128(%rsp)
	.cfi_def_cfa_offset 328
	leaq	_Z17simulate_version1diidddddddPdS_S_i(%rip), %rdi
	pushq	128(%rsp)
	.cfi_def_cfa_offset 336
	movq	164(%rsp), %rcx
	movl	172(%rsp), %r8d
	movq	152(%rsp), %rsi
	movl	160(%rsp), %edx
	leaq	200(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 328
	popq	%rdx
	.cfi_def_cfa_offset 320
.L108:
	movq	296(%rsp), %rax
	subq	%fs:40, %rax
	je	.L110
	call	__stack_chk_fail@PLT
.L110:
	addq	$312, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3946:
	.size	_Z52__device_stub__Z17simulate_version1diidddddddPdS_S_idiidddddddPdS_S_i, .-_Z52__device_stub__Z17simulate_version1diidddddddPdS_S_idiidddddddPdS_S_i
	.globl	_Z17simulate_version1diidddddddPdS_S_i
	.type	_Z17simulate_version1diidddddddPdS_S_i, @function
_Z17simulate_version1diidddddddPdS_S_i:
.LFB3947:
	.cfi_startproc
	endbr64
	jmp	_Z52__device_stub__Z17simulate_version1diidddddddPdS_S_idiidddddddPdS_S_i
	.cfi_endproc
.LFE3947:
	.size	_Z17simulate_version1diidddddddPdS_S_i, .-_Z17simulate_version1diidddddddPdS_S_i
	.section	.rodata.str1.1
.LC38:
	.string	"_Z17simulate_version1diidddddddPdS_S_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3949:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 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	.LC38(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z17simulate_version1diidddddddPdS_S_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3949:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.str1.1
.LC39:
	.string	"n"
.LC40:
	.string	"px"
.LC41:
	.string	"py"
.LC42:
	.string	"tfinal"
.LC43:
	.string	"plot"
.LC44:
	.string	"nocomm"
.LC45:
	.string	"numthreads"
	.section	.data.rel.local,"aw"
	.align 32
	.type	_ZZ7cmdLineiPPcRdRiS2_S2_S2_S2_S2_E12long_options, @object
	.size	_ZZ7cmdLineiPPcRdRiS2_S2_S2_S2_S2_E12long_options, 224
_ZZ7cmdLineiPPcRdRiS2_S2_S2_S2_S2_E12long_options:
	.quad	.LC39
	.long	1
	.zero	4
	.quad	0
	.long	110
	.zero	4
	.quad	.LC40
	.long	1
	.zero	4
	.quad	0
	.long	120
	.zero	4
	.quad	.LC41
	.long	1
	.zero	4
	.quad	0
	.long	121
	.zero	4
	.quad	.LC42
	.long	1
	.zero	4
	.quad	0
	.long	116
	.zero	4
	.quad	.LC43
	.long	1
	.zero	4
	.quad	0
	.long	112
	.zero	4
	.quad	.LC44
	.long	0
	.zero	4
	.quad	0
	.long	107
	.zero	4
	.quad	.LC45
	.long	1
	.zero	4
	.quad	0
	.long	111
	.zero	4
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC2:
	.long	0
	.long	-1074790400
	.align 8
.LC4:
	.long	-1598689907
	.long	1051772663
	.align 8
.LC6:
	.long	0
	.long	1074790400
	.align 8
.LC7:
	.long	0
	.long	1072693248
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC8:
	.long	0
	.long	-2147483648
	.long	0
	.long	0
	.section	.rodata.cst8
	.align 8
.LC9:
	.long	-847035317
	.long	1073378106
	.align 8
.LC11:
	.long	0
	.long	1083129856
	.align 8
.LC12:
	.long	1030792151
	.long	1076156170
	.align 8
.LC13:
	.long	-350469331
	.long	1059731170
	.align 8
.LC14:
	.long	916791160
	.long	1073469340
	.align 8
.LC15:
	.long	1717986918
	.long	1072588390
	.align 8
.LC16:
	.long	-350469331
	.long	1057634018
	.align 8
.LC23:
	.long	-1717986918
	.long	1069128089
	.align 8
.LC24:
	.long	858993459
	.long	1070805811
	.align 8
.LC25:
	.long	515396076
	.long	1068624773
	.align 8
.LC26:
	.long	1202590843
	.long	1065646817
	.align 8
.LC27:
	.long	0
	.long	1075838976
	.align 8
.LC28:
	.long	-400107883
	.long	1041313291
	.align 8
.LC29:
	.long	0
	.long	1077673984
	.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:

/*

	Implement your CUDA kernel in this file

*/



/* 
 * Solves the Panfilov model using an explicit numerical scheme.
 * Based on code orginally provided by Xing Cai, Simula Research Laboratory 
 * and reimplementation by Scott B. Baden, UCSD
 * 
 * Modified and  restructured by Didem Unat, Koc University
 *
 * Refer to "Detailed Numerical Analyses of the Aliev-Panfilov Model on GPGPU"
 * https://www.simula.no/publications/detailed-numerical-analyses-aliev-panfilov-model-gpgpu
 * by Xing Cai, Didem Unat and Scott Baden
 *
*/


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <assert.h>
#include <stdlib.h>
#include <iostream>
#include <iomanip>
#include <string.h>
#include <math.h>
#include <sys/time.h>
#include <getopt.h>

using namespace std;

#define TILE_DIM 16


// External functions
//extern "C" void splot(double **E, double T, int niter, int m, int n);

void cmdLine(int argc, char *argv[], double &T, int &n, int &px, int &py, int &plot_freq, int &no_comm, int &num_threads){
	/// Command line arguments
	// Default value of the domain sizes
	static struct option long_options[] = {
		{"n", required_argument, 0, 'n'},
		{"px", required_argument, 0, 'x'},
		{"py", required_argument, 0, 'y'},
		{"tfinal", required_argument, 0, 't'},
		{"plot", required_argument, 0, 'p'},
		{"nocomm", no_argument, 0, 'k'},
		{"numthreads", required_argument, 0, 'o'},
	};
	// Process command line arguments
	int ac;
	for (ac = 1; ac < argc; ac++)
	{
		int c;
		while ((c = getopt_long(argc, argv, "n:x:y:t:kp:o:", long_options, NULL)) != -1)
		{
			switch (c)
			{

				// Size of the computational box
			case 'n':
				n = atoi(optarg);
				break;

				// X processor geometry
			case 'x':
				px = atoi(optarg);

				// Y processor geometry
			case 'y':
				py = atoi(optarg);

				// Length of simulation, in simulated time units
			case 't':
				T = atof(optarg);
				break;
				// Turn off communication
			case 'k':
				no_comm = 1;
				break;

				// Plot the excitation variable
			case 'p':
				plot_freq = atoi(optarg);
				break;

				// Plot the excitation variable
			case 'o':
				num_threads = atoi(optarg);
				break;

				// Error
			default:
				printf("Usage: a.out [-n <domain size>] [-t <final time >]\n\t [-p <plot frequency>]\n\t[-px <x processor geometry> [-py <y proc. geometry] [-k turn off communication] [-o <Number of OpenMP threads>]\n");
				exit(-1);
			}
		}
	}
}

// Utilities
//

// Timer
// Make successive calls and take a difference to get the elapsed time.
static const double kMicro = 1.0e-6;

double getTime(){
	struct timeval TV;
	struct timezone TZ;

	const int RC = gettimeofday(&TV, &TZ);
	if (RC == -1)
	{
		cerr << "ERROR: Bad call to gettimeofday" << endl;
		return (-1);
	}

	return (((double)TV.tv_sec) + kMicro * ((double)TV.tv_usec));

} 

// end getTime()

// Allocate a 2D array
double **alloc2D(int m, int n){
	double **E;
	int nx = n, ny = m;
	E = (double **)malloc(sizeof(double *) * ny + sizeof(double) * nx * ny);
	assert(E);
	int row;
	for (row = 0; row < ny; row++)
		E[row] = (double *)(E + ny) + row * nx;
	return (E);
}

// Reports statistics about the computation
// These values should not vary (except to within roundoff)
// when we use different numbers of  processes to solve the problem
double stats(double **E, int m, int n, double *_mx){
	double mx = -1;
	double l2norm = 0;
	int col, row;
	for (row = 1; row <= m; row++)
		for (col = 1; col <= n; col++)
		{
			l2norm += E[row][col] * E[row][col];
			if (E[row][col] > mx)
				mx = E[row][col];
		}
	*_mx = mx;
	l2norm /= (double)((m) * (n));
	l2norm = sqrt(l2norm);
	return l2norm;
}

double stats1D(double *E, int m, int n, double *_mx, int WIDTH){
	double mx = -1;
	double l2norm = 0;
	int col, row;
	int index = -1;
	for (row = 1; row <= m; row++) {
		for (col = 1; col <= n; col++){
			index = row * WIDTH + col;
			l2norm += E[index] * E[index];
			if (E[index] > mx) {
				mx = E[index];
			}
		}
	}
		
	*_mx = mx;
	l2norm /= (double)((m) * (n));
	l2norm = sqrt(l2norm);
	return l2norm;
}

__global__ void simulate_version1(const double alpha, const int n, const int m, const double kk, const double dt, const double a, const double epsilon, const double M1, const double M2, const double b, double* E_1D, double* E_prev_1D, double* R_1D, int WIDTH){

   
  int RADIUS = 1;
  int row = blockIdx.y * blockDim.y + threadIdx.y + RADIUS;
  int col = blockIdx.x * blockDim.x + threadIdx.x + RADIUS;
  
  if (row < (WIDTH -1) && col < (WIDTH -1)) {
      E_1D[row*WIDTH+col] = E_prev_1D[row*WIDTH+col] + alpha * (E_prev_1D[row*WIDTH+(col+1)] + E_prev_1D[row*WIDTH + (col-1)] - 4 * E_prev_1D[row*WIDTH+ col] + E_prev_1D[(row + 1)*WIDTH + col] + E_prev_1D[(row - 1) * WIDTH + col]);
  }   

}

void simulate(double **E, double **E_prev, double **R, const double alpha, const int n, const int m, const double kk, const double dt, const double a, const double epsilon, const double M1, const double M2, const double b, double* E_1D, double* E_prev_1D, double* R_1D, int WIDTH){
	
	int col, row;
	/* 
	* Copy data from boundary of the computational box 
	* to the padding region, set up for differencing
	* on the boundary of the computational box
	* Using mirror boundaries
	*/

	for (row = 1; row <= m; row++) {
		E_prev[row][0] = E_prev[row][2];
		E_prev_1D[row* WIDTH + 0] = E_prev_1D[row*WIDTH + 2];
	}
	
	for (row = 1; row <= m; row++) {
		E_prev[row][n + 1] = E_prev[row][n - 1];
		E_prev_1D[row*WIDTH+ (n+1)] = E_prev_1D[row*WIDTH + (n-1)];
	}

	for (col = 1; col <= n; col++) {
		E_prev[0][col] = E_prev[2][col];
		E_prev_1D[0*WIDTH+col] = E_prev_1D[2*WIDTH+ col];
	}

	for (col = 1; col <= n; col++) {
		E_prev[m + 1][col] = E_prev[m - 1][col];
		E_prev_1D[(m+1)*WIDTH+ col] = E_prev_1D[(m-1)*WIDTH+ col];
    }
    
    // ================ Kernel Setup ==============
    // double *d_E_1D, *d_E_prev_1D, *d_R_1D;
    // const dim3 block_size(TILE_DIM, TILE_DIM);
    // const dim3 num_blocks(WIDTH / block_size.x, WIDTH / block_size.y);
    // int Total_Bytes = WIDTH * WIDTH * sizeof(double);

    // // Allocate space on the device
    // cudaMalloc( &d_E_1D, Total_Bytes) ;
    // cudaMalloc( &d_E_prev_1D,Total_Bytes) ;
    // cudaMalloc(&d_R_1D, Total_Bytes);
    
    // // Copy input data to device
    // cudaMemcpy(d_E_1D, E_1D, Total_Bytes, cudaMemcpyHostToDevice) ;
    // cudaMemcpy(d_E_prev_1D, E_prev_1D, Total_Bytes, cudaMemcpyHostToDevice) ;
    // cudaMemcpy(d_R_1D, R_1D, Total_Bytes, cudaMemcpyHostToDevice) ;
    
    // // simulate_version1<<<num_blocks, block_size>>>(alpha, n, m, kk, dt, a, epsilon, M1, M2, b, d_E_1D, d_E_prev_1D, d_R_1D, WIDTH);

    // cudaMemcpy(E_1D ,d_E_1D, Total_Bytes, cudaMemcpyDeviceToHost);

    // cudaFree(d_E_1D);
    // cudaFree(d_E_prev_1D);
    // cudaFree(d_R_1D);
    // =================================================

	// Solve for the excitation, the PDE
	for (row = 1; row <= m; row++){
		for (col = 1; col <= n; col++){
            E[row][col] = E_prev[row][col] + alpha * (E_prev[row][col + 1] + E_prev[row][col - 1] - 4 * E_prev[row][col] + E_prev[row + 1][col] + E_prev[row - 1][col]);
            
			E_1D[row*WIDTH+col] = E_prev_1D[row*WIDTH+col] + alpha * (E_prev_1D[row*WIDTH+(col+1)] + E_prev_1D[row*WIDTH + (col-1)] - 4 * E_prev_1D[row*WIDTH+ col] + E_prev_1D[(row + 1)*WIDTH + col] + E_prev_1D[(row - 1) * WIDTH + col]);
		}
	}

	/* 
	* Solve the ODE, advancing excitation and recovery to the
	*     next timtestep
	*/
	int index = -1;
	for (row = 1; row <= m; row++){
		for (col = 1; col <= n; col++) {
			E[row][col] = E[row][col] - dt * (kk * E[row][col] * (E[row][col] - a) * (E[row][col] - 1) + E[row][col] * R[row][col]);
			index = row * WIDTH + col;	
			E_1D[index] = E_1D[index] - dt * (kk * E_1D[index] * (E_1D[index] - a) * (E_1D[index] - 1) + E_1D[index] * R_1D[index]);

		}
	}

	for (row = 1; row <= m; row++){
		for (col = 1; col <= n; col++) {
			
			R[row][col] = R[row][col] + dt * (epsilon + M1 * R[row][col] / (E[row][col] + M2)) * (-R[row][col] - kk * E[row][col] * (E[row][col] - b - 1));

			index =  row * WIDTH + col;
			R_1D[index] = R_1D[index] + dt * (epsilon + M1 * R_1D[index] / (E_1D[index] + M2)) * (-R_1D[index] - kk * E_1D[index] * (E_1D[index] - b - 1));
		}
	}

}

// Main program
int main(int argc, char **argv){
	/*
	*  Solution arrays
	*   E is the "Excitation" variable, a voltage
	*   R is the "Recovery" variable
	*   E_prev is the Excitation variable for the previous timestep,
	*      and is used in time integration
	*/
    cout<< "\n Hello I am waris"<<endl;
	double **E, **R, **E_prev;

	double * E_1D, *R_1D, *E_prev_1D; 

	// Various constants - these definitions shouldn't change
	const double a = 0.1, b = 0.1, kk = 8.0, M1 = 0.07, M2 = 0.3, epsilon = 0.01, d = 5e-5;

	double T = 1000.0;
	int m = 200, n = 200;
	int plot_freq = 0;
	int px = 1, py = 1;
	int no_comm = 0;
	int num_threads = 1;
	int WIDTH;

	//int size;

	cmdLine(argc, argv, T, n, px, py, plot_freq, no_comm, num_threads);
	m = n;
	// Allocate contiguous memory for solution arrays
	// The computational box is defined on [1:m+1,1:n+1]
	// We pad the arrays in order to facilitate differencing on the
	// boundaries of the computation box
	int Total_Bytes = (m+2) * (n+2) *sizeof(double);
	WIDTH =  m+2;


	E = alloc2D(m + 2, n + 2);
	E_prev = alloc2D(m + 2, n + 2);
	R = alloc2D(m + 2, n + 2);
	
	E_1D = (double *) malloc (Total_Bytes);
	E_prev_1D = (double *) malloc (Total_Bytes);
    R_1D = (double *) malloc (Total_Bytes);
    
    


	int col, row;
	// Initialization
	
	for (row = 1; row <= m; row++) {
		for (col = 1; col <= n; col++) {
			E_prev[row][col] = 0;
			R[row][col] = 0;
			E_prev_1D[row*WIDTH + col] = 0;
			R_1D[row*WIDTH + col] = 0;
		}
	}
		
	for (row = 1; row <= m; row++) {
		for (col = n / 2 + 1; col <= n; col++) {
			E_prev[row][col] = 1.0;
			E_prev_1D[row*WIDTH + col] = 1.0;
		}
	}
		
	for (row = m / 2 + 1; row <= m; row++) {
		for (col = 1; col <= n; col++) {
			R[row][col] = 1.0;
			R_1D[row*WIDTH + col] = 1.0;
		}
	}

    //cout<< "check 1"<<endl;
		

	double dx = 1.0 / n;

	// For time integration, these values shouldn't change
	double rp = kk * (b + 1) * (b + 1) / 4;
	double dte = (dx * dx) / (d * 4 + ((dx * dx)) * (rp + kk));
	double dtr = 1 / (epsilon + ((M1 / M2) * rp));
	double dt = (dte < dtr) ? 0.95 * dte : 0.95 * dtr;
	double alpha = d * dt / (dx * dx);

	cout << "Grid Size       : " << n << endl;
	cout << "Duration of Sim : " << T << endl;
	cout << "Time step dt    : " << dt << endl;
	cout << "Process geometry: " << px << " x " << py << endl;
	
	if (no_comm) {
		cout << "Communication   : DISABLED" << endl;
	}

	cout << endl;

	// Start the timer
	double t0 = getTime();

	// Simulated time is different from the integer timestep number
	// Simulated time
	double t = 0.0;
	// Integer timestep number
	int niter = 0;

	while (t < T){

		t += dt;
		niter++;

		simulate(E, E_prev, R, alpha, n, m, kk, dt, a, epsilon, M1, M2, b, E_1D, E_prev_1D, R_1D, WIDTH);

		//swap current E with previous E
		double **tmp = E;
		E = E_prev;
		E_prev = tmp;

		double *tmp2 = E_1D;
		E_1D = E_prev_1D;
		E_prev_1D = tmp2;


		// if (plot_freq){
		// 	int k = (int)(t / plot_freq);
		// 	if ((t - k * plot_freq) < dt){
		// 		splot(E, t, niter, m + 2, n + 2);
		// 	}
		// }

	} //end of while loop

	double time_elapsed = getTime() - t0;

	double Gflops = (double)(niter * (1E-9 * n * n) * 28.0) / time_elapsed;
	double BW = (double)(niter * 1E-9 * (n * n * sizeof(double) * 4.0)) / time_elapsed;

	cout << "Number of Iterations        : " << niter << endl;
	cout << "Elapsed Time (sec)          : " << time_elapsed << endl;
	cout << "Sustained Gflops Rate       : " << Gflops << endl;
	cout << "Sustained Bandwidth (GB/sec): " << BW << endl << endl;

	double mx;
	double l2norm = stats(E_prev, m, n, &mx);
	cout << "Max: " << mx << " L2norm: " << l2norm << endl;

	l2norm = stats1D(E_prev_1D, m, n, &mx, WIDTH);
	cout << "1D Max : " << mx << " 1D L2norm: " << l2norm << endl;

	// if (plot_freq){
	// 	cout << "\n\nEnter any input to close the program and the plot..." << endl;
	// 	getchar();
	// }

	free(E);
	free(E_prev);
	free(R);

	free(E_1D);
	free(E_prev_1D);
	free(R_1D);

	return 0;
}


/* **********************************************************
*  Author : Urvashi R.V. [04/06/2004]
*      Modified by Didem Unat [03/23/18]
*************************************************************/

//#include <stdio.h>

/* Function to plot the 2D array
* 'gnuplot' is instantiated via a pipe and 
* the values to be plotted are passed through, along 
* with gnuplot commands */

// FILE *gnu = NULL;

// void splot(double **U, double T, int niter, int m, int n){
// 	int col, row;
// 	if (gnu == NULL)
// 		gnu = popen("gnuplot", "w");

// 	double mx = -1, mn = 32768;
// 	for (row = 0; row < m; row++)
// 		for (col = 0; col < n; col++)
// 		{
// 			if (U[row][col] > mx)
// 				mx = U[row][col];
// 			if (U[row][col] < mn)
// 				mn = U[row][col];
// 		}

// 	fprintf(gnu, "set title \"T = %f [niter = %d]\"\n", T, niter);
// 	fprintf(gnu, "set size square\n");
// 	fprintf(gnu, "set key off\n");
// 	fprintf(gnu, "set pm3d map\n");
// 	// Various color schemes
// 	fprintf(gnu, "set palette defined (-3 \"blue\", 0 \"white\", 1 \"red\")\n");

// 	//    fprintf(gnu,"set palette rgbformulae 22, 13, 31\n");
// 	//    fprintf(gnu,"set palette rgbformulae 30, 31, 32\n");

// 	fprintf(gnu, "splot [0:%d] [0:%d][%f:%f] \"-\"\n", m - 1, n - 1, mn, mx);
// 	for (row = 0; row < m; row++)
// 	{
// 		for (col = 0; col < n; col++)
// 		{
// 			fprintf(gnu, "%d %d %f\n", col, row, U[col][row]);
// 		}
// 		fprintf(gnu, "\n");
// 	}
// 	fprintf(gnu, "e\n");
// 	fflush(gnu);
// 	return;
// }
	.text
	.file	"vecAdd_UVM.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_ # -- Begin function _Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_
	.type	_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_,@function
_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_:      # @_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_
	.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	$24, %rsp
	.cfi_def_cfa_offset 80
	.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	%r8, 16(%rsp)                   # 8-byte Spill
	cmpl	$2, %edi
	jl	.LBB0_17
# %bb.1:                                # %.preheader.preheader
	movq	%r9, %rbx
	movq	%rcx, %r15
	movq	%rdx, %r12
	movq	%rsi, %r13
	movl	%edi, %ebp
	movl	$1, %eax
.LBB0_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_3 Depth 2
	movl	%eax, 12(%rsp)                  # 4-byte Spill
.LBB0_3:                                #   Parent Loop BB0_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	$.L.str.7, %edx
	movl	$_ZZ7cmdLineiPPcRdRiS2_S2_S2_S2_S2_E12long_options, %ecx
	movl	%ebp, %edi
	movq	%r13, %rsi
	xorl	%r8d, %r8d
	callq	getopt_long
                                        # kill: def $eax killed $eax def $rax
	leal	-107(%rax), %ecx
	cmpl	$14, %ecx
	ja	.LBB0_4
# %bb.6:                                #   in Loop: Header=BB0_3 Depth=2
	jmpq	*.LJTI0_0(,%rcx,8)
.LBB0_13:                               #   in Loop: Header=BB0_3 Depth=2
	movq	88(%rsp), %rax
	movl	$1, (%rax)
	jmp	.LBB0_3
.LBB0_8:                                # %._crit_edge15
                                        #   in Loop: Header=BB0_3 Depth=2
	movq	optarg(%rip), %r14
	jmp	.LBB0_11
.LBB0_14:                               #   in Loop: Header=BB0_3 Depth=2
	movq	optarg(%rip), %rdi
	callq	atoi
	movq	80(%rsp), %rcx
	movl	%eax, (%rcx)
	jmp	.LBB0_3
.LBB0_7:                                # %._crit_edge16
                                        #   in Loop: Header=BB0_3 Depth=2
	movq	optarg(%rip), %r14
	jmp	.LBB0_12
.LBB0_15:                               #   in Loop: Header=BB0_3 Depth=2
	movq	optarg(%rip), %rdi
	callq	atoi
	movq	96(%rsp), %rcx
	movl	%eax, (%rcx)
	jmp	.LBB0_3
.LBB0_10:                               #   in Loop: Header=BB0_3 Depth=2
	movq	optarg(%rip), %r14
	movq	%r14, %rdi
	callq	atoi
	movq	16(%rsp), %rcx                  # 8-byte Reload
	movl	%eax, (%rcx)
.LBB0_11:                               #   in Loop: Header=BB0_3 Depth=2
	movq	%r14, %rdi
	callq	atoi
	movl	%eax, (%rbx)
.LBB0_12:                               #   in Loop: Header=BB0_3 Depth=2
	movq	%r14, %rdi
	callq	atof
	movsd	%xmm0, (%r12)
	jmp	.LBB0_3
.LBB0_9:                                #   in Loop: Header=BB0_3 Depth=2
	movq	optarg(%rip), %rdi
	callq	atoi
	movl	%eax, (%r15)
	jmp	.LBB0_3
.LBB0_4:                                #   in Loop: Header=BB0_2 Depth=1
	cmpl	$-1, %eax
	jne	.LBB0_5
# %bb.16:                               #   in Loop: Header=BB0_2 Depth=1
	movl	12(%rsp), %eax                  # 4-byte Reload
	incl	%eax
	cmpl	%ebp, %eax
	jne	.LBB0_2
.LBB0_17:                               # %._crit_edge
	addq	$24, %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
.LBB0_5:
	.cfi_def_cfa_offset 80
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$-1, %edi
	callq	exit
.Lfunc_end0:
	.size	_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_, .Lfunc_end0-_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_
	.cfi_endproc
	.section	.rodata,"a",@progbits
	.p2align	3, 0x0
.LJTI0_0:
	.quad	.LBB0_13
	.quad	.LBB0_5
	.quad	.LBB0_5
	.quad	.LBB0_9
	.quad	.LBB0_15
	.quad	.LBB0_14
	.quad	.LBB0_5
	.quad	.LBB0_5
	.quad	.LBB0_5
	.quad	.LBB0_7
	.quad	.LBB0_5
	.quad	.LBB0_5
	.quad	.LBB0_5
	.quad	.LBB0_10
	.quad	.LBB0_8
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z7getTimev
.LCPI1_0:
	.quad	0x3eb0c6f7a0b5ed8d              # double 9.9999999999999995E-7
.LCPI1_1:
	.quad	0xbff0000000000000              # double -1
	.text
	.globl	_Z7getTimev
	.type	_Z7getTimev,@function
_Z7getTimev:                            # @_Z7getTimev
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	leaq	8(%rsp), %rdi
	leaq	24(%rsp), %rsi
	callq	gettimeofday
	cmpl	$-1, %eax
	je	.LBB1_1
# %bb.2:
	cvtsi2sdq	8(%rsp), %xmm1
	cvtsi2sdq	16(%rsp), %xmm0
	mulsd	.LCPI1_0(%rip), %xmm0
	addsd	%xmm1, %xmm0
	jmp	.LBB1_3
.LBB1_1:
	movl	$_ZSt4cerr, %ebx
	movl	$_ZSt4cerr, %edi
	movl	$.L.str.9, %esi
	movl	$31, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cerr(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cerr, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movsd	.LCPI1_1(%rip), %xmm0           # xmm0 = mem[0],zero
.LBB1_3:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z7getTimev, .Lfunc_end1-_Z7getTimev
	.cfi_endproc
                                        # -- End function
	.globl	_Z7alloc2Dii                    # -- Begin function _Z7alloc2Dii
	.type	_Z7alloc2Dii,@function
_Z7alloc2Dii:                           # @_Z7alloc2Dii
	.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
	movslq	%edi, %r14
	movslq	%esi, %rbx
	leaq	8(,%rbx,8), %rdi
	imulq	%r14, %rdi
	callq	malloc
	testl	%r14d, %r14d
	jle	.LBB2_3
# %bb.1:                                # %.lr.ph
	leaq	(%rax,%r14,8), %rcx
	movl	%r14d, %edx
	shlq	$3, %rbx
	xorl	%esi, %esi
.LBB2_2:                                # =>This Inner Loop Header: Depth=1
	movq	%rcx, (%rax,%rsi,8)
	incq	%rsi
	addq	%rbx, %rcx
	cmpq	%rsi, %rdx
	jne	.LBB2_2
.LBB2_3:                                # %._crit_edge
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z7alloc2Dii, .Lfunc_end2-_Z7alloc2Dii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z5statsPPdiiS_
.LCPI3_0:
	.quad	0xbff0000000000000              # double -1
	.text
	.globl	_Z5statsPPdiiS_
	.type	_Z5statsPPdiiS_,@function
_Z5statsPPdiiS_:                        # @_Z5statsPPdiiS_
	.cfi_startproc
# %bb.0:
                                        # kill: def $edx killed $edx def $rdx
                                        # kill: def $esi killed $esi def $rsi
	xorpd	%xmm1, %xmm1
	testl	%esi, %esi
	jle	.LBB3_1
# %bb.2:                                # %.preheader.lr.ph
	leal	1(%rdx), %eax
	leal	1(%rsi), %r8d
	movsd	.LCPI3_0(%rip), %xmm3           # xmm3 = mem[0],zero
	xorpd	%xmm0, %xmm0
	movl	$1, %r9d
.LBB3_3:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_6 Depth 2
	testl	%edx, %edx
	jle	.LBB3_4
# %bb.5:                                # %.lr.ph
                                        #   in Loop: Header=BB3_3 Depth=1
	movq	(%rdi,%r9,8), %r10
	movl	$1, %r11d
.LBB3_6:                                #   Parent Loop BB3_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%r10,%r11,8), %xmm2            # xmm2 = mem[0],zero
	movapd	%xmm2, %xmm4
	mulsd	%xmm2, %xmm4
	addsd	%xmm4, %xmm0
	maxsd	%xmm3, %xmm2
	incq	%r11
	movapd	%xmm2, %xmm3
	cmpq	%r11, %rax
	jne	.LBB3_6
	jmp	.LBB3_7
.LBB3_4:                                #   in Loop: Header=BB3_3 Depth=1
	movapd	%xmm3, %xmm2
.LBB3_7:                                # %._crit_edge
                                        #   in Loop: Header=BB3_3 Depth=1
	incq	%r9
	movapd	%xmm2, %xmm3
	cmpq	%r8, %r9
	jne	.LBB3_3
	jmp	.LBB3_8
.LBB3_1:
	movsd	.LCPI3_0(%rip), %xmm2           # xmm2 = mem[0],zero
	xorpd	%xmm0, %xmm0
.LBB3_8:                                # %._crit_edge42
	imull	%esi, %edx
	xorps	%xmm3, %xmm3
	cvtsi2sd	%edx, %xmm3
	movsd	%xmm2, (%rcx)
	divsd	%xmm3, %xmm0
	ucomisd	%xmm1, %xmm0
	jb	sqrt                            # TAILCALL
# %bb.9:                                # %._crit_edge42.split
	sqrtsd	%xmm0, %xmm0
	retq
.Lfunc_end3:
	.size	_Z5statsPPdiiS_, .Lfunc_end3-_Z5statsPPdiiS_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z7stats1DPdiiS_i
.LCPI4_0:
	.quad	0xbff0000000000000              # double -1
	.text
	.globl	_Z7stats1DPdiiS_i
	.type	_Z7stats1DPdiiS_i,@function
_Z7stats1DPdiiS_i:                      # @_Z7stats1DPdiiS_i
	.cfi_startproc
# %bb.0:
                                        # kill: def $edx killed $edx def $rdx
                                        # kill: def $esi killed $esi def $rsi
	xorpd	%xmm1, %xmm1
	testl	%esi, %esi
	jle	.LBB4_1
# %bb.2:                                # %.preheader.lr.ph
	leal	1(%rdx), %eax
	movslq	%r8d, %r8
	leal	1(%rsi), %r9d
	leaq	(%rdi,%r8,8), %rdi
	shlq	$3, %r8
	movsd	.LCPI4_0(%rip), %xmm3           # xmm3 = mem[0],zero
	xorpd	%xmm0, %xmm0
	movl	$1, %r10d
.LBB4_3:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_6 Depth 2
	testl	%edx, %edx
	jle	.LBB4_4
# %bb.5:                                # %.lr.ph
                                        #   in Loop: Header=BB4_3 Depth=1
	movl	$1, %r11d
.LBB4_6:                                #   Parent Loop BB4_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rdi,%r11,8), %xmm2            # xmm2 = mem[0],zero
	movapd	%xmm2, %xmm4
	mulsd	%xmm2, %xmm4
	addsd	%xmm4, %xmm0
	maxsd	%xmm3, %xmm2
	incq	%r11
	movapd	%xmm2, %xmm3
	cmpq	%r11, %rax
	jne	.LBB4_6
	jmp	.LBB4_7
.LBB4_4:                                #   in Loop: Header=BB4_3 Depth=1
	movapd	%xmm3, %xmm2
.LBB4_7:                                # %._crit_edge
                                        #   in Loop: Header=BB4_3 Depth=1
	incq	%r10
	addq	%r8, %rdi
	movapd	%xmm2, %xmm3
	cmpq	%r9, %r10
	jne	.LBB4_3
	jmp	.LBB4_8
.LBB4_1:
	movsd	.LCPI4_0(%rip), %xmm2           # xmm2 = mem[0],zero
	xorpd	%xmm0, %xmm0
.LBB4_8:                                # %._crit_edge41
	imull	%esi, %edx
	xorps	%xmm3, %xmm3
	cvtsi2sd	%edx, %xmm3
	movsd	%xmm2, (%rcx)
	divsd	%xmm3, %xmm0
	ucomisd	%xmm1, %xmm0
	jb	sqrt                            # TAILCALL
# %bb.9:                                # %._crit_edge41.split
	sqrtsd	%xmm0, %xmm0
	retq
.Lfunc_end4:
	.size	_Z7stats1DPdiiS_i, .Lfunc_end4-_Z7stats1DPdiiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z32__device_stub__simulate_version1diidddddddPdS_S_i # -- Begin function _Z32__device_stub__simulate_version1diidddddddPdS_S_i
	.type	_Z32__device_stub__simulate_version1diidddddddPdS_S_i,@function
_Z32__device_stub__simulate_version1diidddddddPdS_S_i: # @_Z32__device_stub__simulate_version1diidddddddPdS_S_i
	.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	$280, %rsp                      # imm = 0x118
	.cfi_def_cfa_offset 336
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	leaq	8(%rsp), %rax
	movsd	%xmm0, (%rax)
	leaq	28(%rsp), %r10
	movl	%edi, (%r10)
	leaq	24(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	120(%rsp), %rsi
	movsd	%xmm1, (%rsi)
	leaq	112(%rsp), %r11
	movsd	%xmm2, (%r11)
	leaq	104(%rsp), %r14
	movsd	%xmm3, (%r14)
	leaq	96(%rsp), %r15
	movsd	%xmm4, (%r15)
	leaq	88(%rsp), %r12
	movsd	%xmm5, (%r12)
	leaq	80(%rsp), %r13
	movsd	%xmm6, (%r13)
	leaq	72(%rsp), %rbp
	movsd	%xmm7, (%rbp)
	leaq	64(%rsp), %rax
	movq	%rdx, (%rax)
	leaq	56(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	48(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	20(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	160(%rsp), %rbx
	leaq	8(%rsp), %r9
	movq	%r9, (%rbx)
	movq	%r10, 8(%rbx)
	movq	%rdi, 16(%rbx)
	movq	%rsi, 24(%rbx)
	movq	%r11, 32(%rbx)
	movq	%r14, 40(%rbx)
	movq	%r15, 48(%rbx)
	movq	%r12, 56(%rbx)
	movq	%r13, 64(%rbx)
	movq	%rbp, 72(%rbx)
	movq	%rax, 80(%rbx)
	movq	%rdx, 88(%rbx)
	movq	%rcx, 96(%rbx)
	movq	%r8, 104(%rbx)
	leaq	144(%rsp), %r14
	leaq	128(%rsp), %r15
	leaq	40(%rsp), %r12
	leaq	32(%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	$_Z17simulate_version1diidddddddPdS_S_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$296, %rsp                      # imm = 0x128
	.cfi_adjust_cfa_offset -296
	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_end5:
	.size	_Z32__device_stub__simulate_version1diidddddddPdS_S_i, .Lfunc_end5-_Z32__device_stub__simulate_version1diidddddddPdS_S_i
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z8simulatePPdS0_S0_diidddddddS_S_S_i
.LCPI6_0:
	.quad	0x4010000000000000              # double 4
.LCPI6_1:
	.quad	0xbff0000000000000              # double -1
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0
.LCPI6_2:
	.quad	0x8000000000000000              # double -0
	.quad	0x8000000000000000              # double -0
	.text
	.globl	_Z8simulatePPdS0_S0_diidddddddS_S_S_i
	.type	_Z8simulatePPdS0_S0_diidddddddS_S_S_i,@function
_Z8simulatePPdS0_S0_diidddddddS_S_S_i:  # @_Z8simulatePPdS0_S0_diidddddddS_S_S_i
	.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
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
                                        # kill: def $r8d killed $r8d def $r8
                                        # kill: def $ecx killed $ecx def $rcx
	movq	%rdi, -72(%rsp)                 # 8-byte Spill
	movl	72(%rsp), %eax
	movq	56(%rsp), %rbp
	testl	%r8d, %r8d
	jle	.LBB6_6
# %bb.1:                                # %.lr.ph.preheader
	movslq	%eax, %r15
	leal	1(%r8), %edi
	leaq	(,%r15,8), %r11
	leaq	16(,%r15,8), %r14
	addq	%rbp, %r14
	movl	$1, %ebx
.LBB6_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	(%rsi,%rbx,8), %r10
	movsd	16(%r10), %xmm8                 # xmm8 = mem[0],zero
	movsd	%xmm8, (%r10)
	movsd	(%r14), %xmm8                   # xmm8 = mem[0],zero
	movsd	%xmm8, -16(%r14)
	incq	%rbx
	addq	%r11, %r14
	cmpq	%rbx, %rdi
	jne	.LBB6_2
# %bb.3:                                # %.preheader245
	testl	%r8d, %r8d
	jle	.LBB6_6
# %bb.4:                                # %.lr.ph250
	movslq	%ecx, %r11
	leaq	-1(%r11), %r13
	leaq	-8(,%r11,8), %r12
	addq	%rbp, %r12
	incq	%r11
	movl	$1, %r10d
	movq	%r15, %rbx
.LBB6_5:                                # =>This Inner Loop Header: Depth=1
	movq	(%rsi,%r10,8), %r14
	movsd	(%r14,%r13,8), %xmm8            # xmm8 = mem[0],zero
	movsd	%xmm8, (%r14,%r11,8)
	movsd	(%r12,%rbx,8), %xmm8            # xmm8 = mem[0],zero
	movsd	%xmm8, 16(%r12,%rbx,8)
	incq	%r10
	addq	%r15, %rbx
	cmpq	%r10, %rdi
	jne	.LBB6_5
.LBB6_6:                                # %.preheader244
	testl	%ecx, %ecx
	jle	.LBB6_12
# %bb.7:                                # %.lr.ph253
	movq	(%rsi), %r10
	movq	16(%rsi), %rdi
	leal	(%rax,%rax), %r11d
	movslq	%r11d, %r11
	leal	1(%rcx), %ebx
	leaq	(,%r11,8), %r14
	addq	%rbp, %r14
	movl	$1, %r11d
.LBB6_8:                                # =>This Inner Loop Header: Depth=1
	movsd	(%rdi,%r11,8), %xmm8            # xmm8 = mem[0],zero
	movsd	%xmm8, (%r10,%r11,8)
	movsd	(%r14,%r11,8), %xmm8            # xmm8 = mem[0],zero
	movsd	%xmm8, (%rbp,%r11,8)
	incq	%r11
	cmpq	%r11, %rbx
	jne	.LBB6_8
# %bb.9:                                # %.preheader243
	testl	%ecx, %ecx
	jle	.LBB6_12
# %bb.10:                               # %.lr.ph256
	movq	%rbp, %r15
	leal	-1(%r8), %ebp
	movslq	%r8d, %rdi
	movq	-8(%rsi,%rdi,8), %r11
	movq	8(%rsi,%rdi,8), %rdi
	leal	1(%r8), %r10d
	imull	%eax, %ebp
	imull	%eax, %r10d
	movslq	%ebp, %r14
	movq	%r15, %rbp
	movslq	%r10d, %r10
	leaq	(%r15,%r14,8), %r15
	leaq	(,%r10,8), %r14
	addq	%rbp, %r14
	movl	$1, %r10d
.LBB6_11:                               # =>This Inner Loop Header: Depth=1
	movsd	(%r11,%r10,8), %xmm8            # xmm8 = mem[0],zero
	movsd	%xmm8, (%rdi,%r10,8)
	movsd	(%r15,%r10,8), %xmm8            # xmm8 = mem[0],zero
	movsd	%xmm8, (%r14,%r10,8)
	incq	%r10
	cmpq	%r10, %rbx
	jne	.LBB6_11
.LBB6_12:                               # %.preheader242
	testl	%r8d, %r8d
	jle	.LBB6_28
# %bb.13:                               # %.preheader241.lr.ph
	movq	%rdx, -32(%rsp)                 # 8-byte Spill
	leal	1(%rcx), %edx
	cltq
	movq	%r8, -48(%rsp)                  # 8-byte Spill
	leal	1(%r8), %r13d
	leaq	8(%rbp), %rdi
	movq	%rdi, -16(%rsp)                 # 8-byte Spill
	movq	%rax, -64(%rsp)                 # 8-byte Spill
	leaq	(,%rax,8), %rdi
	addq	$16, %rbp
	movq	%rbp, -8(%rsp)                  # 8-byte Spill
	movq	%r9, -40(%rsp)                  # 8-byte Spill
	leaq	8(%r9), %rax
	movq	%rax, -24(%rsp)                 # 8-byte Spill
	movq	%rdx, -56(%rsp)                 # 8-byte Spill
	leaq	-8(,%rdx,8), %rbp
	movl	$1, %ebx
	movsd	.LCPI6_0(%rip), %xmm8           # xmm8 = mem[0],zero
.LBB6_14:                               # %.preheader241
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_30 Depth 2
	testl	%ecx, %ecx
	jle	.LBB6_15
# %bb.29:                               # %.lr.ph259
                                        #   in Loop: Header=BB6_14 Depth=1
	movq	%r13, %r12
	movq	%rcx, %r11
	movq	-8(%rsi,%rbx,8), %r9
	movq	(%rsi,%rbx,8), %r15
	leaq	1(%rbx), %r14
	movq	8(%rsi,%rbx,8), %r13
	leaq	-1(%rbx), %r8
	movq	-72(%rsp), %rax                 # 8-byte Reload
	movq	(%rax,%rbx,8), %rdx
	imulq	%rdi, %r8
	movq	-16(%rsp), %rax                 # 8-byte Reload
	addq	%rax, %r8
	movq	%rdi, %r10
	imulq	%r14, %r10
	addq	%rax, %r10
	imulq	%rdi, %rbx
	movq	-8(%rsp), %rax                  # 8-byte Reload
	leaq	(%rax,%rbx), %rcx
	addq	-24(%rsp), %rbx                 # 8-byte Folded Reload
	xorl	%eax, %eax
.LBB6_30:                               #   Parent Loop BB6_14 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	16(%r15,%rax), %xmm9            # xmm9 = mem[0],zero
	addsd	(%r15,%rax), %xmm9
	movsd	8(%r15,%rax), %xmm10            # xmm10 = mem[0],zero
	movapd	%xmm10, %xmm11
	mulsd	%xmm8, %xmm11
	subsd	%xmm11, %xmm9
	addsd	8(%r13,%rax), %xmm9
	addsd	8(%r9,%rax), %xmm9
	mulsd	%xmm0, %xmm9
	addsd	%xmm10, %xmm9
	movsd	%xmm9, 8(%rdx,%rax)
	movsd	-8(%rcx,%rax), %xmm9            # xmm9 = mem[0],zero
	movsd	(%rcx,%rax), %xmm10             # xmm10 = mem[0],zero
	addsd	-16(%rcx,%rax), %xmm10
	movapd	%xmm9, %xmm11
	mulsd	%xmm8, %xmm11
	subsd	%xmm11, %xmm10
	addsd	(%r10,%rax), %xmm10
	addsd	(%r8,%rax), %xmm10
	mulsd	%xmm0, %xmm10
	addsd	%xmm9, %xmm10
	movsd	%xmm10, (%rbx,%rax)
	addq	$8, %rax
	cmpq	%rax, %rbp
	jne	.LBB6_30
# %bb.31:                               #   in Loop: Header=BB6_14 Depth=1
	movq	%r14, %rbx
	movq	%r11, %rcx
	movq	%r12, %r13
	jmp	.LBB6_32
.LBB6_15:                               # %.preheader241.._crit_edge_crit_edge
                                        #   in Loop: Header=BB6_14 Depth=1
	incq	%rbx
	movq	%rbx, %r14
.LBB6_32:                               # %._crit_edge
                                        #   in Loop: Header=BB6_14 Depth=1
	cmpq	%r13, %r14
	jne	.LBB6_14
# %bb.16:                               # %.preheader240
	movq	-48(%rsp), %r14                 # 8-byte Reload
	testl	%r14d, %r14d
	movq	-32(%rsp), %r10                 # 8-byte Reload
	movq	-40(%rsp), %r9                  # 8-byte Reload
	movq	64(%rsp), %r15
	movq	-56(%rsp), %rbp                 # 8-byte Reload
	jle	.LBB6_28
# %bb.17:                               # %.preheader239.lr.ph
	movq	-64(%rsp), %rax                 # 8-byte Reload
	leaq	(%r9,%rax,8), %rsi
	addq	$8, %rsi
	leaq	(%r15,%rax,8), %r11
	addq	$8, %r11
	leaq	-8(,%rbp,8), %rbx
	movl	$1, %r12d
	movsd	.LCPI6_1(%rip), %xmm0           # xmm0 = mem[0],zero
.LBB6_18:                               # %.preheader239
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_20 Depth 2
	testl	%ecx, %ecx
	jle	.LBB6_21
# %bb.19:                               # %.lr.ph264
                                        #   in Loop: Header=BB6_18 Depth=1
	movq	-72(%rsp), %rax                 # 8-byte Reload
	movq	(%rax,%r12,8), %rdx
	movq	(%r10,%r12,8), %rax
	xorl	%r8d, %r8d
.LBB6_20:                               #   Parent Loop BB6_18 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	8(%rdx,%r8), %xmm8              # xmm8 = mem[0],zero
	movapd	%xmm8, %xmm9
	mulsd	%xmm1, %xmm9
	movapd	%xmm8, %xmm10
	subsd	%xmm3, %xmm10
	mulsd	%xmm9, %xmm10
	movapd	%xmm8, %xmm9
	addsd	%xmm0, %xmm9
	mulsd	%xmm10, %xmm9
	movsd	8(%rax,%r8), %xmm10             # xmm10 = mem[0],zero
	mulsd	%xmm8, %xmm10
	addsd	%xmm9, %xmm10
	mulsd	%xmm2, %xmm10
	subsd	%xmm10, %xmm8
	movsd	%xmm8, 8(%rdx,%r8)
	movsd	(%rsi,%r8), %xmm8               # xmm8 = mem[0],zero
	movapd	%xmm8, %xmm9
	mulsd	%xmm1, %xmm9
	movapd	%xmm8, %xmm10
	subsd	%xmm3, %xmm10
	mulsd	%xmm9, %xmm10
	movapd	%xmm8, %xmm9
	addsd	%xmm0, %xmm9
	mulsd	%xmm10, %xmm9
	movsd	(%r11,%r8), %xmm10              # xmm10 = mem[0],zero
	mulsd	%xmm8, %xmm10
	addsd	%xmm9, %xmm10
	mulsd	%xmm2, %xmm10
	subsd	%xmm10, %xmm8
	movsd	%xmm8, (%rsi,%r8)
	addq	$8, %r8
	cmpq	%r8, %rbx
	jne	.LBB6_20
.LBB6_21:                               # %._crit_edge265
                                        #   in Loop: Header=BB6_18 Depth=1
	incq	%r12
	addq	%rdi, %rsi
	addq	%rdi, %r11
	cmpq	%r13, %r12
	jne	.LBB6_18
# %bb.22:                               # %.preheader238
	testl	%r14d, %r14d
	jle	.LBB6_28
# %bb.23:                               # %.preheader.lr.ph
	movq	-64(%rsp), %rax                 # 8-byte Reload
	leaq	(%r15,%rax,8), %rsi
	addq	$8, %rsi
	leaq	8(%r9,%rax,8), %r8
	leaq	-8(,%rbp,8), %r9
	movl	$1, %r11d
	movapd	.LCPI6_2(%rip), %xmm3           # xmm3 = [-0.0E+0,-0.0E+0]
	movq	-72(%rsp), %rax                 # 8-byte Reload
.LBB6_24:                               # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_26 Depth 2
	testl	%ecx, %ecx
	jle	.LBB6_27
# %bb.25:                               # %.lr.ph270
                                        #   in Loop: Header=BB6_24 Depth=1
	movq	(%r10,%r11,8), %rbx
	movq	(%rax,%r11,8), %r14
	xorl	%r15d, %r15d
.LBB6_26:                               #   Parent Loop BB6_24 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	8(%rbx,%r15), %xmm8             # xmm8 = mem[0],zero
	movapd	%xmm8, %xmm9
	mulsd	%xmm5, %xmm9
	movsd	8(%r14,%r15), %xmm10            # xmm10 = mem[0],zero
	movapd	%xmm10, %xmm11
	addsd	%xmm6, %xmm11
	divsd	%xmm11, %xmm9
	addsd	%xmm4, %xmm9
	mulsd	%xmm2, %xmm9
	movapd	%xmm8, %xmm11
	xorpd	%xmm3, %xmm11
	movapd	%xmm10, %xmm12
	mulsd	%xmm1, %xmm12
	subsd	%xmm7, %xmm10
	addsd	%xmm0, %xmm10
	mulsd	%xmm12, %xmm10
	subsd	%xmm10, %xmm11
	mulsd	%xmm9, %xmm11
	addsd	%xmm8, %xmm11
	movsd	%xmm11, 8(%rbx,%r15)
	movsd	(%rsi,%r15), %xmm8              # xmm8 = mem[0],zero
	movapd	%xmm8, %xmm9
	mulsd	%xmm5, %xmm9
	movsd	(%r8,%r15), %xmm10              # xmm10 = mem[0],zero
	movapd	%xmm10, %xmm11
	addsd	%xmm6, %xmm11
	divsd	%xmm11, %xmm9
	addsd	%xmm4, %xmm9
	mulsd	%xmm2, %xmm9
	movapd	%xmm8, %xmm11
	xorpd	%xmm3, %xmm11
	movapd	%xmm10, %xmm12
	mulsd	%xmm1, %xmm12
	subsd	%xmm7, %xmm10
	addsd	%xmm0, %xmm10
	mulsd	%xmm12, %xmm10
	subsd	%xmm10, %xmm11
	mulsd	%xmm9, %xmm11
	addsd	%xmm8, %xmm11
	movsd	%xmm11, (%rsi,%r15)
	addq	$8, %r15
	cmpq	%r15, %r9
	jne	.LBB6_26
.LBB6_27:                               # %._crit_edge271
                                        #   in Loop: Header=BB6_24 Depth=1
	incq	%r11
	addq	%rdi, %rsi
	addq	%rdi, %r8
	cmpq	%r13, %r11
	jne	.LBB6_24
.LBB6_28:                               # %._crit_edge274
	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_end6:
	.size	_Z8simulatePPdS0_S0_diidddddddS_S_S_i, .Lfunc_end6-_Z8simulatePPdS0_S0_diidddddddS_S_S_i
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI7_0:
	.quad	0x3ff0000000000000              # double 1
.LCPI7_1:
	.quad	0x4024d70a3d70a3d7              # double 10.42
.LCPI7_2:
	.quad	0x3f2a36e2eb1c432d              # double 2.0000000000000001E-4
.LCPI7_3:
	.quad	0x3fee666666666666              # double 0.94999999999999996
.LCPI7_4:
	.quad	0x3ffbd79c36a51f78              # double 1.7401392111368903
.LCPI7_5:
	.quad	0x3ffa733acd83444b              # double 1.6531322505800456
.LCPI7_6:
	.quad	0x3f0a36e2eb1c432d              # double 5.0000000000000002E-5
.LCPI7_7:
	.quad	0x4020000000000000              # double 8
.LCPI7_8:
	.quad	0x3f847ae147ae147b              # double 0.01
.LCPI7_9:
	.quad	0x3fb1eb851eb851ec              # double 0.070000000000000007
.LCPI7_10:
	.quad	0x3fd3333333333333              # double 0.29999999999999999
.LCPI7_11:
	.quad	0x3fb999999999999a              # double 0.10000000000000001
.LCPI7_12:
	.quad	0x3e112e0be826d695              # double 1.0000000000000001E-9
.LCPI7_13:
	.quad	0x403c000000000000              # double 28
.LCPI7_14:
	.quad	0x4010000000000000              # double 4
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.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	$184, %rsp
	.cfi_def_cfa_offset 240
	.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	%rsi, %rbx
	movl	%edi, %ebp
	movl	$_ZSt4cout, %r14d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.10, %esi
	movl	$18, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movabsq	$4652007308841189376, %rax      # imm = 0x408F400000000000
	leaq	160(%rsp), %rdx
	movq	%rax, (%rdx)
	leaq	180(%rsp), %r14
	movl	$200, (%r14)
	xorl	%eax, %eax
	leaq	176(%rsp), %r10
	movl	%eax, (%r10)
	movl	$1, %ecx
	leaq	92(%rsp), %r8
	movl	%ecx, (%r8)
	leaq	88(%rsp), %r9
	movl	%ecx, (%r9)
	leaq	84(%rsp), %r11
	movl	%eax, (%r11)
	leaq	172(%rsp), %rax
	movl	%ecx, (%rax)
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	%ebp, %edi
	movq	%rbx, %rsi
	movq	%r14, %rcx
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	callq	_Z7cmdLineiPPcRdRiS2_S2_S2_S2_S2_
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	movslq	(%r14), %r13
	leaq	2(%r13), %rbp
	movl	%ebp, %ebx
	imull	%ebp, %ebx
	shll	$3, %ebx
	movl	%ebp, %edi
	movl	%ebp, %esi
	callq	_Z7alloc2Dii
	movq	%rax, 104(%rsp)                 # 8-byte Spill
	movl	%ebp, %edi
	movl	%ebp, %esi
	callq	_Z7alloc2Dii
	movq	%rax, %r14
	movl	%ebp, %edi
	movl	%ebp, %esi
	callq	_Z7alloc2Dii
	movq	%rax, 56(%rsp)                  # 8-byte Spill
	movslq	%ebx, %rbx
	movq	%rbx, %rdi
	callq	malloc
	movq	%rax, 96(%rsp)                  # 8-byte Spill
	movq	%rbx, %rdi
	callq	malloc
	movq	%rax, %r15
	movq	%rbx, %rdi
	callq	malloc
	movq	%rax, %r12
	movl	%r13d, %eax
	movq	%rax, (%rsp)                    # 8-byte Spill
	movq	%r13, 64(%rsp)                  # 8-byte Spill
	testq	%r13, %r13
	movq	%r15, 72(%rsp)                  # 8-byte Spill
	movq	%rbp, 40(%rsp)                  # 8-byte Spill
	movq	%r12, 48(%rsp)                  # 8-byte Spill
	jle	.LBB7_10
# %bb.1:                                # %.preheader126.lr.ph
	movq	%r14, 32(%rsp)                  # 8-byte Spill
	movq	(%rsp), %rax                    # 8-byte Reload
	leal	3(%rax), %ecx
	movl	%ecx, 112(%rsp)                 # 4-byte Spill
	leaq	(,%rax,8), %rcx
	movq	%rcx, 16(%rsp)                  # 8-byte Spill
	leaq	1(%rax), %r13
	movl	$1, %r14d
	xorl	%r12d, %r12d
	xorl	%ebp, %ebp
.LBB7_2:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB7_3 Depth 2
	movq	40(%rsp), %rax                  # 8-byte Reload
                                        # kill: def $eax killed $eax killed $rax def $rax
	imull	%ebp, %eax
	addl	112(%rsp), %eax                 # 4-byte Folded Reload
	leaq	(%r15,%rax,8), %rdi
	movq	48(%rsp), %rcx                  # 8-byte Reload
	leaq	(%rcx,%rax,8), %rax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	movq	32(%rsp), %rax                  # 8-byte Reload
	movq	(%rax,%r14,8), %rbx
	movq	56(%rsp), %rax                  # 8-byte Reload
	movq	(%rax,%r14,8), %r15
	xorl	%esi, %esi
	movq	16(%rsp), %rdx                  # 8-byte Reload
	callq	memset@PLT
	movq	8(%rsp), %rdi                   # 8-byte Reload
	xorl	%esi, %esi
	movq	16(%rsp), %rdx                  # 8-byte Reload
	callq	memset@PLT
	movl	$1, %eax
.LBB7_3:                                #   Parent Loop BB7_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%r12, (%rbx,%rax,8)
	movq	%r12, (%r15,%rax,8)
	incq	%rax
	cmpq	%rax, %r13
	jne	.LBB7_3
# %bb.4:                                # %._crit_edge
                                        #   in Loop: Header=BB7_2 Depth=1
	incq	%r14
	incq	%rbp
	movq	(%rsp), %rax                    # 8-byte Reload
	cmpq	%rax, %rbp
	movq	72(%rsp), %r15                  # 8-byte Reload
	jne	.LBB7_2
# %bb.5:                                # %.preheader125
	testl	%eax, %eax
	movq	32(%rsp), %r14                  # 8-byte Reload
	movq	48(%rsp), %r12                  # 8-byte Reload
	movq	40(%rsp), %rbp                  # 8-byte Reload
	jle	.LBB7_10
# %bb.6:                                # %.lr.ph138
	movq	(%rsp), %rax                    # 8-byte Reload
                                        # kill: def $eax killed $eax killed $rax def $rax
	shrl	%eax
	leaq	8(,%rax,8), %rcx
	leaq	(%r15,%rax,8), %rdx
	addq	$8, %rdx
	movl	$1, %esi
	movabsq	$4607182418800017408, %rdi      # imm = 0x3FF0000000000000
	movl	%ebp, %r8d
	movq	64(%rsp), %r15                  # 8-byte Reload
.LBB7_7:                                # %.lr.ph134
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB7_8 Depth 2
	movl	%r8d, %r9d
	leaq	(%rdx,%r9,8), %r9
	movq	(%r14,%rsi,8), %r10
	addq	%rcx, %r10
	xorl	%r11d, %r11d
.LBB7_8:                                #   Parent Loop BB7_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%rdi, (%r10,%r11,8)
	movq	%rdi, (%r9,%r11,8)
	incq	%r11
	leal	(%rax,%r11), %ebx
	cmpl	%ebx, %r15d
	jg	.LBB7_8
# %bb.9:                                # %._crit_edge135
                                        #   in Loop: Header=BB7_7 Depth=1
	incq	%rsi
	addl	%ebp, %r8d
	cmpq	%r13, %rsi
	jne	.LBB7_7
.LBB7_10:                               # %._crit_edge139
	movq	(%rsp), %rbx                    # 8-byte Reload
	movl	%ebx, %ecx
	shrl	$31, %ecx
	addl	%ebx, %ecx
	sarl	%ecx
	cmpl	%ebx, %ecx
	movq	64(%rsp), %r10                  # 8-byte Reload
	movq	56(%rsp), %r11                  # 8-byte Reload
	jge	.LBB7_16
# %bb.11:                               # %.preheader.lr.ph
	leal	1(%rbx), %eax
	movslq	%ecx, %rcx
	leaq	1(%rcx), %rdx
	imulq	%rdx, %rbp
	leaq	(%r12,%rbp,8), %rdx
	leaq	16(,%r10,8), %rsi
	movabsq	$4607182418800017408, %rdi      # imm = 0x3FF0000000000000
.LBB7_13:                               # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB7_15 Depth 2
	incq	%rcx
	testl	%ebx, %ebx
	jle	.LBB7_12
# %bb.14:                               # %.lr.ph142
                                        #   in Loop: Header=BB7_13 Depth=1
	movq	(%r11,%rcx,8), %r8
	movl	$1, %r9d
.LBB7_15:                               #   Parent Loop BB7_13 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%rdi, (%r8,%r9,8)
	movq	%rdi, (%rdx,%r9,8)
	incq	%r9
	cmpq	%r9, %rax
	jne	.LBB7_15
.LBB7_12:                               # %.loopexit
                                        #   in Loop: Header=BB7_13 Depth=1
	addq	%rsi, %rdx
	cmpq	%r10, %rcx
	jne	.LBB7_13
.LBB7_16:                               # %._crit_edge146
	cvtsi2sd	%ebx, %xmm0
	movsd	.LCPI7_0(%rip), %xmm1           # xmm1 = mem[0],zero
	movsd	%xmm0, 144(%rsp)                # 8-byte Spill
	divsd	%xmm0, %xmm1
	mulsd	%xmm1, %xmm1
	movsd	.LCPI7_1(%rip), %xmm0           # xmm0 = mem[0],zero
	mulsd	%xmm1, %xmm0
	addsd	.LCPI7_2(%rip), %xmm0
	movsd	%xmm1, 8(%rsp)                  # 8-byte Spill
	movapd	%xmm1, %xmm2
	divsd	%xmm0, %xmm2
	movsd	.LCPI7_3(%rip), %xmm0           # xmm0 = mem[0],zero
	mulsd	%xmm2, %xmm0
	cmpltsd	.LCPI7_4(%rip), %xmm2
	andpd	%xmm2, %xmm0
	movsd	.LCPI7_5(%rip), %xmm1           # xmm1 = mem[0],zero
	andnpd	%xmm1, %xmm2
	orpd	%xmm0, %xmm2
	movapd	%xmm2, 16(%rsp)                 # 16-byte Spill
	movl	$_ZSt4cout, %r15d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.11, %esi
	movl	$18, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%ebx, %esi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.12, %esi
	movl	$18, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movsd	160(%rsp), %xmm0                # xmm0 = mem[0],zero
	movl	$_ZSt4cout, %edi
	movsd	%xmm0, 32(%rsp)                 # 8-byte Spill
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.13, %esi
	movl	$18, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movapd	16(%rsp), %xmm0                 # 16-byte Reload
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.14, %esi
	movl	$18, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	92(%rsp), %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movl	$.L.str.15, %esi
	movl	$3, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	88(%rsp), %esi
	movq	%rbx, %rdi
	callq	_ZNSolsEi
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	cmpl	$0, 84(%rsp)
	je	.LBB7_18
# %bb.17:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.16, %esi
	movl	$26, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB7_18:
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	callq	_Z7getTimev
	movapd	%xmm0, %xmm1
	xorl	%r15d, %r15d
	xorpd	%xmm2, %xmm2
	movsd	32(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	ucomisd	%xmm2, %xmm0
	movsd	%xmm1, 136(%rsp)                # 8-byte Spill
	jbe	.LBB7_19
# %bb.20:                               # %.lr.ph154.preheader
	movapd	16(%rsp), %xmm0                 # 16-byte Reload
	movapd	%xmm0, %xmm1
	mulsd	.LCPI7_6(%rip), %xmm1
	divsd	8(%rsp), %xmm1                  # 8-byte Folded Reload
	movapd	%xmm1, 112(%rsp)                # 16-byte Spill
	xorl	%r15d, %r15d
	movq	72(%rsp), %r13                  # 8-byte Reload
	movq	(%rsp), %r12                    # 8-byte Reload
	movq	104(%rsp), %rax                 # 8-byte Reload
	movq	96(%rsp), %rcx                  # 8-byte Reload
.LBB7_21:                               # %.lr.ph154
                                        # =>This Inner Loop Header: Depth=1
	movq	%r13, %rbp
	movq	%rcx, %r13
	movq	%r14, %rbx
	movq	%rax, %r14
	addsd	%xmm0, %xmm2
	movsd	%xmm2, 8(%rsp)                  # 8-byte Spill
	incl	%r15d
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movq	%rax, %rdi
	movq	%rbx, %rsi
	movq	64(%rsp), %rdx                  # 8-byte Reload
	movaps	120(%rsp), %xmm0                # 16-byte Reload
	movl	%r12d, %ecx
	movl	%r12d, %r8d
	movsd	.LCPI7_7(%rip), %xmm1           # xmm1 = mem[0],zero
	movaps	24(%rsp), %xmm2                 # 16-byte Reload
	movsd	.LCPI7_11(%rip), %xmm3          # xmm3 = mem[0],zero
	movsd	.LCPI7_8(%rip), %xmm4           # xmm4 = mem[0],zero
	movsd	.LCPI7_9(%rip), %xmm5           # xmm5 = mem[0],zero
	movsd	.LCPI7_10(%rip), %xmm6          # xmm6 = mem[0],zero
	movaps	%xmm3, %xmm7
	movq	%r13, %r9
	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	%rbp
	.cfi_adjust_cfa_offset 8
	callq	_Z8simulatePPdS0_S0_diidddddddS_S_S_i
	movsd	40(%rsp), %xmm2                 # 8-byte Reload
                                        # xmm2 = mem[0],zero
	movapd	48(%rsp), %xmm0                 # 16-byte Reload
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	movsd	32(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	ucomisd	%xmm2, %xmm1
	movq	%rbx, %rax
	movq	%rbp, %rcx
	ja	.LBB7_21
	jmp	.LBB7_22
.LBB7_19:
	movq	96(%rsp), %rbp                  # 8-byte Reload
	movq	104(%rsp), %rbx                 # 8-byte Reload
	movq	72(%rsp), %r13                  # 8-byte Reload
	movq	(%rsp), %r12                    # 8-byte Reload
.LBB7_22:                               # %._crit_edge155
	callq	_Z7getTimev
	xorps	%xmm1, %xmm1
	cvtsi2sd	%r15d, %xmm1
	movsd	.LCPI7_12(%rip), %xmm4          # xmm4 = mem[0],zero
	movsd	144(%rsp), %xmm2                # 8-byte Reload
                                        # xmm2 = mem[0],zero
	movapd	%xmm2, %xmm3
	mulsd	%xmm4, %xmm3
	mulsd	%xmm2, %xmm3
	mulsd	%xmm1, %xmm3
	mulsd	.LCPI7_13(%rip), %xmm3
	subsd	136(%rsp), %xmm0                # 8-byte Folded Reload
	divsd	%xmm0, %xmm3
	movapd	%xmm0, %xmm2
	movsd	%xmm0, (%rsp)                   # 8-byte Spill
	movsd	%xmm3, 8(%rsp)                  # 8-byte Spill
	mulsd	%xmm4, %xmm1
	movl	%r12d, %eax
	imull	%r12d, %eax
	shlq	$3, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	mulsd	.LCPI7_14(%rip), %xmm0
	mulsd	%xmm1, %xmm0
	divsd	%xmm2, %xmm0
	movsd	%xmm0, 16(%rsp)                 # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.17, %esi
	movl	$30, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%r15d, %esi
	callq	_ZNSolsEi
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.18, %esi
	movl	$30, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	(%rsp), %xmm0                   # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.19, %esi
	movl	$30, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.20, %esi
	movl	$30, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movsd	16(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	leaq	152(%rsp), %rax
	movq	%r14, %rdi
	movl	%r12d, %esi
	movq	64(%rsp), %r15                  # 8-byte Reload
	movl	%r15d, %edx
	movq	%rax, %r12
	movq	%rax, %rcx
	callq	_Z5statsPPdiiS_
	movsd	%xmm0, 16(%rsp)                 # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.21, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movsd	(%r12), %xmm0                   # xmm0 = mem[0],zero
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r12
	movl	$.L.str.22, %esi
	movl	$9, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r12, %rdi
	movsd	16(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r12
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r12, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	%r13, %rdi
	movl	%r15d, %esi
	movl	%r15d, %edx
	leaq	152(%rsp), %r15
	movq	%r15, %rcx
	movq	40(%rsp), %r8                   # 8-byte Reload
                                        # kill: def $r8d killed $r8d killed $r8
	callq	_Z7stats1DPdiiS_i
	movsd	%xmm0, 16(%rsp)                 # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.23, %esi
	movl	$9, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movsd	(%r15), %xmm0                   # xmm0 = mem[0],zero
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movl	$.L.str.24, %esi
	movl	$12, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r15, %rdi
	movsd	16(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	%rbx, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	56(%rsp), %rdi                  # 8-byte Reload
	callq	free
	movq	%rbp, %rdi
	callq	free
	movq	%r13, %rdi
	callq	free
	movq	48(%rsp), %rdi                  # 8-byte Reload
	callq	free
	xorl	%eax, %eax
	addq	$184, %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_end7:
	.size	main, .Lfunc_end7-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB8_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB8_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z17simulate_version1diidddddddPdS_S_i, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end8:
	.size	__hip_module_ctor, .Lfunc_end8-__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	.LBB9_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
.LBB9_2:
	retq
.Lfunc_end9:
	.size	__hip_module_dtor, .Lfunc_end9-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_ZZ7cmdLineiPPcRdRiS2_S2_S2_S2_S2_E12long_options,@object # @_ZZ7cmdLineiPPcRdRiS2_S2_S2_S2_S2_E12long_options
	.data
	.p2align	4, 0x0
_ZZ7cmdLineiPPcRdRiS2_S2_S2_S2_S2_E12long_options:
	.quad	.L.str
	.long	1                               # 0x1
	.zero	4
	.quad	0
	.long	110                             # 0x6e
	.zero	4
	.quad	.L.str.1
	.long	1                               # 0x1
	.zero	4
	.quad	0
	.long	120                             # 0x78
	.zero	4
	.quad	.L.str.2
	.long	1                               # 0x1
	.zero	4
	.quad	0
	.long	121                             # 0x79
	.zero	4
	.quad	.L.str.3
	.long	1                               # 0x1
	.zero	4
	.quad	0
	.long	116                             # 0x74
	.zero	4
	.quad	.L.str.4
	.long	1                               # 0x1
	.zero	4
	.quad	0
	.long	112                             # 0x70
	.zero	4
	.quad	.L.str.5
	.long	0                               # 0x0
	.zero	4
	.quad	0
	.long	107                             # 0x6b
	.zero	4
	.quad	.L.str.6
	.long	1                               # 0x1
	.zero	4
	.quad	0
	.long	111                             # 0x6f
	.zero	4
	.size	_ZZ7cmdLineiPPcRdRiS2_S2_S2_S2_S2_E12long_options, 224

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

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"px"
	.size	.L.str.1, 3

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"py"
	.size	.L.str.2, 3

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"tfinal"
	.size	.L.str.3, 7

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"plot"
	.size	.L.str.4, 5

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"nocomm"
	.size	.L.str.5, 7

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"numthreads"
	.size	.L.str.6, 11

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"n:x:y:t:kp:o:"
	.size	.L.str.7, 14

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"ERROR: Bad call to gettimeofday"
	.size	.L.str.9, 32

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

	.type	.L.str.10,@object               # @.str.10
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.10:
	.asciz	"\n Hello I am waris"
	.size	.L.str.10, 19

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Grid Size       : "
	.size	.L.str.11, 19

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"Duration of Sim : "
	.size	.L.str.12, 19

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"Time step dt    : "
	.size	.L.str.13, 19

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"Process geometry: "
	.size	.L.str.14, 19

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	" x "
	.size	.L.str.15, 4

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"Communication   : DISABLED"
	.size	.L.str.16, 27

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"Number of Iterations        : "
	.size	.L.str.17, 31

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"Elapsed Time (sec)          : "
	.size	.L.str.18, 31

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"Sustained Gflops Rate       : "
	.size	.L.str.19, 31

	.type	.L.str.20,@object               # @.str.20
.L.str.20:
	.asciz	"Sustained Bandwidth (GB/sec): "
	.size	.L.str.20, 31

	.type	.L.str.21,@object               # @.str.21
.L.str.21:
	.asciz	"Max: "
	.size	.L.str.21, 6

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	" L2norm: "
	.size	.L.str.22, 10

	.type	.L.str.23,@object               # @.str.23
.L.str.23:
	.asciz	"1D Max : "
	.size	.L.str.23, 10

	.type	.L.str.24,@object               # @.str.24
.L.str.24:
	.asciz	" 1D L2norm: "
	.size	.L.str.24, 13

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z17simulate_version1diidddddddPdS_S_i"
	.size	.L__unnamed_1, 39

	.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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"Usage: a.out [-n <domain size>] [-t <final time >]\n\t [-p <plot frequency>]\n\t[-px <x processor geometry> [-py <y proc. geometry] [-k turn off communication] [-o <Number of OpenMP threads>]"
	.size	.Lstr, 188

	.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 _Z32__device_stub__simulate_version1diidddddddPdS_S_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _ZZ7cmdLineiPPcRdRiS2_S2_S2_S2_S2_E12long_options
	.addrsig_sym _ZSt4cerr
	.addrsig_sym _Z17simulate_version1diidddddddPdS_S_i
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z17simulate_version1diidddddddPdS_S_i ; -- Begin function _Z17simulate_version1diidddddddPdS_S_i
	.globl	_Z17simulate_version1diidddddddPdS_S_i
	.p2align	8
	.type	_Z17simulate_version1diidddddddPdS_S_i,@function
_Z17simulate_version1diidddddddPdS_S_i: ; @_Z17simulate_version1diidddddddPdS_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x74
	s_load_b32 s2, s[0:1], 0x60
	v_bfe_u32 v2, v0, 10, 10
	v_and_b32_e32 v4, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s4, s3, 16
	s_and_b32 s3, s3, 0xffff
	v_mad_u64_u32 v[0:1], null, s15, s4, v[2:3]
	v_mad_u64_u32 v[2:3], null, s14, s3, v[4:5]
	s_add_i32 s3, s2, -1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v3, 1, v0
	v_add_nc_u32_e32 v1, 1, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max_i32_e32 v4, v3, v1
	v_cmp_gt_i32_e32 vcc_lo, s3, v4
	s_and_saveexec_b32 s3, vcc_lo
	s_cbranch_execz .LBB0_2
; %bb.1:
	v_mul_lo_u32 v6, v3, s2
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x48
	s_load_b64 s[0:1], s[0:1], 0x0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, v6, v2
	v_add_nc_u32_e32 v6, v6, v1
	v_add_nc_u32_e32 v4, 2, v2
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_ashrrev_i32_e32 v7, 31, v6
	v_add_nc_u32_e32 v11, s2, v6
	v_ashrrev_i32_e32 v5, 31, v4
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[2:3], 3, v[2:3]
	v_lshlrev_b64 v[7:8], 3, v[6:7]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_ashrrev_i32_e32 v12, 31, v11
	v_lshlrev_b64 v[4:5], 3, v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[11:12], 3, v[11:12]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v4, vcc_lo, s6, v4
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v5, vcc_lo
	v_add_co_u32 v2, vcc_lo, s6, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s7, v3, vcc_lo
	v_add_co_u32 v9, vcc_lo, s6, v7
	s_clause 0x1
	global_load_b64 v[4:5], v[4:5], off
	global_load_b64 v[2:3], v[2:3], off
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v8, vcc_lo
	v_add_co_u32 v11, vcc_lo, s6, v11
	v_add_co_ci_u32_e32 v12, vcc_lo, s7, v12, vcc_lo
	s_clause 0x1
	global_load_b64 v[9:10], v[9:10], off
	global_load_b64 v[11:12], v[11:12], off
	s_waitcnt vmcnt(2)
	v_mad_u64_u32 v[13:14], null, v0, s2, v[1:2]
	v_add_f64 v[2:3], v[4:5], v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v14, 31, v13
	v_lshlrev_b64 v[0:1], 3, v[13:14]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s7, v1, vcc_lo
	global_load_b64 v[0:1], v[0:1], off
	s_waitcnt vmcnt(2)
	v_fma_f64 v[2:3], v[9:10], -4.0, v[2:3]
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_f64 v[2:3], v[11:12], v[2:3]
	s_waitcnt vmcnt(0)
	v_add_f64 v[0:1], v[0:1], v[2:3]
	v_add_co_u32 v2, vcc_lo, s4, v7
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_fma_f64 v[0:1], v[0:1], s[0:1], v[9:10]
	global_store_b64 v[2:3], v[0:1], off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17simulate_version1diidddddddPdS_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 360
		.amdhsa_user_sgpr_count 14
		.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 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 15
		.amdhsa_next_free_sgpr 16
		.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	_Z17simulate_version1diidddddddPdS_S_i, .Lfunc_end0-_Z17simulate_version1diidddddddPdS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 456
; NumSgprs: 18
; NumVgprs: 15
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 15
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; 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: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           8
        .value_kind:     by_value
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         12
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           8
        .value_kind:     by_value
      - .offset:         24
        .size:           8
        .value_kind:     by_value
      - .offset:         32
        .size:           8
        .value_kind:     by_value
      - .offset:         40
        .size:           8
        .value_kind:     by_value
      - .offset:         48
        .size:           8
        .value_kind:     by_value
      - .offset:         56
        .size:           8
        .value_kind:     by_value
      - .offset:         64
        .size:           8
        .value_kind:     by_value
      - .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
      - .offset:         96
        .size:           4
        .value_kind:     by_value
      - .offset:         104
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         108
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         112
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         116
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         118
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         120
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         122
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         124
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         126
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         144
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         152
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         160
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         168
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 360
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z17simulate_version1diidddddddPdS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z17simulate_version1diidddddddPdS_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     15
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
