#include <stdio.h>
#include <math.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>

const char* getfield(char* line, int num){
  const char* tok;
  for (tok = strtok(line, ",");
  tok && *tok;
  tok = strtok(NULL, ",\n"))
  {
    if (!--num)
    return tok;
  }
  return NULL;
}

__global__ void calcPot(double *u, double *unext, int *currRefracIter, double *currTime, int *d_it, double *w, double *lastTime, double *d_refracTime,
int *d_maxRefracIter, int *d_N, int *d_R, double *d_uth, int *d_minMPVIter, double *d_dt, double *d_mi, double *d_sumCoeff){

  double pi=3.14159265359;
  int it = *d_it;
  double refracTime = *d_refracTime;
  int maxRefracIter = *d_maxRefracIter;
  int N = *d_N;
  int R = *d_R;
  double uth = *d_uth;
  int minMPVIter = *d_minMPVIter;
  double dt = *d_dt;
  double mi = *d_mi;
  double sumCoeff = *d_sumCoeff;
  int myId = blockDim.x * blockIdx.x + threadIdx.x;

  /*******Refractory Period*******/
  if (*(u+myId)==0 && *(currRefracIter+myId)<maxRefracIter){
    (*(currRefracIter+myId))++;
    return;
  }
  else{
    *(currRefracIter+myId)=0;
  }

  /*******Sum Calculation*******/
  double sumVar=0.0;
  int k,l;
  int iLeftCorner=N+myId/N-R;
  int jLeftCorner=N+myId%N-R;
  for  (k=iLeftCorner; k<iLeftCorner+2*R+1; k++){
    for (l=jLeftCorner; l<jLeftCorner+2*R+1; l++){
      sumVar+=*(u+myId)-*(u+(k%N)*N+l%N);
    }
  }

  *(unext+myId)=*(u+myId)+dt*(mi-*(u+myId)+sumCoeff*sumVar);
  *(currTime+myId)+=dt;

  if(*(unext+myId)>=uth){   //Threshold crossed
    *(unext+myId)=0.0;
    if (it>=minMPVIter){
      *(w+myId)=((*(w+myId))*(*(lastTime+myId))+2*pi)/((*(lastTime+myId))+(*(currTime+myId))+refracTime);
      *(lastTime+myId)+=(*(currTime+myId))+refracTime;
    }
    *(currTime+myId)=0.0;
  }

  return;
}

int main(int argc, char** argv){

  FILE *file1;
  char filename[100];

  /*******Parameter Declarations*******/
  int N=100;  //Grid dimension
  double dt=0.001; //0.001
  int totalTime=10000;  //Simulation time
  int it=0;
  int totalIter=totalTime/dt; //Total iterations
  int R=22; //Square radius
  double sigma=0.7; //Coupling strength
  double sumCoeff=sigma/((2*R+1)*(2*R+1)-1);  //Potential sum coefficient
  double mi=1.0; //Integrator floor
  double uth=0.98;

  double Ts=log(mi/(mi-uth));
  double refracTime=0.22*Ts;  //Refractory period time
  int maxRefracIter=(int)ceil(refracTime/dt); //Refractory period iterations
  int i,j;

  double u[N*N];
  double unext[N*N];

  int currRefracIter[N*N];  //Current iterations already in refractory period
  int maxMPVIter=30000;
  int minMPVIter=2000000; //bhma meta to opoio me endiaferei na arxisw na upologizw thn syxnothta twn neurwnwn.

  double currTime[N*N];
  double lastTime[N*N];
  double w[N*N];
  double t=0.0;

  for (i=0; i<N; i++){
    for (j=0; j<N; j++){
      (*(unext+i*N+j))=0.0;
      (*(currTime+i*N+j))=0.0;
      (*(lastTime+i*N+j))=0.0;
      (*(currRefracIter+i*N+j))=0.0;
    }
  }
  file1=fopen(argv[1],"r"); //argv[1]
  char line[2048];
  i=0;
  while(fgets(line, 2048, file1)){
    for(j=1;j<=N;j++){
      char* tmp = strdup(line);
      (*(u+N*i+j-1))=atof(getfield(tmp,j));
      free(tmp);
    }
    i++;
  }
  fclose(file1);

  double *d_u, *d_unext, *d_currTime, *d_w, *d_lastTime, *d_refracTime, *d_uth, *d_dt, *d_mi, *d_sumCoeff;
  int *d_currRefracIter, *d_it, *d_maxRefracIter, *d_N, *d_R, *d_minMPVIter;

  cudaMalloc(&d_u, N*N*sizeof(double));
  cudaMalloc(&d_unext, N*N*sizeof(double));
  cudaMalloc(&d_currRefracIter, N*N*sizeof(int));
  cudaMalloc(&d_currTime, N*N*sizeof(double));
  cudaMalloc(&d_it, sizeof(int));
  cudaMalloc(&d_w, N*N*sizeof(double));
  cudaMalloc(&d_lastTime, N*N*sizeof(double));

  cudaMalloc(&d_refracTime, sizeof(double));
  cudaMalloc(&d_maxRefracIter, sizeof(int));
  cudaMalloc(&d_N, sizeof(int));
  cudaMalloc(&d_R, sizeof(int));
  cudaMalloc(&d_uth, sizeof(double));
  cudaMalloc(&d_minMPVIter, sizeof(int));
  cudaMalloc(&d_dt, sizeof(double));
  cudaMalloc(&d_mi, sizeof(double));
  cudaMalloc(&d_sumCoeff, sizeof(double));

  cudaMemcpy(d_refracTime, &refracTime, sizeof(double), cudaMemcpyHostToDevice);
  cudaMemcpy(d_maxRefracIter, &maxRefracIter, sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_N, &N, sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_R, &R, sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_uth, &uth, sizeof(double), cudaMemcpyHostToDevice);
  cudaMemcpy(d_minMPVIter, &minMPVIter, sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_dt, &dt, sizeof(double), cudaMemcpyHostToDevice);
  cudaMemcpy(d_mi, &mi, sizeof(double), cudaMemcpyHostToDevice);
  cudaMemcpy(d_sumCoeff, &sumCoeff, sizeof(double), cudaMemcpyHostToDevice);

  time_t benchBegin = time(NULL);
  /*******Simulation*******/
  while (it<totalIter){

    if (it%10000==0) printf("Iteration %d of %d\n", it, totalIter);

    cudaMemcpy(d_u, u, N*N*sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_unext, unext, N*N*sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_currRefracIter, currRefracIter, N*N*sizeof(int),cudaMemcpyHostToDevice);
    cudaMemcpy(d_currTime, currTime, N*N*sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_it, &it, sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_w, w, N*N*sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(d_lastTime, lastTime, N*N*sizeof(double), cudaMemcpyHostToDevice);

    //printf("STARTING\n");
    calcPot<<<100,100>>>(d_u, d_unext, d_currRefracIter, d_currTime, d_it, d_w, d_lastTime, d_refracTime, d_maxRefracIter, d_N, d_R, d_uth, d_minMPVIter, d_dt, d_mi, d_sumCoeff);
    cudaThreadSynchronize();

    cudaMemcpy(unext, d_unext, N*N*sizeof(double), cudaMemcpyDeviceToHost);
    cudaMemcpy(currRefracIter, d_currRefracIter, N*N*sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(currTime, d_currTime, N*N*sizeof(double), cudaMemcpyDeviceToHost);
    cudaMemcpy(w, d_w, N*N*sizeof(double), cudaMemcpyDeviceToHost);
    cudaMemcpy(lastTime, d_lastTime, N*N*sizeof(double), cudaMemcpyDeviceToHost);

    // for(i=0; i<N; i++){
    //   for(j=0; j<N; j++){
    //     printf("%lf ", *(unext+N*i+j));
    //   }
    //   printf("\n");
    // }
    // printf("FINISHED\n");

    if(it%10000==0){
      sprintf(filename, "ResultsCUDA%s/Results_POT_LIF_2D_Classic_sigma_%lf_R_%d_time_%lf_.dat",argv[2],sigma,R,t);
      file1=fopen(filename,"w");
      for(i=0;i<N;i++){
        for(j=0;j<N;j++){
          fprintf(file1, "%lf,",*(unext+N*i+j));
        }
        fprintf(file1,"\n");
      }
      fclose(file1);
    }
    if (it>minMPVIter){
      if ((it-minMPVIter)%maxMPVIter==0){
        sprintf(filename, "ResultsCUDA%s/Results_MPV_LIF_2D_Classic_sigma_%lf_R_%d_time_%lf_.dat",argv[2],sigma,R,t);
        file1=fopen(filename,"w");
        for(i=0;i<N;i++){
          for(j=0;j<N;j++){
            fprintf(file1,"%lf,",*(w+N*i+j));
          }
          fprintf(file1,"\n");
        }
        fclose(file1);
      }
    }
    if (it == 2000000){
      time_t benchEnd = time(NULL);
      sprintf(filename, "ResultsCUDA%s/execTime.dat",argv[2]);
      file1=fopen(filename,"w");
      fprintf(file1,"Execution time for 2000 time units: %ld seconds\n",benchEnd-benchBegin);
      fclose(file1);
    }
    for (i=0; i<N; i++){
      for (j=0; j<N; j++){
        (*(u+N*i+j))=*(unext+N*i+j);
      }
    }
    t+=dt;
    it++;
  } //edw kleinei h while.

  cudaFree(d_u);
  cudaFree(d_unext);
  cudaFree(d_currRefracIter);
  cudaFree(d_currTime);
  cudaFree(d_it);
  cudaFree(d_w);
  cudaFree(d_lastTime);
  cudaFree(d_refracTime);
  cudaFree(d_maxRefracIter);
  cudaFree(d_N);
  cudaFree(d_R);
  cudaFree(d_uth);
  return(0);
}

	.file	"tmpxft_0027834f_00000000-6_LIF_2D_Classic_CUDA.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0, @function
_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0:
.LFB2174:
	.cfi_startproc
	jmp	cudaMalloc@PLT
	.cfi_endproc
.LFE2174:
	.size	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0, .-_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	","
.LC1:
	.string	",\n"
	.text
	.globl	_Z8getfieldPci
	.type	_Z8getfieldPci, @function
_Z8getfieldPci:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	leaq	.LC1(%rip), %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%esi, %ebx
	leaq	.LC0(%rip), %rsi
	pushq	%rcx
	.cfi_def_cfa_offset 32
	call	strtok@PLT
.L4:
	testq	%rax, %rax
	je	.L3
	cmpb	$0, (%rax)
	je	.L15
	decl	%ebx
	je	.L3
	movq	%rbp, %rsi
	xorl	%edi, %edi
	call	strtok@PLT
	jmp	.L4
.L15:
	xorl	%eax, %eax
.L3:
	popq	%rdx
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z8getfieldPci, .-_Z8getfieldPci
	.globl	_Z61__device_stub__Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_PdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	.type	_Z61__device_stub__Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_PdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_, @function
_Z61__device_stub__Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_PdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$344, %rsp
	.cfi_def_cfa_offset 352
	movq	352(%rsp), %rax
	movq	%rdi, 120(%rsp)
	leaq	152(%rsp), %rdi
	movq	%rsi, 112(%rsp)
	leaq	164(%rsp), %rsi
	movq	%rax, 72(%rsp)
	movq	360(%rsp), %rax
	movq	%rdx, 104(%rsp)
	leaq	136(%rsp), %rdx
	movq	%rax, 64(%rsp)
	movq	368(%rsp), %rax
	movq	%rcx, 96(%rsp)
	leaq	144(%rsp), %rcx
	movq	%rax, 56(%rsp)
	movq	376(%rsp), %rax
	movq	%r8, 88(%rsp)
	movq	%rax, 48(%rsp)
	movq	384(%rsp), %rax
	movq	%r9, 80(%rsp)
	movq	%rax, 40(%rsp)
	movq	392(%rsp), %rax
	movq	%rax, 32(%rsp)
	movq	400(%rsp), %rax
	movq	%rax, 24(%rsp)
	movq	408(%rsp), %rax
	movq	%rax, 16(%rsp)
	movq	416(%rsp), %rax
	movq	%rax, 8(%rsp)
	movq	424(%rsp), %rax
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 328(%rsp)
	xorl	%eax, %eax
	leaq	120(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 312(%rsp)
	movq	%rsp, %rax
	movq	%rax, 320(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 152(%rsp)
	movl	$1, 160(%rsp)
	movq	%rax, 164(%rsp)
	movl	$1, 172(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L16
	pushq	144(%rsp)
	.cfi_def_cfa_offset 360
	leaq	_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_(%rip), %rdi
	pushq	144(%rsp)
	.cfi_def_cfa_offset 368
	movq	180(%rsp), %rcx
	movl	188(%rsp), %r8d
	movq	168(%rsp), %rsi
	movl	176(%rsp), %edx
	leaq	216(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 360
	popq	%rdx
	.cfi_def_cfa_offset 352
.L16:
	movq	328(%rsp), %rax
	subq	%fs:40, %rax
	je	.L18
	call	__stack_chk_fail@PLT
.L18:
	addq	$344, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z61__device_stub__Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_PdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_, .-_Z61__device_stub__Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_PdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	.globl	_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	.type	_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_, @function
_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z61__device_stub__Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_PdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	.cfi_endproc
.LFE2054:
	.size	_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_, .-_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	.section	.rodata.str1.1
.LC8:
	.string	"r"
.LC9:
	.string	"Iteration %d of %d\n"
.LC11:
	.string	"ResultsCUDA%s/Results_POT_LIF_2D_Classic_sigma_%lf_R_%d_time_%lf_.dat"
.LC12:
	.string	"w"
.LC13:
	.string	"%lf,"
.LC14:
	.string	"\n"
.LC15:
	.string	"ResultsCUDA%s/Results_MPV_LIF_2D_Classic_sigma_%lf_R_%d_time_%lf_.dat"
.LC16:
	.string	"ResultsCUDA%s/execTime.dat"
.LC17:
	.string	"Execution time for 2000 time units: %ld seconds\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	.cfi_offset 12, -48
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	subq	$2472, %rsp
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	movq	.LC3(%rip), %rax
	movl	$100, -2420(%rbp)
	movl	$22, -2412(%rbp)
	movq	%rax, -2400(%rbp)
	xorl	%eax, %eax
	movl	%eax, -2416(%rbp)
	movq	.LC4(%rip), %rax
	movl	$861, -2408(%rbp)
	movq	%rax, -2392(%rbp)
	movq	.LC5(%rip), %rax
	movq	%rax, -2384(%rbp)
	movq	.LC6(%rip), %rax
	movq	%rax, -2376(%rbp)
	movq	.LC7(%rip), %rax
	movq	%rax, -2368(%rbp)
	leaq	-77824(%rsp), %rax
.L22:
	cmpq	%rax, %rsp
	je	.L23
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L22
.L23:
	subq	$2176, %rsp
	orq	$0, 2168(%rsp)
	movq	%rsp, -2464(%rbp)
	leaq	-77824(%rsp), %rax
.L25:
	cmpq	%rax, %rsp
	je	.L26
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L25
.L26:
	subq	$2176, %rsp
	orq	$0, 2168(%rsp)
	movq	%rsp, -2440(%rbp)
	leaq	-36864(%rsp), %rax
.L28:
	cmpq	%rax, %rsp
	je	.L29
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L28
.L29:
	subq	$3136, %rsp
	orq	$0, 3128(%rsp)
	movq	%rsp, -2472(%rbp)
	leaq	-77824(%rsp), %rax
	movl	$2000000, -2404(%rbp)
.L31:
	cmpq	%rax, %rsp
	je	.L32
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L31
.L32:
	subq	$2176, %rsp
	orq	$0, 2168(%rsp)
	movq	%rsp, -2480(%rbp)
	leaq	-77824(%rsp), %rax
.L34:
	cmpq	%rax, %rsp
	je	.L35
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L34
.L35:
	subq	$2176, %rsp
	orq	$0, 2168(%rsp)
	movq	%rsp, -2488(%rbp)
	leaq	-77824(%rsp), %rax
.L37:
	cmpq	%rax, %rsp
	je	.L38
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L37
.L38:
	subq	$2176, %rsp
	orq	$0, 2168(%rsp)
	xorl	%eax, %eax
	movl	$20000, %ecx
	leaq	.LC8(%rip), %rsi
	xorl	%r14d, %r14d
	movq	-2440(%rbp), %rdi
	rep stosl
	movq	-2480(%rbp), %rdi
	movl	$20000, %ecx
	rep stosl
	movq	-2488(%rbp), %rdi
	movl	$20000, %ecx
	movq	%rsp, -2496(%rbp)
	leaq	-2104(%rbp), %r15
	rep stosl
	movq	-2472(%rbp), %rdi
	movl	$10000, %ecx
	rep stosl
	movq	8(%rbx), %rdi
	call	fopen@PLT
	movq	%rax, -2448(%rbp)
.L40:
	movq	-2448(%rbp), %rcx
	movl	$2048, %edx
	movl	$2048, %esi
	movq	%r15, %rdi
	call	__fgets_chk@PLT
	testq	%rax, %rax
	je	.L65
	movl	$1, %r12d
.L42:
	movl	%r12d, %esi
	cmpl	%r12d, -2420(%rbp)
	jl	.L66
	movq	%r15, %rdi
	movl	%esi, -2500(%rbp)
	call	strdup@PLT
	movl	-2500(%rbp), %esi
	movq	%rax, %rdi
	movq	%rax, -2456(%rbp)
	call	_Z8getfieldPci
	movl	-2420(%rbp), %ecx
	movq	%rax, %rdi
	movq	-2464(%rbp), %rax
	imull	%r14d, %ecx
	movslq	%ecx, %rcx
	addq	%r12, %rcx
	incq	%r12
	leaq	-8(%rax,%rcx,8), %r13
	call	atof@PLT
	movq	-2456(%rbp), %rdi
	movsd	%xmm0, 0(%r13)
	call	free@PLT
	jmp	.L42
.L66:
	incl	%r14d
	jmp	.L40
.L65:
	movq	-2448(%rbp), %rdi
	call	fclose@PLT
	movl	-2420(%rbp), %esi
	leaq	-2360(%rbp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movl	-2420(%rbp), %esi
	leaq	-2352(%rbp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movl	-2420(%rbp), %esi
	leaq	-2280(%rbp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movl	-2420(%rbp), %esi
	leaq	-2344(%rbp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	-2272(%rbp), %rdi
	movl	$4, %esi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movl	-2420(%rbp), %esi
	leaq	-2336(%rbp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movl	-2420(%rbp), %esi
	leaq	-2328(%rbp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	-2320(%rbp), %rdi
	movl	$8, %esi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	-2264(%rbp), %rdi
	movl	$4, %esi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	-2256(%rbp), %rdi
	movl	$4, %esi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	-2248(%rbp), %rdi
	movl	$4, %esi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	-2312(%rbp), %rdi
	movl	$8, %esi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	-2240(%rbp), %rdi
	movl	$4, %esi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	-2304(%rbp), %rdi
	movl	$8, %esi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	-2296(%rbp), %rdi
	movl	$8, %esi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	leaq	-2288(%rbp), %rdi
	movl	$8, %esi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movl	$1, %ecx
	movl	$8, %edx
	movq	-2320(%rbp), %rdi
	leaq	-2368(%rbp), %rsi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$4, %edx
	movq	-2264(%rbp), %rdi
	leaq	-2408(%rbp), %rsi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$4, %edx
	movq	-2256(%rbp), %rdi
	leaq	-2420(%rbp), %rsi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$4, %edx
	movq	-2248(%rbp), %rdi
	leaq	-2412(%rbp), %rsi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$8, %edx
	movq	-2312(%rbp), %rdi
	leaq	-2376(%rbp), %rsi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$4, %edx
	movq	-2240(%rbp), %rdi
	leaq	-2404(%rbp), %rsi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$8, %edx
	movq	-2304(%rbp), %rdi
	leaq	-2400(%rbp), %rsi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$8, %edx
	movq	-2296(%rbp), %rdi
	leaq	-2384(%rbp), %rsi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$8, %edx
	movq	-2288(%rbp), %rdi
	leaq	-2392(%rbp), %rsi
	call	cudaMemcpy@PLT
	xorl	%edi, %edi
	call	time@PLT
	movq	$0x000000000, -2448(%rbp)
	movq	%rax, %r15
.L43:
	movl	-2416(%rbp), %esi
	cmpl	$9999999, %esi
	jg	.L67
	movl	%esi, %eax
	movl	$10000, %ecx
	cltd
	idivl	%ecx
	testl	%edx, %edx
	jne	.L44
	movl	%esi, %edx
	movl	$10000000, %ecx
	leaq	.LC9(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
.L44:
	movl	-2420(%rbp), %edx
	movq	-2464(%rbp), %rsi
	movl	$1, %ecx
	movq	-2360(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	-2420(%rbp), %edx
	movl	$1, %ecx
	movq	-2440(%rbp), %rsi
	movq	-2352(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	-2420(%rbp), %edx
	movl	$1, %ecx
	movq	-2472(%rbp), %rsi
	movq	-2280(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	-2420(%rbp), %edx
	movl	$1, %ecx
	movq	-2480(%rbp), %rsi
	movq	-2344(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$4, %edx
	movq	-2272(%rbp), %rdi
	leaq	-2416(%rbp), %rsi
	call	cudaMemcpy@PLT
	movl	-2420(%rbp), %edx
	movl	$1, %ecx
	movq	-2496(%rbp), %rsi
	movq	-2336(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	-2420(%rbp), %edx
	movl	$1, %ecx
	movq	-2488(%rbp), %rsi
	movq	-2328(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	$1073741849, %edi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$2, %rdi
	movl	$1, %ecx
	movl	$1, %esi
	movl	$1, -2208(%rbp)
	movq	%rdi, %rdx
	movq	%rdi, -2216(%rbp)
	movq	%rdi, -2228(%rbp)
	movl	$1, -2220(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L45
	pushq	-2288(%rbp)
	pushq	-2296(%rbp)
	pushq	-2304(%rbp)
	pushq	-2240(%rbp)
	pushq	-2312(%rbp)
	pushq	-2248(%rbp)
	pushq	-2256(%rbp)
	pushq	-2264(%rbp)
	movq	-2272(%rbp), %r8
	pushq	-2320(%rbp)
	pushq	-2328(%rbp)
	movq	-2336(%rbp), %r9
	movq	-2344(%rbp), %rcx
	movq	-2280(%rbp), %rdx
	movq	-2352(%rbp), %rsi
	movq	-2360(%rbp), %rdi
	call	_Z61__device_stub__Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_PdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	addq	$80, %rsp
.L45:
	call	cudaThreadSynchronize@PLT
	movl	-2420(%rbp), %edx
	movl	$2, %ecx
	movq	-2352(%rbp), %rsi
	movq	-2440(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	-2420(%rbp), %edx
	movl	$2, %ecx
	movq	-2280(%rbp), %rsi
	movq	-2472(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	-2420(%rbp), %edx
	movl	$2, %ecx
	movq	-2344(%rbp), %rsi
	movq	-2480(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	-2420(%rbp), %edx
	movl	$2, %ecx
	movq	-2336(%rbp), %rsi
	movq	-2496(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	-2420(%rbp), %edx
	movl	$2, %ecx
	movq	-2328(%rbp), %rsi
	movq	-2488(%rbp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	-2416(%rbp), %eax
	movl	$10000, %ecx
	cltd
	idivl	%ecx
	movl	%edx, %r13d
	testl	%edx, %edx
	jne	.L46
	movsd	-2448(%rbp), %xmm1
	movq	16(%rbx), %r8
	movl	$100, %edx
	movb	$2, %al
	movl	-2412(%rbp), %r9d
	movsd	.LC10(%rip), %xmm0
	leaq	-2204(%rbp), %r12
	leaq	.LC11(%rip), %rcx
	movl	$2, %esi
	movq	%r12, %rdi
	leaq	.LC13(%rip), %r14
	call	__sprintf_chk@PLT
	leaq	.LC12(%rip), %rsi
	movq	%r12, %rdi
	call	fopen@PLT
	movq	%rax, -2456(%rbp)
.L47:
	cmpl	%r13d, -2420(%rbp)
	jle	.L68
	xorl	%r12d, %r12d
.L49:
	movl	-2420(%rbp), %eax
	cmpl	%r12d, %eax
	jle	.L69
	imull	%r13d, %eax
	movq	-2440(%rbp), %rcx
	movq	%r14, %rdx
	movl	$2, %esi
	movq	-2456(%rbp), %rdi
	cltq
	addq	%r12, %rax
	incq	%r12
	movsd	(%rcx,%rax,8), %xmm0
	movb	$1, %al
	call	__fprintf_chk@PLT
	jmp	.L49
.L69:
	movq	-2456(%rbp), %rdi
	leaq	.LC14(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	incl	%r13d
	call	__fprintf_chk@PLT
	jmp	.L47
.L68:
	movq	-2456(%rbp), %rdi
	call	fclose@PLT
.L46:
	movl	-2416(%rbp), %eax
	movl	-2404(%rbp), %edx
	cmpl	%edx, %eax
	jle	.L50
	subl	%edx, %eax
	movl	$30000, %ecx
	cltd
	idivl	%ecx
	movl	%edx, %r13d
	testl	%edx, %edx
	jne	.L50
	movsd	-2448(%rbp), %xmm1
	movq	16(%rbx), %r8
	movl	$100, %edx
	movb	$2, %al
	movl	-2412(%rbp), %r9d
	movsd	.LC10(%rip), %xmm0
	leaq	-2204(%rbp), %r12
	leaq	.LC15(%rip), %rcx
	movl	$2, %esi
	movq	%r12, %rdi
	leaq	.LC13(%rip), %r14
	call	__sprintf_chk@PLT
	leaq	.LC12(%rip), %rsi
	movq	%r12, %rdi
	call	fopen@PLT
	movq	%rax, -2456(%rbp)
.L51:
	cmpl	%r13d, -2420(%rbp)
	jle	.L70
	xorl	%r12d, %r12d
.L53:
	movl	-2420(%rbp), %eax
	cmpl	%r12d, %eax
	jle	.L71
	imull	%r13d, %eax
	movq	-2496(%rbp), %rcx
	movq	%r14, %rdx
	movl	$2, %esi
	movq	-2456(%rbp), %rdi
	cltq
	addq	%r12, %rax
	incq	%r12
	movsd	(%rcx,%rax,8), %xmm0
	movb	$1, %al
	call	__fprintf_chk@PLT
	jmp	.L53
.L71:
	movq	-2456(%rbp), %rdi
	leaq	.LC14(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	incl	%r13d
	call	__fprintf_chk@PLT
	jmp	.L51
.L70:
	movq	-2456(%rbp), %rdi
	call	fclose@PLT
.L50:
	cmpl	$2000000, -2416(%rbp)
	jne	.L54
	xorl	%edi, %edi
	leaq	-2204(%rbp), %r13
	call	time@PLT
	movq	16(%rbx), %r8
	movl	$100, %edx
	movq	%r13, %rdi
	leaq	.LC16(%rip), %rcx
	movq	%rax, %r12
	movl	$2, %esi
	xorl	%eax, %eax
	call	__sprintf_chk@PLT
	movq	%r13, %rdi
	leaq	.LC12(%rip), %rsi
	call	fopen@PLT
	movq	%r12, %rcx
	movl	$2, %esi
	leaq	.LC17(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %r13
	subq	%r15, %rcx
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movq	%r13, %rdi
	call	fclose@PLT
.L54:
	movl	-2420(%rbp), %esi
	xorl	%ecx, %ecx
	xorl	%edi, %edi
.L55:
	cmpl	%edi, %esi
	jle	.L57
	movslq	%ecx, %rax
	xorl	%edx, %edx
	salq	$3, %rax
.L56:
	movq	-2440(%rbp), %r10
	movq	-2464(%rbp), %r11
	incl	%edx
	movsd	(%r10,%rax), %xmm0
	movsd	%xmm0, (%r11,%rax)
	addq	$8, %rax
	cmpl	%edx, %esi
	jne	.L56
	incl	%edi
	addl	%esi, %ecx
	jmp	.L55
.L57:
	movsd	-2448(%rbp), %xmm2
	addsd	-2400(%rbp), %xmm2
	incl	-2416(%rbp)
	movsd	%xmm2, -2448(%rbp)
	jmp	.L43
.L67:
	movq	-2360(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2352(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2280(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2344(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2272(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2336(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2328(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2320(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2264(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2256(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2248(%rbp), %rdi
	call	cudaFree@PLT
	movq	-2312(%rbp), %rdi
	call	cudaFree@PLT
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L59
	call	__stack_chk_fail@PLT
.L59:
	leaq	-40(%rbp), %rsp
	xorl	%eax, %eax
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE2028:
	.size	main, .-main
	.section	.rodata.str1.1
.LC18:
	.string	"_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2056:
	.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	.LC18(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_(%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
.LFE2056:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC3:
	.long	-755914244
	.long	1062232653
	.align 8
.LC4:
	.long	-1758729692
	.long	1060547173
	.align 8
.LC5:
	.long	0
	.long	1072693248
	.align 8
.LC6:
	.long	-171798692
	.long	1072651304
	.align 8
.LC7:
	.long	-2098634743
	.long	1072400999
	.align 8
.LC10:
	.long	1717986918
	.long	1072064102
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <math.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>

const char* getfield(char* line, int num){
  const char* tok;
  for (tok = strtok(line, ",");
  tok && *tok;
  tok = strtok(NULL, ",\n"))
  {
    if (!--num)
    return tok;
  }
  return NULL;
}

__global__ void calcPot(double *u, double *unext, int *currRefracIter, double *currTime, int *d_it, double *w, double *lastTime, double *d_refracTime,
int *d_maxRefracIter, int *d_N, int *d_R, double *d_uth, int *d_minMPVIter, double *d_dt, double *d_mi, double *d_sumCoeff){

  double pi=3.14159265359;
  int it = *d_it;
  double refracTime = *d_refracTime;
  int maxRefracIter = *d_maxRefracIter;
  int N = *d_N;
  int R = *d_R;
  double uth = *d_uth;
  int minMPVIter = *d_minMPVIter;
  double dt = *d_dt;
  double mi = *d_mi;
  double sumCoeff = *d_sumCoeff;
  int myId = blockDim.x * blockIdx.x + threadIdx.x;

  /*******Refractory Period*******/
  if (*(u+myId)==0 && *(currRefracIter+myId)<maxRefracIter){
    (*(currRefracIter+myId))++;
    return;
  }
  else{
    *(currRefracIter+myId)=0;
  }

  /*******Sum Calculation*******/
  double sumVar=0.0;
  int k,l;
  int iLeftCorner=N+myId/N-R;
  int jLeftCorner=N+myId%N-R;
  for  (k=iLeftCorner; k<iLeftCorner+2*R+1; k++){
    for (l=jLeftCorner; l<jLeftCorner+2*R+1; l++){
      sumVar+=*(u+myId)-*(u+(k%N)*N+l%N);
    }
  }

  *(unext+myId)=*(u+myId)+dt*(mi-*(u+myId)+sumCoeff*sumVar);
  *(currTime+myId)+=dt;

  if(*(unext+myId)>=uth){   //Threshold crossed
    *(unext+myId)=0.0;
    if (it>=minMPVIter){
      *(w+myId)=((*(w+myId))*(*(lastTime+myId))+2*pi)/((*(lastTime+myId))+(*(currTime+myId))+refracTime);
      *(lastTime+myId)+=(*(currTime+myId))+refracTime;
    }
    *(currTime+myId)=0.0;
  }

  return;
}

int main(int argc, char** argv){

  FILE *file1;
  char filename[100];

  /*******Parameter Declarations*******/
  int N=100;  //Grid dimension
  double dt=0.001; //0.001
  int totalTime=10000;  //Simulation time
  int it=0;
  int totalIter=totalTime/dt; //Total iterations
  int R=22; //Square radius
  double sigma=0.7; //Coupling strength
  double sumCoeff=sigma/((2*R+1)*(2*R+1)-1);  //Potential sum coefficient
  double mi=1.0; //Integrator floor
  double uth=0.98;

  double Ts=log(mi/(mi-uth));
  double refracTime=0.22*Ts;  //Refractory period time
  int maxRefracIter=(int)ceil(refracTime/dt); //Refractory period iterations
  int i,j;

  double u[N*N];
  double unext[N*N];

  int currRefracIter[N*N];  //Current iterations already in refractory period
  int maxMPVIter=30000;
  int minMPVIter=2000000; //bhma meta to opoio me endiaferei na arxisw na upologizw thn syxnothta twn neurwnwn.

  double currTime[N*N];
  double lastTime[N*N];
  double w[N*N];
  double t=0.0;

  for (i=0; i<N; i++){
    for (j=0; j<N; j++){
      (*(unext+i*N+j))=0.0;
      (*(currTime+i*N+j))=0.0;
      (*(lastTime+i*N+j))=0.0;
      (*(currRefracIter+i*N+j))=0.0;
    }
  }
  file1=fopen(argv[1],"r"); //argv[1]
  char line[2048];
  i=0;
  while(fgets(line, 2048, file1)){
    for(j=1;j<=N;j++){
      char* tmp = strdup(line);
      (*(u+N*i+j-1))=atof(getfield(tmp,j));
      free(tmp);
    }
    i++;
  }
  fclose(file1);

  double *d_u, *d_unext, *d_currTime, *d_w, *d_lastTime, *d_refracTime, *d_uth, *d_dt, *d_mi, *d_sumCoeff;
  int *d_currRefracIter, *d_it, *d_maxRefracIter, *d_N, *d_R, *d_minMPVIter;

  hipMalloc(&d_u, N*N*sizeof(double));
  hipMalloc(&d_unext, N*N*sizeof(double));
  hipMalloc(&d_currRefracIter, N*N*sizeof(int));
  hipMalloc(&d_currTime, N*N*sizeof(double));
  hipMalloc(&d_it, sizeof(int));
  hipMalloc(&d_w, N*N*sizeof(double));
  hipMalloc(&d_lastTime, N*N*sizeof(double));

  hipMalloc(&d_refracTime, sizeof(double));
  hipMalloc(&d_maxRefracIter, sizeof(int));
  hipMalloc(&d_N, sizeof(int));
  hipMalloc(&d_R, sizeof(int));
  hipMalloc(&d_uth, sizeof(double));
  hipMalloc(&d_minMPVIter, sizeof(int));
  hipMalloc(&d_dt, sizeof(double));
  hipMalloc(&d_mi, sizeof(double));
  hipMalloc(&d_sumCoeff, sizeof(double));

  hipMemcpy(d_refracTime, &refracTime, sizeof(double), hipMemcpyHostToDevice);
  hipMemcpy(d_maxRefracIter, &maxRefracIter, sizeof(int), hipMemcpyHostToDevice);
  hipMemcpy(d_N, &N, sizeof(int), hipMemcpyHostToDevice);
  hipMemcpy(d_R, &R, sizeof(int), hipMemcpyHostToDevice);
  hipMemcpy(d_uth, &uth, sizeof(double), hipMemcpyHostToDevice);
  hipMemcpy(d_minMPVIter, &minMPVIter, sizeof(int), hipMemcpyHostToDevice);
  hipMemcpy(d_dt, &dt, sizeof(double), hipMemcpyHostToDevice);
  hipMemcpy(d_mi, &mi, sizeof(double), hipMemcpyHostToDevice);
  hipMemcpy(d_sumCoeff, &sumCoeff, sizeof(double), hipMemcpyHostToDevice);

  time_t benchBegin = time(NULL);
  /*******Simulation*******/
  while (it<totalIter){

    if (it%10000==0) printf("Iteration %d of %d\n", it, totalIter);

    hipMemcpy(d_u, u, N*N*sizeof(double), hipMemcpyHostToDevice);
    hipMemcpy(d_unext, unext, N*N*sizeof(double), hipMemcpyHostToDevice);
    hipMemcpy(d_currRefracIter, currRefracIter, N*N*sizeof(int),hipMemcpyHostToDevice);
    hipMemcpy(d_currTime, currTime, N*N*sizeof(double), hipMemcpyHostToDevice);
    hipMemcpy(d_it, &it, sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_w, w, N*N*sizeof(double), hipMemcpyHostToDevice);
    hipMemcpy(d_lastTime, lastTime, N*N*sizeof(double), hipMemcpyHostToDevice);

    //printf("STARTING\n");
    calcPot<<<100,100>>>(d_u, d_unext, d_currRefracIter, d_currTime, d_it, d_w, d_lastTime, d_refracTime, d_maxRefracIter, d_N, d_R, d_uth, d_minMPVIter, d_dt, d_mi, d_sumCoeff);
    hipDeviceSynchronize();

    hipMemcpy(unext, d_unext, N*N*sizeof(double), hipMemcpyDeviceToHost);
    hipMemcpy(currRefracIter, d_currRefracIter, N*N*sizeof(int), hipMemcpyDeviceToHost);
    hipMemcpy(currTime, d_currTime, N*N*sizeof(double), hipMemcpyDeviceToHost);
    hipMemcpy(w, d_w, N*N*sizeof(double), hipMemcpyDeviceToHost);
    hipMemcpy(lastTime, d_lastTime, N*N*sizeof(double), hipMemcpyDeviceToHost);

    // for(i=0; i<N; i++){
    //   for(j=0; j<N; j++){
    //     printf("%lf ", *(unext+N*i+j));
    //   }
    //   printf("\n");
    // }
    // printf("FINISHED\n");

    if(it%10000==0){
      sprintf(filename, "ResultsCUDA%s/Results_POT_LIF_2D_Classic_sigma_%lf_R_%d_time_%lf_.dat",argv[2],sigma,R,t);
      file1=fopen(filename,"w");
      for(i=0;i<N;i++){
        for(j=0;j<N;j++){
          fprintf(file1, "%lf,",*(unext+N*i+j));
        }
        fprintf(file1,"\n");
      }
      fclose(file1);
    }
    if (it>minMPVIter){
      if ((it-minMPVIter)%maxMPVIter==0){
        sprintf(filename, "ResultsCUDA%s/Results_MPV_LIF_2D_Classic_sigma_%lf_R_%d_time_%lf_.dat",argv[2],sigma,R,t);
        file1=fopen(filename,"w");
        for(i=0;i<N;i++){
          for(j=0;j<N;j++){
            fprintf(file1,"%lf,",*(w+N*i+j));
          }
          fprintf(file1,"\n");
        }
        fclose(file1);
      }
    }
    if (it == 2000000){
      time_t benchEnd = time(NULL);
      sprintf(filename, "ResultsCUDA%s/execTime.dat",argv[2]);
      file1=fopen(filename,"w");
      fprintf(file1,"Execution time for 2000 time units: %ld seconds\n",benchEnd-benchBegin);
      fclose(file1);
    }
    for (i=0; i<N; i++){
      for (j=0; j<N; j++){
        (*(u+N*i+j))=*(unext+N*i+j);
      }
    }
    t+=dt;
    it++;
  } //edw kleinei h while.

  hipFree(d_u);
  hipFree(d_unext);
  hipFree(d_currRefracIter);
  hipFree(d_currTime);
  hipFree(d_it);
  hipFree(d_w);
  hipFree(d_lastTime);
  hipFree(d_refracTime);
  hipFree(d_maxRefracIter);
  hipFree(d_N);
  hipFree(d_R);
  hipFree(d_uth);
  return(0);
}

	.text
	.file	"LIF_2D_Classic_CUDA.hip"
	.globl	_Z8getfieldPci                  # -- Begin function _Z8getfieldPci
	.type	_Z8getfieldPci,@function
_Z8getfieldPci:                         # @_Z8getfieldPci
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movl	%esi, %ebx
	movl	$.L.str, %esi
	callq	strtok
	testq	%rax, %rax
	je	.LBB0_1
.LBB0_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	cmpb	$0, (%rax)
	je	.LBB0_1
# %bb.3:                                #   in Loop: Header=BB0_2 Depth=1
	movq	%rax, %rcx
	decl	%ebx
	je	.LBB0_5
# %bb.4:                                #   in Loop: Header=BB0_2 Depth=1
	movl	$.L.str.1, %esi
	xorl	%edi, %edi
	callq	strtok
	movl	$0, %ecx
	testq	%rax, %rax
	jne	.LBB0_2
	jmp	.LBB0_5
.LBB0_1:
	xorl	%ecx, %ecx
.LBB0_5:                                # %.critedge
	movq	%rcx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z8getfieldPci, .Lfunc_end0-_Z8getfieldPci
	.cfi_endproc
                                        # -- End function
	.globl	_Z22__device_stub__calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_ # -- Begin function _Z22__device_stub__calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	.type	_Z22__device_stub__calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_,@function
_Z22__device_stub__calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_: # @_Z22__device_stub__calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$224, %rsp
	.cfi_def_cfa_offset 272
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	16(%rsp), %r8
	movq	%r9, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	272(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	280(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	288(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	296(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	304(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	312(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	320(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	328(%rsp), %rax
	movq	%rax, 104(%rbx)
	leaq	336(%rsp), %rax
	movq	%rax, 112(%rbx)
	leaq	344(%rsp), %rax
	movq	%rax, 120(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$240, %rsp
	.cfi_adjust_cfa_offset -240
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z22__device_stub__calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_, .Lfunc_end1-_Z22__device_stub__calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI2_0:
	.quad	0x3fe6666666666666              # double 0.69999999999999996
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:                                # %.preheader94.lr.ph
	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	$442392, %rsp                   # imm = 0x6C018
	.cfi_def_cfa_offset 442448
	.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, 24(%rsp)                  # 8-byte Spill
	movq	%rsp, 184(%rsp)                 # 8-byte Spill
	movl	$100, (%rsp)
	movabsq	$4562254508917369340, %rax      # imm = 0x3F50624DD2F1A9FC
	movq	%rax, 176(%rsp)
	movl	$0, 4(%rsp)
	movl	$22, 16(%rsp)
	movabsq	$4555015426436491812, %rax      # imm = 0x3F36AA65972BEA24
	movq	%rax, 216(%rsp)
	movabsq	$4607182418800017408, %rax      # imm = 0x3FF0000000000000
	movq	%rax, 208(%rsp)
	movabsq	$4607002274814922588, %rax      # imm = 0x3FEF5C28F5C28F5C
	movq	%rax, 200(%rsp)
	movabsq	$4605927221099061257, %rax      # imm = 0x3FEB8A6782E96009
	movq	%rax, 192(%rsp)
	movl	$861, 100(%rsp)                 # imm = 0x35D
	movl	$2000000, 20(%rsp)              # imm = 0x1E8480
	xorl	%r12d, %r12d
.LBB2_1:                                # %._crit_edge
                                        # =>This Inner Loop Header: Depth=1
	leaq	(%rsp,%r12,2), %rdi
	addq	$42384, %rdi                    # imm = 0xA590
	leaq	282384(%rsp,%r12,2), %rbx
	leaq	(%rsp,%r12,2), %r14
	addq	$202384, %r14                   # imm = 0x31690
	leaq	(%rsp,%r12), %r15
	addq	$2384, %r15                     # imm = 0x950
	movl	$800, %edx                      # imm = 0x320
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$800, %edx                      # imm = 0x320
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$800, %edx                      # imm = 0x320
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$400, %edx                      # imm = 0x190
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	addq	$400, %r12                      # imm = 0x190
	cmpq	$40000, %r12                    # imm = 0x9C40
	jne	.LBB2_1
# %bb.2:                                # %._crit_edge97
	movq	24(%rsp), %rax                  # 8-byte Reload
	movq	8(%rax), %rdi
	movl	$.L.str.2, %esi
	callq	fopen
	leaq	336(%rsp), %rdi
	movl	$2048, %esi                     # imm = 0x800
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movq	%rax, %rdx
	callq	fgets
	testq	%rax, %rax
	je	.LBB2_8
# %bb.3:                                # %.preheader93.preheader
	xorl	%eax, %eax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	leaq	336(%rsp), %r14
.LBB2_4:                                # %.preheader93
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_6 Depth 2
	cmpl	$0, (%rsp)
	jle	.LBB2_7
# %bb.5:                                # %.lr.ph100.preheader
                                        #   in Loop: Header=BB2_4 Depth=1
	movslq	8(%rsp), %rbp                   # 4-byte Folded Reload
	xorl	%r12d, %r12d
.LBB2_6:                                # %.lr.ph100
                                        #   Parent Loop BB2_4 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%r14, %rdi
	callq	strdup
	movq	%rax, %r15
	leaq	1(%r12), %rbx
	movq	%rax, %rdi
	movl	%ebx, %esi
	callq	_Z8getfieldPci
	movq	%rax, %rdi
	callq	atof
	movslq	(%rsp), %r13
	movq	%rbp, %rax
	imulq	%r13, %rax
	leaq	(%rsp,%rax,8), %rax
	addq	$362384, %rax                   # imm = 0x58790
	movsd	%xmm0, (%rax,%r12,8)
	movq	%r15, %rdi
	callq	free
	movq	%rbx, %r12
	cmpq	%r13, %rbx
	jl	.LBB2_6
.LBB2_7:                                # %._crit_edge101
                                        #   in Loop: Header=BB2_4 Depth=1
	movq	8(%rsp), %rax                   # 8-byte Reload
	incl	%eax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	movq	%r14, %rdi
	movl	$2048, %esi                     # imm = 0x800
	movq	32(%rsp), %rdx                  # 8-byte Reload
	callq	fgets
	testq	%rax, %rax
	jne	.LBB2_4
.LBB2_8:                                # %._crit_edge104
	movq	32(%rsp), %rdi                  # 8-byte Reload
	callq	fclose
	movq	%rsp, %rbx
	movl	(%rbx), %esi
	imull	%esi, %esi
	shlq	$3, %rsi
	leaq	88(%rsp), %rdi
	callq	hipMalloc
	movl	(%rbx), %esi
	imull	%esi, %esi
	shlq	$3, %rsi
	leaq	72(%rsp), %rdi
	callq	hipMalloc
	movl	(%rbx), %esi
	imull	%esi, %esi
	shlq	$2, %rsi
	leaq	40(%rsp), %rdi
	callq	hipMalloc
	movl	(%rbx), %esi
	imull	%esi, %esi
	shlq	$3, %rsi
	leaq	64(%rsp), %rdi
	callq	hipMalloc
	leaq	80(%rsp), %rdi
	movl	$4, %esi
	callq	hipMalloc
	movl	(%rbx), %esi
	imull	%esi, %esi
	shlq	$3, %rsi
	leaq	56(%rsp), %rdi
	callq	hipMalloc
	movl	(%rbx), %esi
	imull	%esi, %esi
	shlq	$3, %rsi
	leaq	48(%rsp), %rdi
	callq	hipMalloc
	leaq	168(%rsp), %r13
	movl	$8, %esi
	movq	%r13, %rdi
	callq	hipMalloc
	leaq	128(%rsp), %r12
	movl	$4, %esi
	movq	%r12, %rdi
	callq	hipMalloc
	leaq	120(%rsp), %r15
	movl	$4, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	leaq	112(%rsp), %r14
	movl	$4, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	leaq	160(%rsp), %rbp
	movl	$8, %esi
	movq	%rbp, %rdi
	callq	hipMalloc
	leaq	104(%rsp), %rdi
	movl	$4, %esi
	callq	hipMalloc
	leaq	152(%rsp), %rdi
	movl	$8, %esi
	callq	hipMalloc
	leaq	144(%rsp), %rdi
	movl	$8, %esi
	callq	hipMalloc
	leaq	136(%rsp), %rdi
	movl	$8, %esi
	callq	hipMalloc
	movq	(%r13), %rdi
	leaq	192(%rsp), %rsi
	movl	$8, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r12), %rdi
	leaq	100(%rsp), %rsi
	movl	$4, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r15), %rdi
	movl	$4, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r14), %rdi
	leaq	16(%rsp), %rsi
	movl	$4, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbp), %rdi
	leaq	200(%rsp), %rsi
	movl	$8, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	104(%rsp), %rax
	movq	(%rax), %rdi
	leaq	20(%rsp), %rsi
	movl	$4, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	152(%rsp), %rax
	movq	(%rax), %rdi
	leaq	176(%rsp), %rsi
	movl	$8, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	144(%rsp), %rax
	movq	(%rax), %rdi
	leaq	208(%rsp), %rsi
	movl	$8, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	136(%rsp), %rax
	movq	(%rax), %rdi
	leaq	216(%rsp), %rsi
	movl	$8, %edx
	movl	$1, %ecx
	callq	hipMemcpy
	xorl	%edi, %edi
	callq	time
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movl	4(%rsp), %esi
	cmpl	$9999999, %esi                  # imm = 0x98967F
	jg	.LBB2_37
# %bb.9:                                # %.lr.ph122
	xorpd	%xmm0, %xmm0
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	leaq	42384(%rsp), %r14
	leaq	2384(%rsp), %rbx
	leaq	282384(%rsp), %r15
	leaq	122384(%rsp), %r12
	leaq	202384(%rsp), %r13
.LBB2_10:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_17 Depth 2
                                        #       Child Loop BB2_19 Depth 3
                                        #     Child Loop BB2_26 Depth 2
                                        #       Child Loop BB2_28 Depth 3
                                        #     Child Loop BB2_35 Depth 2
	imull	$989560465, %esi, %eax          # imm = 0x3AFB7E91
	addl	$3435968, %eax                  # imm = 0x346DC0
	rorl	$4, %eax
	cmpl	$429496, %eax                   # imm = 0x68DB8
	ja	.LBB2_12
# %bb.11:                               #   in Loop: Header=BB2_10 Depth=1
	movl	$.L.str.3, %edi
	movl	$10000000, %edx                 # imm = 0x989680
	xorl	%eax, %eax
	callq	printf
.LBB2_12:                               #   in Loop: Header=BB2_10 Depth=1
	movq	88(%rsp), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	leaq	362384(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	72(%rsp), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	40(%rsp), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	64(%rsp), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	80(%rsp), %rdi
	movl	$4, %edx
	leaq	4(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	56(%rsp), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	48(%rsp), %rdi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%r13, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967396, %rdi               # imm = 0x100000064
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_14
# %bb.13:                               #   in Loop: Header=BB2_10 Depth=1
	movq	88(%rsp), %rdi
	movq	72(%rsp), %rsi
	movq	40(%rsp), %rdx
	movq	64(%rsp), %rcx
	movq	80(%rsp), %r8
	movq	56(%rsp), %r9
	pushq	136(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	152(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	168(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	128(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	192(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	152(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	168(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	184(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	232(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	120(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z22__device_stub__calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	addq	$80, %rsp
	.cfi_adjust_cfa_offset -80
.LBB2_14:                               #   in Loop: Header=BB2_10 Depth=1
	callq	hipDeviceSynchronize
	movq	72(%rsp), %rsi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	40(%rsp), %rsi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	64(%rsp), %rsi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	56(%rsp), %rsi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	48(%rsp), %rsi
	movl	(%rsp), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%r13, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	4(%rsp), %eax
	imull	$989560465, %eax, %ecx          # imm = 0x3AFB7E91
	addl	$3435968, %ecx                  # imm = 0x346DC0
	rorl	$4, %ecx
	cmpl	$429496, %ecx                   # imm = 0x68DB8
	ja	.LBB2_22
# %bb.15:                               #   in Loop: Header=BB2_10 Depth=1
	movq	24(%rsp), %rax                  # 8-byte Reload
	movq	16(%rax), %rdx
	movl	16(%rsp), %ecx
	movl	$.L.str.4, %esi
	leaq	224(%rsp), %r15
	movq	%r15, %rdi
	movsd	.LCPI2_0(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	8(%rsp), %xmm1                  # 8-byte Reload
                                        # xmm1 = mem[0],zero
	movb	$2, %al
	callq	sprintf
	movl	$.L.str.5, %esi
	movq	%r15, %rdi
	callq	fopen
	movq	%rax, %r13
	movl	(%rsp), %eax
	testl	%eax, %eax
	jle	.LBB2_21
# %bb.16:                               # %.preheader92.preheader
                                        #   in Loop: Header=BB2_10 Depth=1
	xorl	%ebp, %ebp
.LBB2_17:                               # %.preheader92
                                        #   Parent Loop BB2_10 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB2_19 Depth 3
	testl	%eax, %eax
	jle	.LBB2_20
# %bb.18:                               # %.lr.ph106.preheader
                                        #   in Loop: Header=BB2_17 Depth=2
	xorl	%r12d, %r12d
.LBB2_19:                               # %.lr.ph106
                                        #   Parent Loop BB2_10 Depth=1
                                        #     Parent Loop BB2_17 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	imull	%ebp, %eax
	cltq
	leaq	(%rsp,%rax,8), %rax
	addq	$42384, %rax                    # imm = 0xA590
	movsd	(%rax,%r12,8), %xmm0            # xmm0 = mem[0],zero
	movl	$.L.str.6, %esi
	movq	%r13, %rdi
	movb	$1, %al
	callq	fprintf
	incq	%r12
	movslq	(%rsp), %rax
	cmpq	%rax, %r12
	jl	.LBB2_19
.LBB2_20:                               # %._crit_edge107
                                        #   in Loop: Header=BB2_17 Depth=2
	movl	$10, %edi
	movq	%r13, %rsi
	callq	fputc@PLT
	incl	%ebp
	movl	(%rsp), %eax
	cmpl	%eax, %ebp
	jl	.LBB2_17
.LBB2_21:                               # %._crit_edge109
                                        #   in Loop: Header=BB2_10 Depth=1
	movq	%r13, %rdi
	callq	fclose
	movl	4(%rsp), %eax
.LBB2_22:                               #   in Loop: Header=BB2_10 Depth=1
	movl	%eax, %ecx
	subl	20(%rsp), %ecx
	jle	.LBB2_31
# %bb.23:                               #   in Loop: Header=BB2_10 Depth=1
	imull	$-1101802277, %ecx, %ecx        # imm = 0xBE53D4DB
	addl	$1145312, %ecx                  # imm = 0x1179E0
	rorl	$4, %ecx
	cmpl	$143164, %ecx                   # imm = 0x22F3C
	ja	.LBB2_31
# %bb.24:                               #   in Loop: Header=BB2_10 Depth=1
	movq	24(%rsp), %rax                  # 8-byte Reload
	movq	16(%rax), %rdx
	movl	16(%rsp), %ecx
	movl	$.L.str.8, %esi
	leaq	224(%rsp), %r15
	movq	%r15, %rdi
	movsd	.LCPI2_0(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	8(%rsp), %xmm1                  # 8-byte Reload
                                        # xmm1 = mem[0],zero
	movb	$2, %al
	callq	sprintf
	movl	$.L.str.5, %esi
	movq	%r15, %rdi
	callq	fopen
	movq	%rax, %r13
	movl	(%rsp), %eax
	testl	%eax, %eax
	jle	.LBB2_30
# %bb.25:                               # %.preheader91.preheader
                                        #   in Loop: Header=BB2_10 Depth=1
	xorl	%ebp, %ebp
.LBB2_26:                               # %.preheader91
                                        #   Parent Loop BB2_10 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB2_28 Depth 3
	testl	%eax, %eax
	jle	.LBB2_29
# %bb.27:                               # %.lr.ph111.preheader
                                        #   in Loop: Header=BB2_26 Depth=2
	xorl	%r12d, %r12d
.LBB2_28:                               # %.lr.ph111
                                        #   Parent Loop BB2_10 Depth=1
                                        #     Parent Loop BB2_26 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	imull	%ebp, %eax
	cltq
	leaq	(%rsp,%rax,8), %rax
	addq	$122384, %rax                   # imm = 0x1DE10
	movsd	(%rax,%r12,8), %xmm0            # xmm0 = mem[0],zero
	movl	$.L.str.6, %esi
	movq	%r13, %rdi
	movb	$1, %al
	callq	fprintf
	incq	%r12
	movslq	(%rsp), %rax
	cmpq	%rax, %r12
	jl	.LBB2_28
.LBB2_29:                               # %._crit_edge112
                                        #   in Loop: Header=BB2_26 Depth=2
	movl	$10, %edi
	movq	%r13, %rsi
	callq	fputc@PLT
	incl	%ebp
	movl	(%rsp), %eax
	cmpl	%eax, %ebp
	jl	.LBB2_26
.LBB2_30:                               # %._crit_edge114
                                        #   in Loop: Header=BB2_10 Depth=1
	movq	%r13, %rdi
	callq	fclose
	movl	4(%rsp), %eax
.LBB2_31:                               # %thread-pre-split
                                        #   in Loop: Header=BB2_10 Depth=1
	movq	%rbx, %r15
	movq	%r14, %rbx
	cmpl	$2000000, %eax                  # imm = 0x1E8480
	jne	.LBB2_33
# %bb.32:                               #   in Loop: Header=BB2_10 Depth=1
	xorl	%edi, %edi
	callq	time
	movq	%rax, %r13
	movq	24(%rsp), %rax                  # 8-byte Reload
	movq	16(%rax), %rdx
	movl	$.L.str.9, %esi
	leaq	224(%rsp), %r14
	movq	%r14, %rdi
	xorl	%eax, %eax
	callq	sprintf
	movl	$.L.str.5, %esi
	movq	%r14, %rdi
	callq	fopen
	movq	%rax, %rbp
	subq	32(%rsp), %r13                  # 8-byte Folded Reload
	movl	$.L.str.10, %esi
	movq	%rax, %rdi
	movq	%r13, %rdx
	xorl	%eax, %eax
	callq	fprintf
	movq	%rbp, %rdi
	callq	fclose
.LBB2_33:                               #   in Loop: Header=BB2_10 Depth=1
	movl	(%rsp), %r14d
	testl	%r14d, %r14d
	jle	.LBB2_36
# %bb.34:                               # %.preheader.preheader
                                        #   in Loop: Header=BB2_10 Depth=1
	leaq	(,%r14,8), %r13
	xorl	%r12d, %r12d
	movq	%r14, %rbp
.LBB2_35:                               # %.preheader
                                        #   Parent Loop BB2_10 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	%r12d, %eax
	leaq	(%rsp,%rax,8), %rdi
	addq	$362384, %rdi                   # imm = 0x58790
	leaq	(%rsp,%rax,8), %rsi
	addq	$42384, %rsi                    # imm = 0xA590
	movq	%r13, %rdx
	callq	memcpy@PLT
	addl	%r14d, %r12d
	decq	%rbp
	jne	.LBB2_35
.LBB2_36:                               # %._crit_edge119
                                        #   in Loop: Header=BB2_10 Depth=1
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	addsd	176(%rsp), %xmm0
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movl	4(%rsp), %eax
	leal	1(%rax), %esi
	movl	%esi, 4(%rsp)
	cmpl	$9999999, %eax                  # imm = 0x98967F
	movq	%rbx, %r14
	movq	%r15, %rbx
	leaq	282384(%rsp), %r15
	leaq	122384(%rsp), %r12
	leaq	202384(%rsp), %r13
	jl	.LBB2_10
.LBB2_37:                               # %._crit_edge123
	movq	88(%rsp), %rdi
	callq	hipFree
	movq	72(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	64(%rsp), %rdi
	callq	hipFree
	movq	80(%rsp), %rdi
	callq	hipFree
	movq	56(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	168(%rsp), %rdi
	callq	hipFree
	movq	128(%rsp), %rdi
	callq	hipFree
	movq	120(%rsp), %rdi
	callq	hipFree
	movq	112(%rsp), %rdi
	callq	hipFree
	movq	160(%rsp), %rdi
	callq	hipFree
	movq	184(%rsp), %rsp                 # 8-byte Reload
	xorl	%eax, %eax
	addq	$442392, %rsp                   # imm = 0x6C018
	.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_end2:
	.size	main, .Lfunc_end2-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	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_, %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_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__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	.LBB4_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
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	","
	.size	.L.str, 2

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

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

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

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Iteration %d of %d\n"
	.size	.L.str.3, 20

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"ResultsCUDA%s/Results_POT_LIF_2D_Classic_sigma_%lf_R_%d_time_%lf_.dat"
	.size	.L.str.4, 70

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"w"
	.size	.L.str.5, 2

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"%lf,"
	.size	.L.str.6, 5

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"ResultsCUDA%s/Results_MPV_LIF_2D_Classic_sigma_%lf_R_%d_time_%lf_.dat"
	.size	.L.str.8, 70

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"ResultsCUDA%s/execTime.dat"
	.size	.L.str.9, 27

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Execution time for 2000 time units: %ld seconds\n"
	.size	.L.str.10, 49

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_"
	.size	.L__unnamed_1, 48

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

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

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

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_ ; -- Begin function _Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	.globl	_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
	.p2align	8
	.type	_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_,@function
_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_: ; @_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x8c
	s_load_b512 s[16:31], s[0:1], 0x0
	s_mov_b32 s33, 0
	s_mov_b32 s34, exec_lo
                                        ; implicit-def: $vgpr10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[6:7], null, s15, s2, v[0:1]
	s_load_b512 s[0:15], s[0:1], 0x40
	v_ashrrev_i32_e32 v7, 31, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[2:3], 3, v[6:7]
	v_lshlrev_b64 v[7:8], 2, v[6:7]
	v_add_co_u32 v0, vcc_lo, s16, v2
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v1, vcc_lo, s17, v3, vcc_lo
	global_load_b64 v[4:5], v[0:1], off
                                        ; implicit-def: $vgpr0_vgpr1
	s_waitcnt vmcnt(0)
	v_cmp_neq_f64_e64 s35, 0, v[4:5]
	v_cmpx_eq_f64_e32 0, v[4:5]
	s_cbranch_execz .LBB0_2
; %bb.1:
	v_add_co_u32 v0, vcc_lo, s20, v7
	v_add_co_ci_u32_e32 v1, vcc_lo, s21, v8, vcc_lo
	s_waitcnt lgkmcnt(0)
	s_load_b32 s0, s[0:1], 0x0
	s_mov_b32 s33, exec_lo
	global_load_b32 v10, v[0:1], off
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_cmp_le_i32_e32 vcc_lo, s0, v10
	s_and_not1_b32 s0, s35, exec_lo
	s_and_b32 s1, vcc_lo, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s35, s0, s1
.LBB0_2:                                ; %Flow164
	s_or_b32 exec_lo, exec_lo, s34
	s_delay_alu instid0(VALU_DEP_2)
	s_and_saveexec_b32 s34, s35
	s_cbranch_execz .LBB0_14
; %bb.3:
	s_waitcnt lgkmcnt(0)
	s_load_b32 s35, s[2:3], 0x0
	s_load_b32 s24, s[24:25], 0x0
	s_load_b64 s[0:1], s[30:31], 0x0
	s_load_b32 s30, s[4:5], 0x0
	s_load_b64 s[2:3], s[6:7], 0x0
	s_load_b32 s25, s[8:9], 0x0
	s_load_b64 s[4:5], s[10:11], 0x0
	s_load_b64 s[6:7], s[12:13], 0x0
	s_load_b64 s[8:9], s[14:15], 0x0
	v_add_co_u32 v11, vcc_lo, s20, v7
	v_add_co_ci_u32_e32 v12, vcc_lo, s21, v8, vcc_lo
	v_mov_b32_e32 v7, 0
	v_dual_mov_b32 v9, 0 :: v_dual_mov_b32 v8, 0
	s_mov_b32 s10, 0
	global_store_b32 v[11:12], v9, off
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_i32 s30, 0
	s_cbranch_scc1 .LBB0_9
; %bb.4:                                ; %.preheader.lr.ph
	s_ashr_i32 s12, s35, 31
	v_ashrrev_i32_e32 v11, 31, v6
	s_add_i32 s11, s35, s12
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_xor_b32 s11, s11, s12
	v_add_nc_u32_e32 v9, v6, v11
	v_cvt_f32_u32_e32 v7, s11
	s_sub_i32 s13, 0, s11
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v12, v9, v11
	v_rcp_iflag_f32_e32 v7, v7
	v_xor_b32_e32 v11, s12, v11
	s_sub_i32 s12, s35, s30
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v7, 0x4f7ffffe, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v7, v7
	v_mul_lo_u32 v8, s13, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v8, v7, v8
	v_add_nc_u32_e32 v9, v7, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v7, v12, v9
	v_mul_lo_u32 v8, v7, s11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v8, v12, v8
	v_add_nc_u32_e32 v12, 1, v7
	v_subrev_nc_u32_e32 v13, s11, v8
	v_cmp_le_u32_e32 vcc_lo, s11, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v7, v7, v12 :: v_dual_cndmask_b32 v8, v8, v13
	v_add_nc_u32_e32 v12, 1, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s11, v8
	v_cndmask_b32_e32 v7, v7, v12, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v7, v7, v11
	v_sub_nc_u32_e32 v7, v7, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v8, v7, s35
	v_add_nc_u32_e32 v12, s12, v7
	v_sub_nc_u32_e32 v6, v6, v8
	v_mov_b32_e32 v7, 0
	v_mov_b32_e32 v8, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_nc_u32_e32 v6, s12, v6
	s_lshl_b32 s12, s30, 1
	v_add_nc_u32_e32 v11, s12, v12
	s_delay_alu instid0(VALU_DEP_2)
	v_add_nc_u32_e32 v13, s12, v6
.LBB0_5:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_6 Depth 2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v14, 31, v12
	s_mov_b32 s12, 0
	v_add_nc_u32_e32 v15, v12, v14
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v15, v15, v14
	v_mul_hi_u32 v16, v15, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v16, v16, s11
	v_sub_nc_u32_e32 v15, v15, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v16, s11, v15
	v_cmp_le_u32_e32 vcc_lo, s11, v15
	v_cndmask_b32_e32 v15, v15, v16, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v16, s11, v15
	v_cmp_le_u32_e32 vcc_lo, s11, v15
	v_dual_cndmask_b32 v15, v15, v16 :: v_dual_mov_b32 v16, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v15, v15, v14
	v_sub_nc_u32_e32 v14, v15, v14
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v14, v14, s35
	v_ashrrev_i32_e32 v15, 31, v14
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[14:15], 3, v[14:15]
	v_add_co_u32 v14, vcc_lo, s16, v14
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v15, vcc_lo, s17, v15, vcc_lo
.LBB0_6:                                ;   Parent Loop BB0_5 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v17, 31, v16
	v_add_nc_u32_e32 v18, v16, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v18, v18, v17
	v_mul_hi_u32 v19, v18, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v19, v19, s11
	v_sub_nc_u32_e32 v18, v18, v19
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v19, s11, v18
	v_cmp_le_u32_e32 vcc_lo, s11, v18
	v_cndmask_b32_e32 v18, v18, v19, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v19, s11, v18
	v_cmp_le_u32_e32 vcc_lo, s11, v18
	v_cndmask_b32_e32 v18, v18, v19, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v18, v18, v17
	v_sub_nc_u32_e32 v17, v18, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v18, 31, v17
	v_lshlrev_b64 v[17:18], 3, v[17:18]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v17, vcc_lo, v14, v17
	v_add_co_ci_u32_e32 v18, vcc_lo, v15, v18, vcc_lo
	v_cmp_ge_i32_e32 vcc_lo, v16, v13
	global_load_b64 v[17:18], v[17:18], off
	s_or_b32 s12, vcc_lo, s12
	s_waitcnt vmcnt(0)
	v_add_f64 v[17:18], v[4:5], -v[17:18]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_f64 v[7:8], v[7:8], v[17:18]
	v_add_nc_u32_e32 v17, 1, v16
	v_mov_b32_e32 v16, v17
	s_and_not1_b32 exec_lo, exec_lo, s12
	s_cbranch_execnz .LBB0_6
; %bb.7:                                ; %._crit_edge
                                        ;   in Loop: Header=BB0_5 Depth=1
	s_or_b32 exec_lo, exec_lo, s12
	v_add_nc_u32_e32 v14, 1, v12
	v_cmp_ge_i32_e32 vcc_lo, v12, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_mov_b32_e32 v12, v14
	s_or_b32 s10, vcc_lo, s10
	s_and_not1_b32 exec_lo, exec_lo, s10
	s_cbranch_execnz .LBB0_5
; %bb.8:                                ; %Flow161
	s_or_b32 exec_lo, exec_lo, s10
.LBB0_9:                                ; %._crit_edge111
	v_add_f64 v[11:12], s[6:7], -v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f64 v[6:7], s[8:9], v[7:8], v[11:12]
	v_add_co_u32 v8, vcc_lo, s18, v2
	v_add_co_ci_u32_e32 v9, vcc_lo, s19, v3, vcc_lo
	v_fma_f64 v[6:7], s[4:5], v[6:7], v[4:5]
	v_add_co_u32 v4, vcc_lo, s22, v2
	v_add_co_ci_u32_e32 v5, vcc_lo, s23, v3, vcc_lo
	global_store_b64 v[8:9], v[6:7], off
	global_load_b64 v[6:7], v[4:5], off
	s_waitcnt vmcnt(0)
	v_add_f64 v[6:7], s[4:5], v[6:7]
	global_store_b64 v[4:5], v[6:7], off
	global_load_b64 v[6:7], v[8:9], off
	s_waitcnt vmcnt(0)
	v_cmp_le_f64_e32 vcc_lo, s[2:3], v[6:7]
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_13
; %bb.10:
	v_mov_b32_e32 v6, 0
	s_cmp_lt_i32 s24, s25
	s_delay_alu instid0(VALU_DEP_1)
	v_mov_b32_e32 v7, v6
	global_store_b64 v[8:9], v[6:7], off
	s_cbranch_scc1 .LBB0_12
; %bb.11:
	v_add_co_u32 v8, vcc_lo, s28, v2
	v_add_co_ci_u32_e32 v9, vcc_lo, s29, v3, vcc_lo
	v_add_co_u32 v2, vcc_lo, s26, v2
	global_load_b64 v[11:12], v[4:5], off
	global_load_b64 v[13:14], v[8:9], off
	v_add_co_ci_u32_e32 v3, vcc_lo, s27, v3, vcc_lo
	s_mov_b32 s4, 0x54442eea
	s_mov_b32 s5, 0x401921fb
	global_load_b64 v[15:16], v[2:3], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[11:12], v[13:14], v[11:12]
	s_waitcnt vmcnt(0)
	v_fma_f64 v[13:14], v[15:16], v[13:14], s[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[11:12], s[0:1], v[11:12]
	v_div_scale_f64 v[15:16], null, v[11:12], v[11:12], v[13:14]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[17:18], v[15:16]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[19:20], -v[15:16], v[17:18], 1.0
	v_fma_f64 v[17:18], v[17:18], v[19:20], v[17:18]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[19:20], -v[15:16], v[17:18], 1.0
	v_fma_f64 v[17:18], v[17:18], v[19:20], v[17:18]
	v_div_scale_f64 v[19:20], vcc_lo, v[13:14], v[11:12], v[13:14]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[21:22], v[19:20], v[17:18]
	v_fma_f64 v[15:16], -v[15:16], v[21:22], v[19:20]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[15:16], v[15:16], v[17:18], v[21:22]
	v_div_fixup_f64 v[11:12], v[15:16], v[11:12], v[13:14]
	global_store_b64 v[2:3], v[11:12], off
	global_load_b64 v[2:3], v[4:5], off
	global_load_b64 v[11:12], v[8:9], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[2:3], s[0:1], v[2:3]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[2:3], v[11:12], v[2:3]
	global_store_b64 v[8:9], v[2:3], off
.LBB0_12:
	global_store_b64 v[4:5], v[6:7], off
.LBB0_13:                               ; %Flow
	s_or_b32 exec_lo, exec_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 s33, s33, exec_lo
.LBB0_14:                               ; %Flow165
	s_or_b32 exec_lo, exec_lo, s34
	s_waitcnt lgkmcnt(0)
	s_and_saveexec_b32 s0, s33
	s_cbranch_execz .LBB0_16
; %bb.15:
	v_add_nc_u32_e32 v2, 1, v10
	global_store_b32 v[0:1], v2, off
.LBB0_16:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 384
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 23
		.amdhsa_next_free_sgpr 36
		.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	_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_, .Lfunc_end0-_Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1372
; NumSgprs: 38
; NumVgprs: 23
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 4
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 38
; NumVGPRsForWavesPerEU: 23
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.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:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         80
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         88
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         96
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         104
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         112
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         120
        .size:           8
        .value_kind:     global_buffer
      - .offset:         128
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         132
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         136
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         140
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         142
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         144
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         146
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         148
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         150
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         168
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         176
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         184
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         192
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 384
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     38
    .sgpr_spill_count: 0
    .symbol:         _Z7calcPotPdS_PiS_S0_S_S_S_S0_S0_S0_S_S0_S_S_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     23
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
