#include <stdio.h>
#include <time.h>

#define PI 3.1415926535897932384
#define mu0 4*PI*1e-7
//Threads per block is capped at 1024 for hardware reasons
//In some cases using a smaller number of threads per block will be more efficient
#define threadsPerBlock 1024
//Max grid points is to defined in order to allocate shared memory for each block
#define MaxGridPoints 6144


/*
NOTE:
grid is technically a 3-dimensional variable length array. In order to copy
this 3d array, it needed to be compressed into a single dimension.
Rather than decompressing this and wasting valuable resources all computation is done in line.
i.e. all blocks will write to adjacent memory
the position in memory is determined by three factors: the block numnber, the r
segment value and the z segment value
Each block is given an initial offset. This offset gives the block a unique space inside
of the overall shared array in which to write its data. Each block's grid is
r driven meaning grid(r,z) = grid[r*zMax + z], thus the position of each point in the
overall array is given by grids[offset + r*zMax + z]
*/


//This function will be called to initialize all data in the grid, Each thread will
//calculate the initial condion for a variable number of points and each block will have a different
//set of initial conditions all passed to device memory already
__global__ void init(double *grids, double *Ils, double *dIs, double *ldrs, double *rlengths, long int *grid_sizes, int *zsegs, long int *offsets){
  // get needed data from memory, all memory is block adjacent
  //block index determines which problem you are working on

  //Current on the left side of the grid
  double Il = Ils[blockIdx.x];
  //amount the current increments per segment in the z direction
  //This model assumes linear change
  double dI = dIs[blockIdx.x];
  //Physical distance change between r segments
  double ldr = ldrs[blockIdx.x];
  //total r length
  double rlength = rlengths[blockIdx.x];
  //Total number of grid points
  long int grid_size = grid_sizes[blockIdx.x];
  //number of z segmnents
  int zseg = zsegs[blockIdx.x];
  //Offset of the current grid
  //Because data needs to be copied to the gpu, it is compressed into one dimension
  //thus all grid data lives in the same array and the offset for that grid determines the satrt of this grid
  //This is explained in more detail above
  long int gridStart = offsets[blockIdx.x];

  //Find how many points per thread will be executing

  //remainder tells you how many threads will have one extra point to cover
  int rem = grid_size%threadsPerBlock;
  //division tells you how many points ALL theads will cover
  int divi = grid_size/threadsPerBlock;
  long int start, fin;
  //If threadIdx.x is less than rem the current thread will recieve one extra point to calculate
  if(threadIdx.x<rem){
    start = threadIdx.x*(divi+1);
    fin = start + divi + 1;
  }
  else{
    //rem is added here because at this point a "rem" number of threads will have one extra point to compute
    //overall this is "rem" extra points and this is accounted for by adding rem
    start = threadIdx.x*divi + rem;
    fin = start + divi;
  }

  //get initial conditions for all points per thread

  //Loops over each point a thread takes care of
  //NOTE: If a thread is not supposed to execute anything it will have the same
  //value for "start" and "fin" and will not enter the loop
  for(int i = start; i<fin; i++){
    //Gets current r and z segmenty values
    //Because grid is "r-driven", dividing by zMax = the r value and moding by
    //zMax gves the current z value
    //For more info check out the "note" at the top
    long int r_seg_val = i/(zseg+2);
    long int z_seg_val = i%(zseg+2);
    //Sets initial condition at the given point based on an equation given by Dr. Sankaran
    grids[i + gridStart] = (1-(r_seg_val*r_seg_val*ldr*ldr/(3*rlength*rlength)))*3*mu0*(Il + z_seg_val*dI)*r_seg_val*ldr/(4*PI*rlength*rlength);
  }
}

//This function funs until the number of timesteps needed is completed
//again this function passes data by memory copy in the form of arrays
//each block excercises a different set of input params
__global__ void run(double *grids, double *r_augs, double *z_augs, long int *allMaxSteps, long int *grid_sizes, int *rsegs, int *zsegs, long int *offsets){
  //get data from initialization arrays
  //block index determines the problem that is being worked on

  //r_aug and z_aug are constants that multiple values need to update to 1 timestep ahead
  //r_aug is used for changes in r and z_aug is used for changes in z
  double r_aug = r_augs[blockIdx.x];
  double z_aug = z_augs[blockIdx.x];
  //max steps is the number of steps needed to be completed for the simulation
  //to be finished
  long int maxSteps = allMaxSteps[blockIdx.x];
  //total grid size for this block's problem
  long int grid_size = grid_sizes[blockIdx.x];
  //number of segments in r and z directions
  int rseg = rsegs[blockIdx.x];
  int zseg = zsegs[blockIdx.x];
  //offset for block's memory locations
  //More info at top
  long int gridStart = offsets[blockIdx.x];

  //find how many points each thread will execute for

  //rem represents the number of threads that will have one more than divi points
  //to take care of
  int rem = grid_size%threadsPerBlock;
  //divi is the amount of points each thread will have to take care of at a minimum
  int divi = grid_size/threadsPerBlock;
  int start, fin;
  if(threadIdx.x<rem){
    start = threadIdx.x*(divi+1);
    fin = start + divi + 1;
  }
  else{
    //rem is added here because at this point a "rem" number of threads will have one extra point to compute
    //overall this is "rem" extra points and this is accounted for by adding rem
    start = threadIdx.x*divi + rem;
    fin = start + divi;
  }
  long int steps = 0;

  //This is shared memory
  //It is quick access but there is a limited amount and can only be
  //accessed inside each block
  //the extern keyword states that the size of this memory is passed in at the
  //function call. This allows a variable amount of shared memory to be used
  extern __shared__ double grid_new_s[];
  extern __shared__ double grid_old_s[];

  //Initialize shared memory
  //gridStart is the offset for this block
  for(int i = start; i<fin; i++){
    grid_new_s[i] = grids[i + gridStart];
  }
  //wait for all threads
  __syncthreads();

  while(steps<maxSteps){
    //copy all points from new to old
    for(int i = start; i<fin; i++){
      grid_old_s[i] = grid_new_s[i];
    }
    //wait for all threads
    __syncthreads();
    for(int i = start; i<fin; i++){
      //because the grid is row driven
      //i / zMax = rvalue
      //i % zMax = zvalue
      int r = i/(zseg+2);
      int z = i%(zseg+2);

      //leave boundary conditions alone as specified by the problem
      if(z != 0 && z != zseg+1){
        if(r!= 0 && r!= rseg+1){
          //when r = 1 a phontom point needs to be used as the slope change from
          //0 -> 1 is too large
          //in this case the B field at -1 is assumed to be of the same magnitude as at 1
          //but in the opposite direction
          //this yields the folloeing equation
          if(r==1){
            grid_new_s[i] += r_aug*(2*grid_old_s[i+(zseg+2)] - 4*grid_old_s[i]) +
              z_aug * (grid_old_s[i+1] - 2*grid_old_s[i] + grid_old_s[i-1]);
          }
          //normal update function
          else{
            grid_new_s[i] += r_aug*((1+(1/(2*r)))*grid_old_s[i+(zseg+2)] + (-2-(1/(r*r)))*grid_old_s[i] + (1-(1/(2*r)))*grid_old_s[i-(zseg+2)])
              +z_aug*(grid_old_s[i+1] - 2*grid_old_s[i] + grid_old_s[i-1]);
          }
        }
      }
    }
    steps++;
    //wait for all threads
    __syncthreads();
  }
  //copy all values back to main memory
  for(int i = start; i<fin; i++){
    grids[i + gridStart] = grid_new_s[i];
  }
}

int main(){
  //string label in input FILE
  //no use just for clarity
  char label[256];

  double *Il_h, *Ir_h, *rlength_h, *eta_h, *tstep_h, *ldr_h, *ldz_h, *zlength_h, *dI_h, *r_aug_h, *z_aug_h, *grids_h;
  int *rseg_h, *zseg_h;
  long int *totsteps_h, *grid_size_h, *offsets_h;

  double *Il_d, *rlength_d, *ldr_d, *dI_d, *r_aug_d, *z_aug_d, *grids_d;
  int *rseg_d, *zseg_d;
  long int *totsteps_d, *grid_size_d, *offsets_d;

  int testcases;
  //printf("How many test cases? ");
  //gets number of test cases
  scanf("%d", &testcases);

  //these sizes are used for malloc ops
  //the size is based on the amount of test cases and the
  //type of data this array holds
  size_t doubleSize = testcases*sizeof(double);
  size_t intSize = testcases*sizeof(int);
  size_t longIntSize = testcases*sizeof(long int);

  //host memory allocation
  Il_h = (double*)malloc(doubleSize);
  Ir_h = (double*)malloc(doubleSize);
  rlength_h = (double*)malloc(doubleSize);
  eta_h = (double*)malloc(doubleSize);
  tstep_h = (double*)malloc(doubleSize);
  ldr_h = (double*)malloc(doubleSize);
  ldz_h = (double*)malloc(doubleSize);
  zlength_h = (double*)malloc(doubleSize);
  dI_h = (double*)malloc(doubleSize);
  r_aug_h = (double*)malloc(doubleSize);
  z_aug_h = (double*)malloc(doubleSize);
  rseg_h = (int*)malloc(intSize);
  zseg_h = (int*)malloc(intSize);
  totsteps_h = (long int*)malloc(longIntSize);
  grid_size_h = (long int*)malloc(longIntSize);
  offsets_h = (long int*)malloc(longIntSize);

  //device memory allocation
  cudaMalloc(&Il_d, doubleSize);
  cudaMalloc(&rlength_d, doubleSize);
  cudaMalloc(&ldr_d, doubleSize);
  cudaMalloc(&dI_d, doubleSize);
  cudaMalloc(&r_aug_d, doubleSize);
  cudaMalloc(&z_aug_d, doubleSize);
  cudaMalloc(&rseg_d, intSize);
  cudaMalloc(&zseg_d, intSize);
  cudaMalloc(&totsteps_d, longIntSize);
  cudaMalloc(&grid_size_d, longIntSize);
  cudaMalloc(&offsets_d, longIntSize);

  long int total_grid_size = 0;
  for(int counter = 0; counter < testcases; counter++){
    //label is unused
    scanf("%s", &label);
    //value of current when z = 0
    //printf("What is your left I? ");
    scanf("%lf", Il_h + counter);
    //value of current when z = zMax
    //printf("What is your right I? ");
    scanf("%lf", Ir_h + counter);
    //total r value
    //printf("What is the radius of your rod? ");
    scanf("%lf", rlength_h + counter);
    //total z value
    //printf("What is the length of your rod? ");
    scanf("%lf", zlength_h + counter);
    //value of eta (proportional to diffusivity)
    //printf("What is eta? ");
    scanf("%lf", eta_h + counter);
    //number of r segments
    //printf("How many segments would you like per radius? ");
    scanf("%d", rseg_h + counter);
    //number of z segments
    //printf("How many segments would you like per length? ");
    scanf("%d", zseg_h + counter);
    //length of each r segment
    ldr_h[counter] = rlength_h[counter]/(rseg_h[counter]+1);
    //length of each z segment
    ldz_h[counter] = zlength_h[counter]/(zseg_h[counter]+1);
    double smallest = ldr_h[counter];
    if(ldz_h[counter] < smallest)
      smallest = ldz_h[counter];
    //determines tstep that ensures stability
    //0.125 is derived from 2^-(num dimensions + 1)
    tstep_h[counter] = 0.125*smallest*smallest*mu0/eta_h[counter];
    //gets total run time
    double tottime;
    //printf("How long would you like to run? ");
    scanf("%lf", &tottime);
    //total steps is an integer that truncates remainder
    totsteps_h[counter] = tottime/tstep_h[counter];
    //dI is change in I per each z segment
    //zseg + 2 to account for boundary conditions
    dI_h[counter] = (Ir_h[counter] - Il_h[counter]) / (zseg_h[counter]+2);
    //r and z aug are used for updating values for increasing time
    //as per numerical update
    r_aug_h[counter] = eta_h[counter]*tstep_h[counter]/(mu0*ldr_h[counter]*ldr_h[counter]);
    z_aug_h[counter] = eta_h[counter]*tstep_h[counter]/(mu0*ldz_h[counter]*ldz_h[counter]);
    //grid size accounts for boundary conditions
    grid_size_h[counter] = (rseg_h[counter] + 2)*(zseg_h[counter] + 2);
    //offsets says where in memory this problems data starts
    //more info above
    offsets_h[counter] = total_grid_size;
    //Total grid size is the size of all needed grid memory
    total_grid_size += grid_size_h[counter];
  }


  size_t gridsSize = total_grid_size*sizeof(double);

  //allocate variable amount of grid memory on both host and device
  grids_h = (double*)malloc(gridsSize);
  cudaMalloc(&grids_d, gridsSize);

  //copy memory down to device
  cudaMemcpy(Il_d, Il_h, doubleSize, cudaMemcpyHostToDevice);
  cudaMemcpy(rlength_d, rlength_h, doubleSize, cudaMemcpyHostToDevice);
  cudaMemcpy(ldr_d, ldr_h, doubleSize, cudaMemcpyHostToDevice);
  cudaMemcpy(dI_d, dI_h, doubleSize, cudaMemcpyHostToDevice);
  cudaMemcpy(r_aug_d, r_aug_h, doubleSize, cudaMemcpyHostToDevice);
  cudaMemcpy(z_aug_d, z_aug_h, doubleSize, cudaMemcpyHostToDevice);
  cudaMemcpy(rseg_d, rseg_h, intSize, cudaMemcpyHostToDevice);
  cudaMemcpy(zseg_d, zseg_h, intSize, cudaMemcpyHostToDevice);
  cudaMemcpy(totsteps_d, totsteps_h, longIntSize, cudaMemcpyHostToDevice);
  cudaMemcpy(grid_size_d, grid_size_h, longIntSize, cudaMemcpyHostToDevice);
  cudaMemcpy(offsets_d, offsets_h, longIntSize, cudaMemcpyHostToDevice);

  //initialize all problems with initial values as per Dr. Sankaran's instructions
  init<<<testcases,threadsPerBlock>>>(grids_d, Il_d, dI_d, ldr_d, rlength_d, grid_size_d, zseg_d, offsets_d);
  //wait for all device implementation to finish
  cudaDeviceSynchronize();

  //copy initial grid data back to host
  cudaMemcpy(grids_h, grids_d, gridsSize, cudaMemcpyDeviceToHost);


  //outputs all problems initial conditions
  FILE *myfile;
  long int i;
  for(int counter = 0; counter < testcases; counter++){
    long int gridStart = offsets_h[counter];
    char init[] = "initXX.txt";
    int tens = counter / 10;
    int ones = counter % 10;
    init[4] = tens + '0';
    init[5] = ones + '0';
    myfile = fopen(init, "w");
    for(i = 0; i< zseg_h[counter]+1; i++)
      fprintf(myfile, "%lf ", i*ldz_h[counter]);
    fprintf(myfile, "%lf\n", i*ldz_h[counter]);

    for(i = 0; i< rseg_h[counter]+1; i++)
      fprintf(myfile, "%lf ", i*ldr_h[counter]);
    fprintf(myfile, "%lf\n", i*ldr_h[counter]);

    for(i = 0; i< (rseg_h[counter] + 2)*(zseg_h[counter]+2); i++){
      if(i%(zseg_h[counter]+2)==zseg_h[counter]+1)
        fprintf(myfile, "%lf\n", grids_h[i + gridStart]);
      else
        fprintf(myfile, "%lf ", grids_h[i + gridStart]);
    }

    fclose(myfile);
  }



  clock_t begin, end;
  double time_spent;
  begin = clock();

  //run
  //size capped at MaxGridPoints, grid points per problem
  run<<<testcases,threadsPerBlock, MaxGridPoints*sizeof(double)>>>(grids_d, r_aug_d, z_aug_d, totsteps_d, grid_size_d, rseg_d, zseg_d, offsets_d);
  //wait for all device implementation to finish
  cudaDeviceSynchronize();

  end = clock();
  time_spent = (double)(end - begin) / CLOCKS_PER_SEC;

  //copy final grid data back to host
  cudaMemcpy(grids_h, grids_d, gridsSize, cudaMemcpyDeviceToHost);

  //output all problems final grid data
  for(int counter = 0; counter < testcases; counter++){
    long int gridStart = offsets_h[counter];
    char init[] = "resXX.txt";
    int tens = counter / 10;
    int ones = counter % 10;
    init[3] = tens + '0';
    init[4] = ones + '0';
    myfile = fopen(init, "w");
    for(i = 0; i< zseg_h[counter]+1; i++)
      fprintf(myfile, "%lf ", i*ldz_h[counter]);
    fprintf(myfile, "%lf\n", i*ldz_h[counter]);

    for(i = 0; i< rseg_h[counter]+1; i++)
      fprintf(myfile, "%lf ", i*ldr_h[counter]);
    fprintf(myfile, "%lf\n", i*ldr_h[counter]);

    for(i = 0; i< (rseg_h[counter] + 2)*(zseg_h[counter]+2); i++){
      if(i%(zseg_h[counter]+2)==zseg_h[counter]+1)
        fprintf(myfile, "%lf\n", grids_h[i + gridStart]);
      else
        fprintf(myfile, "%lf ", grids_h[i + gridStart]);
    }

    fclose(myfile);
  }

  //free host memory
  free(Il_h);
  free(Ir_h);
  free(rlength_h);
  free(eta_h);
  free(tstep_h);
  free(ldr_h);
  free(ldz_h);
  free(zlength_h);
  free(dI_h);
  free(r_aug_h);
  free(z_aug_h);
  free(rseg_h);
  free(zseg_h);
  free(totsteps_h);
  free(grid_size_h);
  free(offsets_h);

  //free device memory
  cudaFree(Il_d);
  cudaFree(rlength_d);
  cudaFree(ldr_d);
  cudaFree(dI_d);
  cudaFree(r_aug_d);
  cudaFree(z_aug_d);
  cudaFree(rseg_d);
  cudaFree(zseg_d);
  cudaFree(totsteps_d);
  cudaFree(grid_size_d);
  cudaFree(offsets_d);



  printf("\n------------------------------------\nExecution took: %lf sec\n", time_spent);
  return 0;
}

	.file	"tmpxft_00296313_00000000-6_run.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0, @function
_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0:
.LFB2176:
	.cfi_startproc
	jmp	cudaMalloc@PLT
	.cfi_endproc
.LFE2176:
	.size	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0, .-_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	.globl	_Z38__device_stub__Z4initPdS_S_S_S_PlPiS0_PdS_S_S_S_PlPiS0_
	.type	_Z38__device_stub__Z4initPdS_S_S_S_PlPiS0_PdS_S_S_S_PlPiS0_, @function
_Z38__device_stub__Z4initPdS_S_S_S_PlPiS0_PdS_S_S_S_PlPiS0_:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movq	%rax, 8(%rsp)
	movq	232(%rsp), %rax
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 184(%rsp)
	movq	%rsp, %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L3
	pushq	80(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z4initPdS_S_S_S_PlPiS0_(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 240
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L3:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L5
	call	__stack_chk_fail@PLT
.L5:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z38__device_stub__Z4initPdS_S_S_S_PlPiS0_PdS_S_S_S_PlPiS0_, .-_Z38__device_stub__Z4initPdS_S_S_S_PlPiS0_PdS_S_S_S_PlPiS0_
	.globl	_Z4initPdS_S_S_S_PlPiS0_
	.type	_Z4initPdS_S_S_S_PlPiS0_, @function
_Z4initPdS_S_S_S_PlPiS0_:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z4initPdS_S_S_S_PlPiS0_PdS_S_S_S_PlPiS0_
	.cfi_endproc
.LFE2053:
	.size	_Z4initPdS_S_S_S_PlPiS0_, .-_Z4initPdS_S_S_S_PlPiS0_
	.globl	_Z39__device_stub__Z3runPdS_S_PlS0_PiS1_S0_PdS_S_PlS0_PiS1_S0_
	.type	_Z39__device_stub__Z3runPdS_S_PlS0_PiS1_S0_PdS_S_PlS0_PiS1_S0_, @function
_Z39__device_stub__Z3runPdS_S_PlS0_PiS1_S0_PdS_S_PlS0_PiS1_S0_:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movq	%rax, 8(%rsp)
	movq	232(%rsp), %rax
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 184(%rsp)
	movq	%rsp, %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	pushq	80(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z3runPdS_S_PlS0_PiS1_S0_(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 240
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L9:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z39__device_stub__Z3runPdS_S_PlS0_PiS1_S0_PdS_S_PlS0_PiS1_S0_, .-_Z39__device_stub__Z3runPdS_S_PlS0_PiS1_S0_PdS_S_PlS0_PiS1_S0_
	.globl	_Z3runPdS_S_PlS0_PiS1_S0_
	.type	_Z3runPdS_S_PlS0_PiS1_S0_, @function
_Z3runPdS_S_PlS0_PiS1_S0_:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z3runPdS_S_PlS0_PiS1_S0_PdS_S_PlS0_PiS1_S0_
	.cfi_endproc
.LFE2055:
	.size	_Z3runPdS_S_PlS0_PiS1_S0_, .-_Z3runPdS_S_PlS0_PiS1_S0_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"%d"
.LC1:
	.string	"%s"
.LC2:
	.string	"%lf"
.LC8:
	.string	"w"
.LC9:
	.string	"%lf "
.LC10:
	.string	"%lf\n"
.LC12:
	.string	"\n------------------------------------\nExecution took: %lf sec\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC0(%rip), %rdi
	xorl	%r15d, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$584, %rsp
	.cfi_def_cfa_offset 640
	movq	%fs:40, %rax
	movq	%rax, 568(%rsp)
	xorl	%eax, %eax
	leaq	172(%rsp), %rsi
	call	__isoc23_scanf@PLT
	movslq	172(%rsp), %rax
	leaq	0(,%rax,8), %rbx
	salq	$2, %rax
	movq	%rbx, %rdi
	movq	%rbx, (%rsp)
	movq	%rax, 8(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 40(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 88(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 48(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 56(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 152(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 16(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 64(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 96(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 104(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 112(%rsp)
	call	malloc@PLT
	movq	8(%rsp), %rdi
	movq	%rax, 120(%rsp)
	call	malloc@PLT
	movq	8(%rsp), %rdi
	movq	%rax, 24(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, %rbx
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rbx, %r13
	movq	%rax, 128(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 136(%rsp)
	call	malloc@PLT
	movq	(%rsp), %rsi
	leaq	176(%rsp), %rdi
	movq	%rax, 32(%rsp)
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	(%rsp), %rsi
	leaq	184(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	(%rsp), %rsi
	leaq	192(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	(%rsp), %rsi
	leaq	200(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	(%rsp), %rsi
	leaq	208(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	(%rsp), %rsi
	leaq	216(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rsi
	leaq	232(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rsi
	leaq	240(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	(%rsp), %rsi
	leaq	248(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	(%rsp), %rsi
	leaq	256(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	(%rsp), %rsi
	leaq	264(%rsp), %rdi
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	24(%rsp), %rax
	xorl	%ebp, %ebp
	xorl	%r14d, %r14d
	movq	%rax, 144(%rsp)
	movq	%rax, %r12
.L15:
	cmpl	%ebp, 172(%rsp)
	jle	.L45
	leaq	312(%rsp), %rsi
	leaq	.LC1(%rip), %rdi
	xorl	%eax, %eax
	incl	%ebp
	call	__isoc23_scanf@PLT
	movq	40(%rsp), %rax
	leaq	.LC2(%rip), %rdi
	leaq	(%rax,%r15), %rsi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movq	88(%rsp), %rax
	leaq	.LC2(%rip), %rdi
	leaq	(%rax,%r15), %rsi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movq	48(%rsp), %rax
	leaq	.LC2(%rip), %rdi
	leaq	(%rax,%r15), %rsi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movq	96(%rsp), %rax
	leaq	.LC2(%rip), %rdi
	leaq	(%rax,%r15), %rsi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movq	56(%rsp), %rax
	leaq	.LC2(%rip), %rdi
	leaq	(%rax,%r15), %rsi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movq	%r12, %rsi
	leaq	.LC0(%rip), %rdi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movq	%r13, %rsi
	xorl	%eax, %eax
	addq	$4, %r13
	leaq	.LC0(%rip), %rdi
	addq	$4, %r12
	call	__isoc23_scanf@PLT
	movl	-4(%r12), %eax
	leaq	288(%rsp), %rsi
	leaq	.LC2(%rip), %rdi
	incl	%eax
	cvtsi2sdl	%eax, %xmm0
	movq	48(%rsp), %rax
	movsd	(%rax,%r15), %xmm1
	movq	16(%rsp), %rax
	divsd	%xmm0, %xmm1
	movsd	%xmm1, (%rax,%r15)
	movl	-4(%r13), %eax
	incl	%eax
	cvtsi2sdl	%eax, %xmm0
	movq	96(%rsp), %rax
	movsd	(%rax,%r15), %xmm2
	movq	64(%rsp), %rax
	divsd	%xmm0, %xmm2
	movsd	.LC3(%rip), %xmm0
	movaps	%xmm2, %xmm5
	movsd	%xmm2, (%rax,%r15)
	movq	56(%rsp), %rax
	minsd	%xmm1, %xmm5
	movsd	%xmm2, 80(%rsp)
	mulsd	%xmm5, %xmm0
	mulsd	%xmm5, %xmm0
	mulsd	.LC4(%rip), %xmm0
	mulsd	.LC5(%rip), %xmm0
	mulsd	.LC6(%rip), %xmm0
	divsd	(%rax,%r15), %xmm0
	movq	152(%rsp), %rax
	movsd	%xmm0, 72(%rsp)
	movsd	%xmm0, (%rax,%r15)
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movsd	72(%rsp), %xmm0
	movsd	288(%rsp), %xmm1
	movq	128(%rsp), %rcx
	movsd	80(%rsp), %xmm2
	divsd	%xmm0, %xmm1
	cvttsd2siq	%xmm1, %rax
	movq	%rax, (%rcx,%r15)
	movl	-4(%r13), %eax
	movq	136(%rsp), %rcx
	leal	2(%rax), %esi
	movq	88(%rsp), %rax
	cvtsi2sdl	%esi, %xmm3
	movsd	(%rax,%r15), %xmm1
	movq	40(%rsp), %rax
	subsd	(%rax,%r15), %xmm1
	movq	104(%rsp), %rax
	divsd	%xmm3, %xmm1
	movsd	%xmm1, (%rax,%r15)
	movq	56(%rsp), %rax
	movsd	.LC7(%rip), %xmm1
	mulsd	(%rax,%r15), %xmm0
	movq	16(%rsp), %rax
	movsd	(%rax,%r15), %xmm4
	movq	112(%rsp), %rax
	movaps	%xmm4, %xmm3
	mulsd	%xmm1, %xmm3
	mulsd	%xmm2, %xmm1
	mulsd	%xmm4, %xmm3
	movaps	%xmm0, %xmm4
	mulsd	%xmm2, %xmm1
	divsd	%xmm3, %xmm4
	divsd	%xmm1, %xmm0
	movsd	%xmm4, (%rax,%r15)
	movq	120(%rsp), %rax
	movsd	%xmm0, (%rax,%r15)
	movl	-4(%r12), %eax
	addl	$2, %eax
	imull	%esi, %eax
	cltq
	movq	%rax, (%rcx,%r15)
	movq	32(%rsp), %rcx
	movq	%r14, (%rcx,%r15)
	addq	%rax, %r14
	addq	$8, %r15
	jmp	.L15
.L45:
	leaq	0(,%r14,8), %rax
	movq	%rax, %rdi
	movq	%rax, 72(%rsp)
	call	malloc@PLT
	movq	72(%rsp), %rsi
	leaq	224(%rsp), %rdi
	movq	%rax, 80(%rsp)
	call	_Z10cudaMallocIdE9cudaErrorPPT_m.isra.0
	movq	(%rsp), %rdx
	movq	40(%rsp), %rsi
	movl	$1, %ecx
	movq	176(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	48(%rsp), %rsi
	movl	$1, %ecx
	movq	184(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	$1, %ecx
	movq	192(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	104(%rsp), %rsi
	movl	$1, %ecx
	movq	200(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	112(%rsp), %rsi
	movl	$1, %ecx
	movq	208(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	120(%rsp), %rsi
	movl	$1, %ecx
	movq	216(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rsi
	movl	$1, %ecx
	movq	232(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	8(%rsp), %rdx
	movl	$1, %ecx
	movq	%rbx, %rsi
	movq	240(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movl	$1, %ecx
	movq	128(%rsp), %rsi
	movq	248(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movl	$1, %ecx
	movq	136(%rsp), %rsi
	movq	256(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	32(%rsp), %rsi
	movl	$1, %ecx
	movq	264(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	172(%rsp), %eax
	salq	$10, %rdx
	movl	$1, %ecx
	movl	$1, 296(%rsp)
	movq	%rdx, 288(%rsp)
	movl	%eax, 276(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 280(%rsp)
	movq	276(%rsp), %rdi
	movl	284(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	264(%rsp)
	.cfi_def_cfa_offset 648
	pushq	248(%rsp)
	.cfi_def_cfa_offset 656
	movq	272(%rsp), %r9
	movq	200(%rsp), %r8
	movq	208(%rsp), %rcx
	movq	216(%rsp), %rdx
	movq	192(%rsp), %rsi
	movq	240(%rsp), %rdi
	call	_Z38__device_stub__Z4initPdS_S_S_S_PlPiS0_PdS_S_S_S_PlPiS0_
	popq	%rcx
	.cfi_def_cfa_offset 648
	popq	%rsi
	.cfi_def_cfa_offset 640
.L18:
	call	cudaDeviceSynchronize@PLT
	movq	72(%rsp), %rdx
	xorl	%ebp, %ebp
	movq	224(%rsp), %rsi
	movq	80(%rsp), %rdi
	movl	$2, %ecx
	movabsq	$8371725894261763689, %r13
	call	cudaMemcpy@PLT
	movq	24(%rsp), %r12
.L19:
	movl	%ebp, %eax
	cmpl	%ebp, 172(%rsp)
	jle	.L46
	movq	32(%rsp), %rcx
	cltd
	leaq	.LC8(%rip), %rsi
	xorl	%r15d, %r15d
	leaq	301(%rsp), %rdi
	movq	%r13, 301(%rsp)
	movl	$7633012, 308(%rsp)
	movq	(%rcx,%rbp,8), %r14
	movl	$10, %ecx
	idivl	%ecx
	addl	$48, %eax
	addl	$48, %edx
	movb	%al, 305(%rsp)
	movb	%dl, 306(%rsp)
	call	fopen@PLT
	movq	%rax, (%rsp)
.L20:
	movq	64(%rsp), %rax
	cvtsi2sdq	%r15, %xmm0
	mulsd	(%rax,%rbp,8), %xmm0
	movl	(%rbx,%rbp,4), %eax
	incl	%eax
	cltq
	cmpq	%r15, %rax
	jle	.L47
	movq	(%rsp), %rdi
	leaq	.LC9(%rip), %rdx
	movb	$1, %al
	incq	%r15
	movl	$2, %esi
	call	__fprintf_chk@PLT
	jmp	.L20
.L47:
	movq	(%rsp), %rdi
	leaq	.LC10(%rip), %rdx
	movb	$1, %al
	xorl	%r15d, %r15d
	movl	$2, %esi
	call	__fprintf_chk@PLT
.L22:
	movq	16(%rsp), %rax
	cvtsi2sdq	%r15, %xmm0
	mulsd	(%rax,%rbp,8), %xmm0
	movl	(%r12), %eax
	incl	%eax
	cltq
	cmpq	%r15, %rax
	jle	.L48
	movq	(%rsp), %rdi
	leaq	.LC9(%rip), %rdx
	movb	$1, %al
	incq	%r15
	movl	$2, %esi
	call	__fprintf_chk@PLT
	jmp	.L22
.L48:
	movq	(%rsp), %rdi
	leaq	.LC10(%rip), %rdx
	movb	$1, %al
	xorl	%r15d, %r15d
	movl	$2, %esi
	call	__fprintf_chk@PLT
	movq	80(%rsp), %rax
	leaq	(%rax,%r14,8), %r14
.L24:
	movl	(%rbx,%rbp,4), %esi
	movl	(%r12), %eax
	leal	2(%rsi), %edi
	addl	$2, %eax
	imull	%edi, %eax
	cltq
	cmpq	%r15, %rax
	jle	.L49
	movq	%r15, %rax
	movslq	%edi, %rdi
	incl	%esi
	movsd	(%r14,%r15,8), %xmm0
	cqto
	movslq	%esi, %rsi
	idivq	%rdi
	cmpq	%rsi, %rdx
	jne	.L25
	leaq	.LC10(%rip), %rdx
	jmp	.L43
.L25:
	leaq	.LC9(%rip), %rdx
.L43:
	movq	(%rsp), %rdi
	movl	$2, %esi
	movb	$1, %al
	incq	%r15
	call	__fprintf_chk@PLT
	jmp	.L24
.L49:
	movq	(%rsp), %rdi
	incq	%rbp
	addq	$4, %r12
	call	fclose@PLT
	jmp	.L19
.L46:
	call	clock@PLT
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	movl	$49152, %r8d
	movq	%rax, %rbp
	salq	$10, %rdx
	movl	172(%rsp), %eax
	movl	$1, %ecx
	movq	%rdx, 288(%rsp)
	movl	%eax, 276(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 280(%rsp)
	movq	276(%rsp), %rdi
	movl	284(%rsp), %esi
	movl	$1, 296(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L29
	pushq	264(%rsp)
	.cfi_def_cfa_offset 648
	pushq	248(%rsp)
	.cfi_def_cfa_offset 656
	movq	248(%rsp), %r9
	movq	272(%rsp), %r8
	movq	264(%rsp), %rcx
	movq	232(%rsp), %rdx
	movq	224(%rsp), %rsi
	movq	240(%rsp), %rdi
	call	_Z39__device_stub__Z3runPdS_S_PlS0_PiS1_S0_PdS_S_PlS0_PiS1_S0_
	popq	%rax
	.cfi_def_cfa_offset 648
	popq	%rdx
	.cfi_def_cfa_offset 640
.L29:
	call	cudaDeviceSynchronize@PLT
	movabsq	$8679613338825745778, %r12
	call	clock@PLT
	movq	72(%rsp), %rdx
	movq	80(%rsp), %rdi
	movl	$2, %ecx
	movq	224(%rsp), %rsi
	subq	%rbp, %rax
	xorl	%ebp, %ebp
	cvtsi2sdq	%rax, %xmm0
	divsd	.LC11(%rip), %xmm0
	movsd	%xmm0, (%rsp)
	call	cudaMemcpy@PLT
.L30:
	movl	%ebp, %eax
	cmpl	%ebp, 172(%rsp)
	jle	.L50
	movq	32(%rsp), %rcx
	cltd
	leaq	.LC8(%rip), %rsi
	xorl	%r14d, %r14d
	leaq	301(%rsp), %rdi
	movq	%r12, 301(%rsp)
	movw	$116, 309(%rsp)
	movq	(%rcx,%rbp,8), %r13
	movl	$10, %ecx
	idivl	%ecx
	addl	$48, %eax
	addl	$48, %edx
	movb	%al, 304(%rsp)
	movb	%dl, 305(%rsp)
	call	fopen@PLT
	movq	%rax, %r15
.L31:
	movq	64(%rsp), %rax
	cvtsi2sdq	%r14, %xmm0
	mulsd	(%rax,%rbp,8), %xmm0
	movl	(%rbx,%rbp,4), %eax
	incl	%eax
	cltq
	cmpq	%r14, %rax
	jle	.L51
	leaq	.LC9(%rip), %rdx
	movq	%r15, %rdi
	movb	$1, %al
	incq	%r14
	movl	$2, %esi
	call	__fprintf_chk@PLT
	jmp	.L31
.L51:
	leaq	.LC10(%rip), %rdx
	movq	%r15, %rdi
	movb	$1, %al
	xorl	%r14d, %r14d
	movl	$2, %esi
	call	__fprintf_chk@PLT
.L33:
	movq	16(%rsp), %rax
	cvtsi2sdq	%r14, %xmm0
	mulsd	(%rax,%rbp,8), %xmm0
	movq	144(%rsp), %rax
	movl	(%rax), %eax
	movl	%eax, 8(%rsp)
	incl	%eax
	cltq
	cmpq	%r14, %rax
	jle	.L52
	leaq	.LC9(%rip), %rdx
	movq	%r15, %rdi
	movb	$1, %al
	incq	%r14
	movl	$2, %esi
	call	__fprintf_chk@PLT
	jmp	.L33
.L52:
	leaq	.LC10(%rip), %rdx
	movq	%r15, %rdi
	movb	$1, %al
	xorl	%r14d, %r14d
	movl	$2, %esi
	call	__fprintf_chk@PLT
	movq	80(%rsp), %rax
	leaq	(%rax,%r13,8), %r13
.L35:
	movq	144(%rsp), %rax
	movl	(%rbx,%rbp,4), %esi
	movl	(%rax), %eax
	leal	2(%rsi), %edi
	movl	%eax, 8(%rsp)
	addl	$2, %eax
	imull	%edi, %eax
	cltq
	cmpq	%r14, %rax
	jle	.L53
	movq	%r14, %rax
	movslq	%edi, %rdi
	incl	%esi
	movsd	0(%r13,%r14,8), %xmm0
	cqto
	movslq	%esi, %rsi
	idivq	%rdi
	cmpq	%rsi, %rdx
	jne	.L36
	leaq	.LC10(%rip), %rdx
	jmp	.L44
.L36:
	leaq	.LC9(%rip), %rdx
.L44:
	movl	$2, %esi
	movq	%r15, %rdi
	movb	$1, %al
	incq	%r14
	call	__fprintf_chk@PLT
	jmp	.L35
.L53:
	movq	%r15, %rdi
	incq	%rbp
	call	fclose@PLT
	addq	$4, 144(%rsp)
	jmp	.L30
.L50:
	movq	40(%rsp), %rdi
	call	free@PLT
	movq	88(%rsp), %rdi
	call	free@PLT
	movq	48(%rsp), %rdi
	call	free@PLT
	movq	56(%rsp), %rdi
	call	free@PLT
	movq	152(%rsp), %rdi
	call	free@PLT
	movq	16(%rsp), %rdi
	call	free@PLT
	movq	64(%rsp), %rdi
	call	free@PLT
	movq	96(%rsp), %rdi
	call	free@PLT
	movq	104(%rsp), %rdi
	call	free@PLT
	movq	112(%rsp), %rdi
	call	free@PLT
	movq	120(%rsp), %rdi
	call	free@PLT
	movq	24(%rsp), %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	128(%rsp), %rdi
	call	free@PLT
	movq	136(%rsp), %rdi
	call	free@PLT
	movq	32(%rsp), %rdi
	call	free@PLT
	movq	176(%rsp), %rdi
	call	cudaFree@PLT
	movq	184(%rsp), %rdi
	call	cudaFree@PLT
	movq	192(%rsp), %rdi
	call	cudaFree@PLT
	movq	200(%rsp), %rdi
	call	cudaFree@PLT
	movq	208(%rsp), %rdi
	call	cudaFree@PLT
	movq	216(%rsp), %rdi
	call	cudaFree@PLT
	movq	232(%rsp), %rdi
	call	cudaFree@PLT
	movq	240(%rsp), %rdi
	call	cudaFree@PLT
	movq	248(%rsp), %rdi
	call	cudaFree@PLT
	movq	256(%rsp), %rdi
	call	cudaFree@PLT
	movq	264(%rsp), %rdi
	call	cudaFree@PLT
	movsd	(%rsp), %xmm0
	leaq	.LC12(%rip), %rsi
	movb	$1, %al
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	568(%rsp), %rax
	subq	%fs:40, %rax
	je	.L40
	call	__stack_chk_fail@PLT
.L40:
	addq	$584, %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
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC13:
	.string	"_Z3runPdS_S_PlS0_PiS1_S0_"
.LC14:
	.string	"_Z4initPdS_S_S_S_PlPiS0_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2057:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC13(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z3runPdS_S_PlS0_PiS1_S0_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC14(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z4initPdS_S_S_S_PlPiS0_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2057:
	.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	0
	.long	1069547520
	.align 8
.LC4:
	.long	0
	.long	1074790400
	.align 8
.LC5:
	.long	1413754136
	.long	1074340347
	.align 8
.LC6:
	.long	-1698910392
	.long	1048238066
	.align 8
.LC7:
	.long	261748427
	.long	1052054839
	.align 8
.LC11:
	.long	0
	.long	1093567616
	.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 <time.h>

#define PI 3.1415926535897932384
#define mu0 4*PI*1e-7
//Threads per block is capped at 1024 for hardware reasons
//In some cases using a smaller number of threads per block will be more efficient
#define threadsPerBlock 1024
//Max grid points is to defined in order to allocate shared memory for each block
#define MaxGridPoints 6144


/*
NOTE:
grid is technically a 3-dimensional variable length array. In order to copy
this 3d array, it needed to be compressed into a single dimension.
Rather than decompressing this and wasting valuable resources all computation is done in line.
i.e. all blocks will write to adjacent memory
the position in memory is determined by three factors: the block numnber, the r
segment value and the z segment value
Each block is given an initial offset. This offset gives the block a unique space inside
of the overall shared array in which to write its data. Each block's grid is
r driven meaning grid(r,z) = grid[r*zMax + z], thus the position of each point in the
overall array is given by grids[offset + r*zMax + z]
*/


//This function will be called to initialize all data in the grid, Each thread will
//calculate the initial condion for a variable number of points and each block will have a different
//set of initial conditions all passed to device memory already
__global__ void init(double *grids, double *Ils, double *dIs, double *ldrs, double *rlengths, long int *grid_sizes, int *zsegs, long int *offsets){
  // get needed data from memory, all memory is block adjacent
  //block index determines which problem you are working on

  //Current on the left side of the grid
  double Il = Ils[blockIdx.x];
  //amount the current increments per segment in the z direction
  //This model assumes linear change
  double dI = dIs[blockIdx.x];
  //Physical distance change between r segments
  double ldr = ldrs[blockIdx.x];
  //total r length
  double rlength = rlengths[blockIdx.x];
  //Total number of grid points
  long int grid_size = grid_sizes[blockIdx.x];
  //number of z segmnents
  int zseg = zsegs[blockIdx.x];
  //Offset of the current grid
  //Because data needs to be copied to the gpu, it is compressed into one dimension
  //thus all grid data lives in the same array and the offset for that grid determines the satrt of this grid
  //This is explained in more detail above
  long int gridStart = offsets[blockIdx.x];

  //Find how many points per thread will be executing

  //remainder tells you how many threads will have one extra point to cover
  int rem = grid_size%threadsPerBlock;
  //division tells you how many points ALL theads will cover
  int divi = grid_size/threadsPerBlock;
  long int start, fin;
  //If threadIdx.x is less than rem the current thread will recieve one extra point to calculate
  if(threadIdx.x<rem){
    start = threadIdx.x*(divi+1);
    fin = start + divi + 1;
  }
  else{
    //rem is added here because at this point a "rem" number of threads will have one extra point to compute
    //overall this is "rem" extra points and this is accounted for by adding rem
    start = threadIdx.x*divi + rem;
    fin = start + divi;
  }

  //get initial conditions for all points per thread

  //Loops over each point a thread takes care of
  //NOTE: If a thread is not supposed to execute anything it will have the same
  //value for "start" and "fin" and will not enter the loop
  for(int i = start; i<fin; i++){
    //Gets current r and z segmenty values
    //Because grid is "r-driven", dividing by zMax = the r value and moding by
    //zMax gves the current z value
    //For more info check out the "note" at the top
    long int r_seg_val = i/(zseg+2);
    long int z_seg_val = i%(zseg+2);
    //Sets initial condition at the given point based on an equation given by Dr. Sankaran
    grids[i + gridStart] = (1-(r_seg_val*r_seg_val*ldr*ldr/(3*rlength*rlength)))*3*mu0*(Il + z_seg_val*dI)*r_seg_val*ldr/(4*PI*rlength*rlength);
  }
}

//This function funs until the number of timesteps needed is completed
//again this function passes data by memory copy in the form of arrays
//each block excercises a different set of input params
__global__ void run(double *grids, double *r_augs, double *z_augs, long int *allMaxSteps, long int *grid_sizes, int *rsegs, int *zsegs, long int *offsets){
  //get data from initialization arrays
  //block index determines the problem that is being worked on

  //r_aug and z_aug are constants that multiple values need to update to 1 timestep ahead
  //r_aug is used for changes in r and z_aug is used for changes in z
  double r_aug = r_augs[blockIdx.x];
  double z_aug = z_augs[blockIdx.x];
  //max steps is the number of steps needed to be completed for the simulation
  //to be finished
  long int maxSteps = allMaxSteps[blockIdx.x];
  //total grid size for this block's problem
  long int grid_size = grid_sizes[blockIdx.x];
  //number of segments in r and z directions
  int rseg = rsegs[blockIdx.x];
  int zseg = zsegs[blockIdx.x];
  //offset for block's memory locations
  //More info at top
  long int gridStart = offsets[blockIdx.x];

  //find how many points each thread will execute for

  //rem represents the number of threads that will have one more than divi points
  //to take care of
  int rem = grid_size%threadsPerBlock;
  //divi is the amount of points each thread will have to take care of at a minimum
  int divi = grid_size/threadsPerBlock;
  int start, fin;
  if(threadIdx.x<rem){
    start = threadIdx.x*(divi+1);
    fin = start + divi + 1;
  }
  else{
    //rem is added here because at this point a "rem" number of threads will have one extra point to compute
    //overall this is "rem" extra points and this is accounted for by adding rem
    start = threadIdx.x*divi + rem;
    fin = start + divi;
  }
  long int steps = 0;

  //This is shared memory
  //It is quick access but there is a limited amount and can only be
  //accessed inside each block
  //the extern keyword states that the size of this memory is passed in at the
  //function call. This allows a variable amount of shared memory to be used
  extern __shared__ double grid_new_s[];
  extern __shared__ double grid_old_s[];

  //Initialize shared memory
  //gridStart is the offset for this block
  for(int i = start; i<fin; i++){
    grid_new_s[i] = grids[i + gridStart];
  }
  //wait for all threads
  __syncthreads();

  while(steps<maxSteps){
    //copy all points from new to old
    for(int i = start; i<fin; i++){
      grid_old_s[i] = grid_new_s[i];
    }
    //wait for all threads
    __syncthreads();
    for(int i = start; i<fin; i++){
      //because the grid is row driven
      //i / zMax = rvalue
      //i % zMax = zvalue
      int r = i/(zseg+2);
      int z = i%(zseg+2);

      //leave boundary conditions alone as specified by the problem
      if(z != 0 && z != zseg+1){
        if(r!= 0 && r!= rseg+1){
          //when r = 1 a phontom point needs to be used as the slope change from
          //0 -> 1 is too large
          //in this case the B field at -1 is assumed to be of the same magnitude as at 1
          //but in the opposite direction
          //this yields the folloeing equation
          if(r==1){
            grid_new_s[i] += r_aug*(2*grid_old_s[i+(zseg+2)] - 4*grid_old_s[i]) +
              z_aug * (grid_old_s[i+1] - 2*grid_old_s[i] + grid_old_s[i-1]);
          }
          //normal update function
          else{
            grid_new_s[i] += r_aug*((1+(1/(2*r)))*grid_old_s[i+(zseg+2)] + (-2-(1/(r*r)))*grid_old_s[i] + (1-(1/(2*r)))*grid_old_s[i-(zseg+2)])
              +z_aug*(grid_old_s[i+1] - 2*grid_old_s[i] + grid_old_s[i-1]);
          }
        }
      }
    }
    steps++;
    //wait for all threads
    __syncthreads();
  }
  //copy all values back to main memory
  for(int i = start; i<fin; i++){
    grids[i + gridStart] = grid_new_s[i];
  }
}

int main(){
  //string label in input FILE
  //no use just for clarity
  char label[256];

  double *Il_h, *Ir_h, *rlength_h, *eta_h, *tstep_h, *ldr_h, *ldz_h, *zlength_h, *dI_h, *r_aug_h, *z_aug_h, *grids_h;
  int *rseg_h, *zseg_h;
  long int *totsteps_h, *grid_size_h, *offsets_h;

  double *Il_d, *rlength_d, *ldr_d, *dI_d, *r_aug_d, *z_aug_d, *grids_d;
  int *rseg_d, *zseg_d;
  long int *totsteps_d, *grid_size_d, *offsets_d;

  int testcases;
  //printf("How many test cases? ");
  //gets number of test cases
  scanf("%d", &testcases);

  //these sizes are used for malloc ops
  //the size is based on the amount of test cases and the
  //type of data this array holds
  size_t doubleSize = testcases*sizeof(double);
  size_t intSize = testcases*sizeof(int);
  size_t longIntSize = testcases*sizeof(long int);

  //host memory allocation
  Il_h = (double*)malloc(doubleSize);
  Ir_h = (double*)malloc(doubleSize);
  rlength_h = (double*)malloc(doubleSize);
  eta_h = (double*)malloc(doubleSize);
  tstep_h = (double*)malloc(doubleSize);
  ldr_h = (double*)malloc(doubleSize);
  ldz_h = (double*)malloc(doubleSize);
  zlength_h = (double*)malloc(doubleSize);
  dI_h = (double*)malloc(doubleSize);
  r_aug_h = (double*)malloc(doubleSize);
  z_aug_h = (double*)malloc(doubleSize);
  rseg_h = (int*)malloc(intSize);
  zseg_h = (int*)malloc(intSize);
  totsteps_h = (long int*)malloc(longIntSize);
  grid_size_h = (long int*)malloc(longIntSize);
  offsets_h = (long int*)malloc(longIntSize);

  //device memory allocation
  hipMalloc(&Il_d, doubleSize);
  hipMalloc(&rlength_d, doubleSize);
  hipMalloc(&ldr_d, doubleSize);
  hipMalloc(&dI_d, doubleSize);
  hipMalloc(&r_aug_d, doubleSize);
  hipMalloc(&z_aug_d, doubleSize);
  hipMalloc(&rseg_d, intSize);
  hipMalloc(&zseg_d, intSize);
  hipMalloc(&totsteps_d, longIntSize);
  hipMalloc(&grid_size_d, longIntSize);
  hipMalloc(&offsets_d, longIntSize);

  long int total_grid_size = 0;
  for(int counter = 0; counter < testcases; counter++){
    //label is unused
    scanf("%s", &label);
    //value of current when z = 0
    //printf("What is your left I? ");
    scanf("%lf", Il_h + counter);
    //value of current when z = zMax
    //printf("What is your right I? ");
    scanf("%lf", Ir_h + counter);
    //total r value
    //printf("What is the radius of your rod? ");
    scanf("%lf", rlength_h + counter);
    //total z value
    //printf("What is the length of your rod? ");
    scanf("%lf", zlength_h + counter);
    //value of eta (proportional to diffusivity)
    //printf("What is eta? ");
    scanf("%lf", eta_h + counter);
    //number of r segments
    //printf("How many segments would you like per radius? ");
    scanf("%d", rseg_h + counter);
    //number of z segments
    //printf("How many segments would you like per length? ");
    scanf("%d", zseg_h + counter);
    //length of each r segment
    ldr_h[counter] = rlength_h[counter]/(rseg_h[counter]+1);
    //length of each z segment
    ldz_h[counter] = zlength_h[counter]/(zseg_h[counter]+1);
    double smallest = ldr_h[counter];
    if(ldz_h[counter] < smallest)
      smallest = ldz_h[counter];
    //determines tstep that ensures stability
    //0.125 is derived from 2^-(num dimensions + 1)
    tstep_h[counter] = 0.125*smallest*smallest*mu0/eta_h[counter];
    //gets total run time
    double tottime;
    //printf("How long would you like to run? ");
    scanf("%lf", &tottime);
    //total steps is an integer that truncates remainder
    totsteps_h[counter] = tottime/tstep_h[counter];
    //dI is change in I per each z segment
    //zseg + 2 to account for boundary conditions
    dI_h[counter] = (Ir_h[counter] - Il_h[counter]) / (zseg_h[counter]+2);
    //r and z aug are used for updating values for increasing time
    //as per numerical update
    r_aug_h[counter] = eta_h[counter]*tstep_h[counter]/(mu0*ldr_h[counter]*ldr_h[counter]);
    z_aug_h[counter] = eta_h[counter]*tstep_h[counter]/(mu0*ldz_h[counter]*ldz_h[counter]);
    //grid size accounts for boundary conditions
    grid_size_h[counter] = (rseg_h[counter] + 2)*(zseg_h[counter] + 2);
    //offsets says where in memory this problems data starts
    //more info above
    offsets_h[counter] = total_grid_size;
    //Total grid size is the size of all needed grid memory
    total_grid_size += grid_size_h[counter];
  }


  size_t gridsSize = total_grid_size*sizeof(double);

  //allocate variable amount of grid memory on both host and device
  grids_h = (double*)malloc(gridsSize);
  hipMalloc(&grids_d, gridsSize);

  //copy memory down to device
  hipMemcpy(Il_d, Il_h, doubleSize, hipMemcpyHostToDevice);
  hipMemcpy(rlength_d, rlength_h, doubleSize, hipMemcpyHostToDevice);
  hipMemcpy(ldr_d, ldr_h, doubleSize, hipMemcpyHostToDevice);
  hipMemcpy(dI_d, dI_h, doubleSize, hipMemcpyHostToDevice);
  hipMemcpy(r_aug_d, r_aug_h, doubleSize, hipMemcpyHostToDevice);
  hipMemcpy(z_aug_d, z_aug_h, doubleSize, hipMemcpyHostToDevice);
  hipMemcpy(rseg_d, rseg_h, intSize, hipMemcpyHostToDevice);
  hipMemcpy(zseg_d, zseg_h, intSize, hipMemcpyHostToDevice);
  hipMemcpy(totsteps_d, totsteps_h, longIntSize, hipMemcpyHostToDevice);
  hipMemcpy(grid_size_d, grid_size_h, longIntSize, hipMemcpyHostToDevice);
  hipMemcpy(offsets_d, offsets_h, longIntSize, hipMemcpyHostToDevice);

  //initialize all problems with initial values as per Dr. Sankaran's instructions
  init<<<testcases,threadsPerBlock>>>(grids_d, Il_d, dI_d, ldr_d, rlength_d, grid_size_d, zseg_d, offsets_d);
  //wait for all device implementation to finish
  hipDeviceSynchronize();

  //copy initial grid data back to host
  hipMemcpy(grids_h, grids_d, gridsSize, hipMemcpyDeviceToHost);


  //outputs all problems initial conditions
  FILE *myfile;
  long int i;
  for(int counter = 0; counter < testcases; counter++){
    long int gridStart = offsets_h[counter];
    char init[] = "initXX.txt";
    int tens = counter / 10;
    int ones = counter % 10;
    init[4] = tens + '0';
    init[5] = ones + '0';
    myfile = fopen(init, "w");
    for(i = 0; i< zseg_h[counter]+1; i++)
      fprintf(myfile, "%lf ", i*ldz_h[counter]);
    fprintf(myfile, "%lf\n", i*ldz_h[counter]);

    for(i = 0; i< rseg_h[counter]+1; i++)
      fprintf(myfile, "%lf ", i*ldr_h[counter]);
    fprintf(myfile, "%lf\n", i*ldr_h[counter]);

    for(i = 0; i< (rseg_h[counter] + 2)*(zseg_h[counter]+2); i++){
      if(i%(zseg_h[counter]+2)==zseg_h[counter]+1)
        fprintf(myfile, "%lf\n", grids_h[i + gridStart]);
      else
        fprintf(myfile, "%lf ", grids_h[i + gridStart]);
    }

    fclose(myfile);
  }



  clock_t begin, end;
  double time_spent;
  begin = clock();

  //run
  //size capped at MaxGridPoints, grid points per problem
  run<<<testcases,threadsPerBlock, MaxGridPoints*sizeof(double)>>>(grids_d, r_aug_d, z_aug_d, totsteps_d, grid_size_d, rseg_d, zseg_d, offsets_d);
  //wait for all device implementation to finish
  hipDeviceSynchronize();

  end = clock();
  time_spent = (double)(end - begin) / CLOCKS_PER_SEC;

  //copy final grid data back to host
  hipMemcpy(grids_h, grids_d, gridsSize, hipMemcpyDeviceToHost);

  //output all problems final grid data
  for(int counter = 0; counter < testcases; counter++){
    long int gridStart = offsets_h[counter];
    char init[] = "resXX.txt";
    int tens = counter / 10;
    int ones = counter % 10;
    init[3] = tens + '0';
    init[4] = ones + '0';
    myfile = fopen(init, "w");
    for(i = 0; i< zseg_h[counter]+1; i++)
      fprintf(myfile, "%lf ", i*ldz_h[counter]);
    fprintf(myfile, "%lf\n", i*ldz_h[counter]);

    for(i = 0; i< rseg_h[counter]+1; i++)
      fprintf(myfile, "%lf ", i*ldr_h[counter]);
    fprintf(myfile, "%lf\n", i*ldr_h[counter]);

    for(i = 0; i< (rseg_h[counter] + 2)*(zseg_h[counter]+2); i++){
      if(i%(zseg_h[counter]+2)==zseg_h[counter]+1)
        fprintf(myfile, "%lf\n", grids_h[i + gridStart]);
      else
        fprintf(myfile, "%lf ", grids_h[i + gridStart]);
    }

    fclose(myfile);
  }

  //free host memory
  free(Il_h);
  free(Ir_h);
  free(rlength_h);
  free(eta_h);
  free(tstep_h);
  free(ldr_h);
  free(ldz_h);
  free(zlength_h);
  free(dI_h);
  free(r_aug_h);
  free(z_aug_h);
  free(rseg_h);
  free(zseg_h);
  free(totsteps_h);
  free(grid_size_h);
  free(offsets_h);

  //free device memory
  hipFree(Il_d);
  hipFree(rlength_d);
  hipFree(ldr_d);
  hipFree(dI_d);
  hipFree(r_aug_d);
  hipFree(z_aug_d);
  hipFree(rseg_d);
  hipFree(zseg_d);
  hipFree(totsteps_d);
  hipFree(grid_size_d);
  hipFree(offsets_d);



  printf("\n------------------------------------\nExecution took: %lf sec\n", time_spent);
  return 0;
}

	.text
	.file	"run.hip"
	.globl	_Z19__device_stub__initPdS_S_S_S_PlPiS0_ # -- Begin function _Z19__device_stub__initPdS_S_S_S_PlPiS0_
	.type	_Z19__device_stub__initPdS_S_S_S_PlPiS0_,@function
_Z19__device_stub__initPdS_S_S_S_PlPiS0_: # @_Z19__device_stub__initPdS_S_S_S_PlPiS0_
	.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	$160, %rsp
	.cfi_def_cfa_offset 208
	.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	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	216(%rsp), %rax
	movq	%rax, 56(%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	$_Z4initPdS_S_S_S_PlPiS0_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z19__device_stub__initPdS_S_S_S_PlPiS0_, .Lfunc_end0-_Z19__device_stub__initPdS_S_S_S_PlPiS0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z18__device_stub__runPdS_S_PlS0_PiS1_S0_ # -- Begin function _Z18__device_stub__runPdS_S_PlS0_PiS1_S0_
	.type	_Z18__device_stub__runPdS_S_PlS0_PiS1_S0_,@function
_Z18__device_stub__runPdS_S_PlS0_PiS1_S0_: # @_Z18__device_stub__runPdS_S_PlS0_PiS1_S0_
	.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	$160, %rsp
	.cfi_def_cfa_offset 208
	.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	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	216(%rsp), %rax
	movq	%rax, 56(%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	$_Z3runPdS_S_PlS0_PiS1_S0_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	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	_Z18__device_stub__runPdS_S_PlS0_PiS1_S0_, .Lfunc_end1-_Z18__device_stub__runPdS_S_PlS0_PiS1_S0_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI2_0:
	.quad	0x3fc0000000000000              # double 0.125
.LCPI2_1:
	.quad	0x4010000000000000              # double 4
.LCPI2_2:
	.quad	0x400921fb54442d18              # double 3.1415926535897931
.LCPI2_3:
	.quad	0x3e7ad7f29abcaf48              # double 9.9999999999999995E-8
.LCPI2_4:
	.quad	0x3eb515370f99f6cb              # double 1.2566370614359173E-6
.LCPI2_5:
	.quad	0x412e848000000000              # double 1.0E+6
	.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	$536, %rsp                      # imm = 0x218
	.cfi_def_cfa_offset 592
	.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	4(%rsp), %rbx
	movl	$.L.str, %edi
	movq	%rbx, %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movslq	(%rbx), %r15
	leaq	(,%r15,8), %rbp
	shlq	$2, %r15
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 80(%rsp)                  # 8-byte Spill
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 200(%rsp)                 # 8-byte Spill
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, %r12
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 192(%rsp)                 # 8-byte Spill
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, %r13
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 88(%rsp)                  # 8-byte Spill
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 184(%rsp)                 # 8-byte Spill
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 72(%rsp)                  # 8-byte Spill
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 168(%rsp)                 # 8-byte Spill
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 160(%rsp)                 # 8-byte Spill
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %r14
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %rbx
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 152(%rsp)                 # 8-byte Spill
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 144(%rsp)                 # 8-byte Spill
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 56(%rsp)                  # 8-byte Spill
	leaq	264(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	256(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	248(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	240(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	232(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	224(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	216(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	112(%rsp), %rdi
	movq	%r15, 128(%rsp)                 # 8-byte Spill
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	208(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	104(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	96(%rsp), %rdi
	movq	%rbp, 136(%rsp)                 # 8-byte Spill
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	4(%rsp), %rax
	cmpl	$0, (%rax)
	movq	%r12, 64(%rsp)                  # 8-byte Spill
	movq	%r13, 24(%rsp)                  # 8-byte Spill
	jle	.LBB2_1
# %bb.9:                                # %.lr.ph.preheader
	movq	%r14, %rbp
	movq	%rbx, %r13
	xorl	%r12d, %r12d
	xorl	%ecx, %ecx
	xorl	%edi, %edi
.LBB2_10:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	%rcx, 32(%rsp)                  # 8-byte Spill
	movq	%rdi, 48(%rsp)                  # 8-byte Spill
	movl	$.L.str.1, %edi
	leaq	272(%rsp), %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movq	80(%rsp), %rax                  # 8-byte Reload
	leaq	(%rax,%r12), %rsi
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movq	200(%rsp), %rax                 # 8-byte Reload
	leaq	(%rax,%r12), %rsi
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movq	64(%rsp), %rax                  # 8-byte Reload
	leaq	(%rax,%r12), %rsi
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movq	184(%rsp), %r15                 # 8-byte Reload
	leaq	(%r15,%r12), %rsi
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movq	192(%rsp), %rax                 # 8-byte Reload
	leaq	(%rax,%r12), %rsi
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movl	$.L.str, %edi
	movq	%rbp, %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movl	$.L.str, %edi
	movq	%r13, %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movl	(%rbp), %eax
	incl	%eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	movq	64(%rsp), %rax                  # 8-byte Reload
	movsd	(%rax,%r12), %xmm1              # xmm1 = mem[0],zero
	divsd	%xmm0, %xmm1
	movq	24(%rsp), %rax                  # 8-byte Reload
	movsd	%xmm1, (%rax,%r12)
	movsd	(%r15,%r12), %xmm2              # xmm2 = mem[0],zero
	movl	(%r13), %eax
	incl	%eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	divsd	%xmm0, %xmm2
	movsd	%xmm2, 40(%rsp)                 # 8-byte Spill
	movq	88(%rsp), %rax                  # 8-byte Reload
	movsd	%xmm2, (%rax,%r12)
	movapd	%xmm2, %xmm0
	minsd	%xmm1, %xmm0
	movapd	%xmm0, %xmm1
	mulsd	.LCPI2_0(%rip), %xmm1
	mulsd	%xmm0, %xmm1
	mulsd	.LCPI2_1(%rip), %xmm1
	mulsd	.LCPI2_2(%rip), %xmm1
	mulsd	.LCPI2_3(%rip), %xmm1
	movq	192(%rsp), %r15                 # 8-byte Reload
	divsd	(%r15,%r12), %xmm1
	movsd	%xmm1, 176(%rsp)                # 8-byte Spill
	movl	$.L.str.2, %edi
	leaq	8(%rsp), %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movsd	8(%rsp), %xmm0                  # xmm0 = mem[0],zero
	movsd	176(%rsp), %xmm3                # 8-byte Reload
                                        # xmm3 = mem[0],zero
	divsd	%xmm3, %xmm0
	cvttsd2si	%xmm0, %rax
	movslq	(%r13), %rcx
	addq	$2, %rcx
	xorps	%xmm0, %xmm0
	cvtsi2sd	%ecx, %xmm0
	movq	152(%rsp), %rdx                 # 8-byte Reload
	movq	%rax, (%rdx,%r12)
	movslq	(%rbp), %rdi
	addq	$2, %rdi
	imulq	%rcx, %rdi
	movq	32(%rsp), %rcx                  # 8-byte Reload
	movq	200(%rsp), %rax                 # 8-byte Reload
	movsd	(%rax,%r12), %xmm1              # xmm1 = mem[0],zero
	movq	80(%rsp), %rax                  # 8-byte Reload
	subsd	(%rax,%r12), %xmm1
	divsd	%xmm0, %xmm1
	movq	72(%rsp), %rax                  # 8-byte Reload
	movsd	%xmm1, (%rax,%r12)
	mulsd	(%r15,%r12), %xmm3
	movq	24(%rsp), %rax                  # 8-byte Reload
	movsd	(%rax,%r12), %xmm0              # xmm0 = mem[0],zero
	movapd	%xmm0, %xmm1
	movsd	.LCPI2_4(%rip), %xmm2           # xmm2 = mem[0],zero
	mulsd	%xmm2, %xmm1
	mulsd	%xmm0, %xmm1
	movapd	%xmm3, %xmm0
	divsd	%xmm1, %xmm0
	movq	168(%rsp), %rax                 # 8-byte Reload
	movsd	%xmm0, (%rax,%r12)
	movsd	40(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	movapd	%xmm1, %xmm0
	mulsd	%xmm2, %xmm0
	mulsd	%xmm1, %xmm0
	movapd	%xmm3, %xmm1
	divsd	%xmm0, %xmm1
	movq	160(%rsp), %rax                 # 8-byte Reload
	movsd	%xmm1, (%rax,%r12)
	movq	144(%rsp), %rax                 # 8-byte Reload
	movq	%rdi, (%rax,%r12)
	movq	56(%rsp), %rax                  # 8-byte Reload
	movq	48(%rsp), %rdx                  # 8-byte Reload
	movq	%rdx, (%rax,%r12)
	addq	%rdx, %rdi
	incq	%rcx
	movslq	4(%rsp), %rax
	addq	$8, %r12
	addq	$4, %r13
	addq	$4, %rbp
	cmpq	%rax, %rcx
	jl	.LBB2_10
# %bb.2:                                # %._crit_edge.loopexit
	shlq	$3, %rdi
	movq	64(%rsp), %r12                  # 8-byte Reload
	movq	24(%rsp), %r13                  # 8-byte Reload
	jmp	.LBB2_3
.LBB2_1:
	xorl	%edi, %edi
.LBB2_3:                                # %._crit_edge
	movq	%rdi, %r15
	callq	malloc
	movq	%rax, 40(%rsp)                  # 8-byte Spill
	leaq	120(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movq	264(%rsp), %rdi
	movq	80(%rsp), %rsi                  # 8-byte Reload
	movq	136(%rsp), %rbp                 # 8-byte Reload
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	256(%rsp), %rdi
	movq	%r12, %rsi
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	248(%rsp), %rdi
	movq	%r13, %rsi
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	240(%rsp), %rdi
	movq	72(%rsp), %rsi                  # 8-byte Reload
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	232(%rsp), %rdi
	movq	168(%rsp), %rsi                 # 8-byte Reload
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	224(%rsp), %rdi
	movq	160(%rsp), %rsi                 # 8-byte Reload
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	216(%rsp), %rdi
	movq	%r14, %rsi
	movq	128(%rsp), %r12                 # 8-byte Reload
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	112(%rsp), %rdi
	movq	%rbx, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	208(%rsp), %rdi
	movq	152(%rsp), %rsi                 # 8-byte Reload
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	104(%rsp), %rdi
	movq	144(%rsp), %rsi                 # 8-byte Reload
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	96(%rsp), %rdi
	movq	56(%rsp), %rsi                  # 8-byte Reload
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	4(%rsp), %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rax               # imm = 0x100000000
	leaq	1024(%rax), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_5
# %bb.4:
	movq	120(%rsp), %rdi
	movq	264(%rsp), %rsi
	movq	240(%rsp), %rdx
	movq	248(%rsp), %rcx
	movq	256(%rsp), %r8
	movq	104(%rsp), %r9
	pushq	96(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	120(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z19__device_stub__initPdS_S_S_S_PlPiS0_
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
.LBB2_5:
	callq	hipDeviceSynchronize
	movq	120(%rsp), %rsi
	movq	40(%rsp), %rdi                  # 8-byte Reload
	movq	%r15, 176(%rsp)                 # 8-byte Spill
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	cmpl	$0, 4(%rsp)
	movq	88(%rsp), %r13                  # 8-byte Reload
	jle	.LBB2_23
# %bb.6:                                # %.lr.ph309
	xorl	%r12d, %r12d
.LBB2_7:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_12 Depth 2
                                        #     Child Loop BB2_16 Depth 2
                                        #     Child Loop BB2_19 Depth 2
	movq	56(%rsp), %rax                  # 8-byte Reload
	movq	(%rax,%r12,8), %rax
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movl	$1953066601, 8(%rsp)            # imm = 0x74696E69
	movl	$1954051118, 14(%rsp)           # imm = 0x7478742E
	movb	$0, 18(%rsp)
	movl	%r12d, %eax
	movl	$3435973837, %ecx               # imm = 0xCCCCCCCD
	imulq	%rcx, %rax
	shrq	$35, %rax
	leal	(%rax,%rax), %ecx
	leal	(%rcx,%rcx,4), %ecx
	movl	%r12d, %edx
	subl	%ecx, %edx
	addb	$48, %al
	movb	%al, 12(%rsp)
	orb	$48, %dl
	movb	%dl, 13(%rsp)
	movl	$.L.str.3, %esi
	leaq	8(%rsp), %rdi
	callq	fopen
	movq	%rax, %r15
	cmpl	$0, (%rbx,%r12,4)
	movsd	(%r13,%r12,8), %xmm1            # xmm1 = mem[0],zero
	js	.LBB2_8
# %bb.11:                               # %.lr.ph293.preheader
                                        #   in Loop: Header=BB2_7 Depth=1
	xorl	%r13d, %r13d
	movsd	%xmm1, 32(%rsp)                 # 8-byte Spill
.LBB2_12:                               # %.lr.ph293
                                        #   Parent Loop BB2_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r13, %xmm0
	mulsd	%xmm1, %xmm0
	movl	$.L.str.4, %esi
	movq	%r15, %rdi
	movb	$1, %al
	callq	fprintf
	movsd	32(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	incq	%r13
	movslq	(%rbx,%r12,4), %rax
	incq	%rax
	cmpq	%rax, %r13
	jl	.LBB2_12
	jmp	.LBB2_13
.LBB2_8:                                #   in Loop: Header=BB2_7 Depth=1
	xorl	%r13d, %r13d
.LBB2_13:                               # %._crit_edge294
                                        #   in Loop: Header=BB2_7 Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r13, %xmm0
	mulsd	%xmm0, %xmm1
	movl	$.L.str.5, %esi
	movq	%r15, %rdi
	movapd	%xmm1, %xmm0
	movb	$1, %al
	callq	fprintf
	cmpl	$0, (%r14,%r12,4)
	js	.LBB2_14
# %bb.15:                               # %.lr.ph298
                                        #   in Loop: Header=BB2_7 Depth=1
	xorl	%r13d, %r13d
	movq	24(%rsp), %rbp                  # 8-byte Reload
.LBB2_16:                               #   Parent Loop BB2_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r13, %xmm0
	mulsd	(%rbp,%r12,8), %xmm0
	movl	$.L.str.4, %esi
	movq	%r15, %rdi
	movb	$1, %al
	callq	fprintf
	incq	%r13
	movslq	(%r14,%r12,4), %rax
	incq	%rax
	cmpq	%rax, %r13
	jl	.LBB2_16
	jmp	.LBB2_17
.LBB2_14:                               #   in Loop: Header=BB2_7 Depth=1
	xorl	%r13d, %r13d
	movq	24(%rsp), %rbp                  # 8-byte Reload
.LBB2_17:                               # %._crit_edge299
                                        #   in Loop: Header=BB2_7 Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r13, %xmm0
	mulsd	(%rbp,%r12,8), %xmm0
	movl	$.L.str.5, %esi
	movq	%r15, %rdi
	movb	$1, %al
	callq	fprintf
	movl	(%r14,%r12,4), %edx
	addl	$2, %edx
	movl	(%rbx,%r12,4), %ecx
	leal	2(%rcx), %eax
	imull	%eax, %edx
	testl	%edx, %edx
	jle	.LBB2_22
# %bb.18:                               # %.lr.ph305.preheader
                                        #   in Loop: Header=BB2_7 Depth=1
	movq	40(%rsp), %rdx                  # 8-byte Reload
	movq	48(%rsp), %rsi                  # 8-byte Reload
	leaq	(%rdx,%rsi,8), %r13
	xorl	%ebp, %ebp
.LBB2_19:                               # %.lr.ph305
                                        #   Parent Loop BB2_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movslq	%eax, %rsi
	movq	%rbp, %rax
	cqto
	idivq	%rsi
	incl	%ecx
	movslq	%ecx, %rax
	movl	$.L.str.5, %esi
	cmpq	%rax, %rdx
	je	.LBB2_21
# %bb.20:                               # %.lr.ph305
                                        #   in Loop: Header=BB2_19 Depth=2
	movl	$.L.str.4, %esi
.LBB2_21:                               # %.lr.ph305
                                        #   in Loop: Header=BB2_19 Depth=2
	movsd	(%r13,%rbp,8), %xmm0            # xmm0 = mem[0],zero
	movq	%r15, %rdi
	movb	$1, %al
	callq	fprintf
	incq	%rbp
	movslq	(%r14,%r12,4), %rdx
	addq	$2, %rdx
	movslq	(%rbx,%r12,4), %rcx
	movq	%rcx, %rax
	addq	$2, %rax
	movq	%rax, %rsi
	imulq	%rdx, %rsi
	cmpq	%rsi, %rbp
	jl	.LBB2_19
.LBB2_22:                               # %._crit_edge306
                                        #   in Loop: Header=BB2_7 Depth=1
	movq	%r15, %rdi
	callq	fclose
	incq	%r12
	movslq	4(%rsp), %rax
	cmpq	%rax, %r12
	movq	88(%rsp), %r13                  # 8-byte Reload
	jl	.LBB2_7
.LBB2_23:                               # %._crit_edge310
	callq	clock
	movq	%rax, 136(%rsp)                 # 8-byte Spill
	movl	4(%rsp), %edi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	%rdx, %rdi
	addq	$1024, %rdx                     # imm = 0x400
	movl	$49152, %r8d                    # imm = 0xC000
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movq	184(%rsp), %r15                 # 8-byte Reload
	movq	72(%rsp), %rbp                  # 8-byte Reload
	jne	.LBB2_25
# %bb.24:
	movq	120(%rsp), %rdi
	movq	232(%rsp), %rsi
	movq	224(%rsp), %rdx
	movq	208(%rsp), %rcx
	movq	104(%rsp), %r8
	movq	216(%rsp), %r9
	pushq	96(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	120(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z18__device_stub__runPdS_S_PlS0_PiS1_S0_
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
.LBB2_25:
	callq	hipDeviceSynchronize
	callq	clock
	movq	%rax, 128(%rsp)                 # 8-byte Spill
	movq	120(%rsp), %rsi
	movq	40(%rsp), %rdi                  # 8-byte Reload
	movq	176(%rsp), %rdx                 # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	cmpl	$0, 4(%rsp)
	jle	.LBB2_41
# %bb.26:                               # %.lr.ph331
	xorl	%r12d, %r12d
.LBB2_27:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_30 Depth 2
                                        #     Child Loop BB2_34 Depth 2
                                        #     Child Loop BB2_37 Depth 2
	movq	56(%rsp), %rax                  # 8-byte Reload
	movq	(%rax,%r12,8), %rax
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movw	$25970, 8(%rsp)                 # imm = 0x6572
	movb	$115, 10(%rsp)
	movl	$1954051118, 13(%rsp)           # imm = 0x7478742E
	movb	$0, 17(%rsp)
	movl	%r12d, %eax
	movl	$3435973837, %ecx               # imm = 0xCCCCCCCD
	imulq	%rcx, %rax
	shrq	$35, %rax
	leal	(%rax,%rax), %ecx
	leal	(%rcx,%rcx,4), %ecx
	movl	%r12d, %edx
	subl	%ecx, %edx
	addb	$48, %al
	movb	%al, 11(%rsp)
	orb	$48, %dl
	movb	%dl, 12(%rsp)
	movl	$.L.str.3, %esi
	leaq	8(%rsp), %rdi
	callq	fopen
	movq	%rax, %r15
	cmpl	$0, (%rbx,%r12,4)
	movsd	(%r13,%r12,8), %xmm1            # xmm1 = mem[0],zero
	js	.LBB2_28
# %bb.29:                               # %.lr.ph313.preheader
                                        #   in Loop: Header=BB2_27 Depth=1
	xorl	%r13d, %r13d
	movsd	%xmm1, 32(%rsp)                 # 8-byte Spill
.LBB2_30:                               # %.lr.ph313
                                        #   Parent Loop BB2_27 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r13, %xmm0
	mulsd	%xmm1, %xmm0
	movl	$.L.str.4, %esi
	movq	%r15, %rdi
	movb	$1, %al
	callq	fprintf
	movsd	32(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	incq	%r13
	movslq	(%rbx,%r12,4), %rax
	incq	%rax
	cmpq	%rax, %r13
	jl	.LBB2_30
	jmp	.LBB2_31
.LBB2_28:                               #   in Loop: Header=BB2_27 Depth=1
	xorl	%r13d, %r13d
.LBB2_31:                               # %._crit_edge314
                                        #   in Loop: Header=BB2_27 Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r13, %xmm0
	mulsd	%xmm0, %xmm1
	movl	$.L.str.5, %esi
	movq	%r15, %rdi
	movapd	%xmm1, %xmm0
	movb	$1, %al
	callq	fprintf
	cmpl	$0, (%r14,%r12,4)
	js	.LBB2_32
# %bb.33:                               # %.lr.ph318
                                        #   in Loop: Header=BB2_27 Depth=1
	xorl	%r13d, %r13d
	movq	24(%rsp), %rbp                  # 8-byte Reload
.LBB2_34:                               #   Parent Loop BB2_27 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r13, %xmm0
	mulsd	(%rbp,%r12,8), %xmm0
	movl	$.L.str.4, %esi
	movq	%r15, %rdi
	movb	$1, %al
	callq	fprintf
	incq	%r13
	movslq	(%r14,%r12,4), %rax
	incq	%rax
	cmpq	%rax, %r13
	jl	.LBB2_34
	jmp	.LBB2_35
.LBB2_32:                               #   in Loop: Header=BB2_27 Depth=1
	xorl	%r13d, %r13d
	movq	24(%rsp), %rbp                  # 8-byte Reload
.LBB2_35:                               # %._crit_edge319
                                        #   in Loop: Header=BB2_27 Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r13, %xmm0
	mulsd	(%rbp,%r12,8), %xmm0
	movl	$.L.str.5, %esi
	movq	%r15, %rdi
	movb	$1, %al
	callq	fprintf
	movl	(%r14,%r12,4), %edx
	addl	$2, %edx
	movl	(%rbx,%r12,4), %ecx
	leal	2(%rcx), %eax
	imull	%eax, %edx
	testl	%edx, %edx
	jle	.LBB2_40
# %bb.36:                               # %.lr.ph327.preheader
                                        #   in Loop: Header=BB2_27 Depth=1
	movq	40(%rsp), %rdx                  # 8-byte Reload
	movq	48(%rsp), %rsi                  # 8-byte Reload
	leaq	(%rdx,%rsi,8), %r13
	xorl	%ebp, %ebp
.LBB2_37:                               # %.lr.ph327
                                        #   Parent Loop BB2_27 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movslq	%eax, %rsi
	movq	%rbp, %rax
	cqto
	idivq	%rsi
	incl	%ecx
	movslq	%ecx, %rax
	movl	$.L.str.5, %esi
	cmpq	%rax, %rdx
	je	.LBB2_39
# %bb.38:                               # %.lr.ph327
                                        #   in Loop: Header=BB2_37 Depth=2
	movl	$.L.str.4, %esi
.LBB2_39:                               # %.lr.ph327
                                        #   in Loop: Header=BB2_37 Depth=2
	movsd	(%r13,%rbp,8), %xmm0            # xmm0 = mem[0],zero
	movq	%r15, %rdi
	movb	$1, %al
	callq	fprintf
	incq	%rbp
	movslq	(%r14,%r12,4), %rdx
	addq	$2, %rdx
	movslq	(%rbx,%r12,4), %rcx
	movq	%rcx, %rax
	addq	$2, %rax
	movq	%rax, %rsi
	imulq	%rdx, %rsi
	cmpq	%rsi, %rbp
	jl	.LBB2_37
.LBB2_40:                               # %._crit_edge328
                                        #   in Loop: Header=BB2_27 Depth=1
	movq	%r15, %rdi
	callq	fclose
	incq	%r12
	movslq	4(%rsp), %rax
	cmpq	%rax, %r12
	movq	88(%rsp), %r13                  # 8-byte Reload
	movq	184(%rsp), %r15                 # 8-byte Reload
	movq	72(%rsp), %rbp                  # 8-byte Reload
	jl	.LBB2_27
.LBB2_41:                               # %._crit_edge332
	movq	128(%rsp), %rax                 # 8-byte Reload
	subq	136(%rsp), %rax                 # 8-byte Folded Reload
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI2_5(%rip), %xmm0
	movsd	%xmm0, 32(%rsp)                 # 8-byte Spill
	movq	80(%rsp), %rdi                  # 8-byte Reload
	callq	free
	movq	200(%rsp), %rdi                 # 8-byte Reload
	callq	free
	movq	64(%rsp), %rdi                  # 8-byte Reload
	callq	free
	movq	192(%rsp), %rdi                 # 8-byte Reload
	callq	free
	movq	24(%rsp), %rdi                  # 8-byte Reload
	callq	free
	movq	%r13, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movq	%rbp, %rdi
	callq	free
	movq	168(%rsp), %rdi                 # 8-byte Reload
	callq	free
	movq	160(%rsp), %rdi                 # 8-byte Reload
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	%rbx, %rdi
	callq	free
	movq	152(%rsp), %rdi                 # 8-byte Reload
	callq	free
	movq	144(%rsp), %rdi                 # 8-byte Reload
	callq	free
	movq	56(%rsp), %rdi                  # 8-byte Reload
	callq	free
	movq	264(%rsp), %rdi
	callq	hipFree
	movq	256(%rsp), %rdi
	callq	hipFree
	movq	248(%rsp), %rdi
	callq	hipFree
	movq	240(%rsp), %rdi
	callq	hipFree
	movq	232(%rsp), %rdi
	callq	hipFree
	movq	224(%rsp), %rdi
	callq	hipFree
	movq	216(%rsp), %rdi
	callq	hipFree
	movq	112(%rsp), %rdi
	callq	hipFree
	movq	208(%rsp), %rdi
	callq	hipFree
	movq	104(%rsp), %rdi
	callq	hipFree
	movq	96(%rsp), %rdi
	callq	hipFree
	movl	$.L.str.6, %edi
	movsd	32(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$1, %al
	callq	printf
	xorl	%eax, %eax
	addq	$536, %rsp                      # imm = 0x218
	.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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z4initPdS_S_S_S_PlPiS0_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z3runPdS_S_PlS0_PiS1_S0_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.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	_Z4initPdS_S_S_S_PlPiS0_,@object # @_Z4initPdS_S_S_S_PlPiS0_
	.section	.rodata,"a",@progbits
	.globl	_Z4initPdS_S_S_S_PlPiS0_
	.p2align	3, 0x0
_Z4initPdS_S_S_S_PlPiS0_:
	.quad	_Z19__device_stub__initPdS_S_S_S_PlPiS0_
	.size	_Z4initPdS_S_S_S_PlPiS0_, 8

	.type	_Z3runPdS_S_PlS0_PiS1_S0_,@object # @_Z3runPdS_S_PlS0_PiS1_S0_
	.globl	_Z3runPdS_S_PlS0_PiS1_S0_
	.p2align	3, 0x0
_Z3runPdS_S_PlS0_PiS1_S0_:
	.quad	_Z18__device_stub__runPdS_S_PlS0_PiS1_S0_
	.size	_Z3runPdS_S_PlS0_PiS1_S0_, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"%d"
	.size	.L.str, 3

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

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"%lf"
	.size	.L.str.2, 4

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

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

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

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"\n------------------------------------\nExecution took: %lf sec\n"
	.size	.L.str.6, 63

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z4initPdS_S_S_S_PlPiS0_"
	.size	.L__unnamed_1, 25

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z3runPdS_S_PlS0_PiS1_S0_"
	.size	.L__unnamed_2, 26

	.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 _Z19__device_stub__initPdS_S_S_S_PlPiS0_
	.addrsig_sym _Z18__device_stub__runPdS_S_PlS0_PiS1_S0_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z4initPdS_S_S_S_PlPiS0_
	.addrsig_sym _Z3runPdS_S_PlS0_PiS1_S0_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z4initPdS_S_S_S_PlPiS0_ ; -- Begin function _Z4initPdS_S_S_S_PlPiS0_
	.globl	_Z4initPdS_S_S_S_PlPiS0_
	.p2align	8
	.type	_Z4initPdS_S_S_S_PlPiS0_,@function
_Z4initPdS_S_S_S_PlPiS0_:               ; @_Z4initPdS_S_S_S_PlPiS0_
; %bb.0:
	s_mov_b32 s18, s15
	s_load_b512 s[0:15], s[0:1], 0x0
	s_mov_b32 s19, 0
                                        ; implicit-def: $vgpr1
                                        ; implicit-def: $vgpr3_vgpr4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b64 s[20:21], s[18:19], 3
	s_waitcnt lgkmcnt(0)
	s_add_u32 s16, s2, s20
	s_addc_u32 s17, s3, s21
	s_add_u32 s4, s4, s20
	s_addc_u32 s5, s5, s21
	s_add_u32 s6, s6, s20
	s_addc_u32 s7, s7, s21
	s_add_u32 s2, s8, s20
	s_addc_u32 s3, s9, s21
	s_add_u32 s8, s10, s20
	s_addc_u32 s9, s11, s21
	s_load_b64 s[22:23], s[8:9], 0x0
	s_lshl_b64 s[8:9], s[18:19], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s8, s12, s8
	s_addc_u32 s9, s13, s9
	s_add_u32 s10, s14, s20
	s_addc_u32 s11, s15, s21
	s_waitcnt lgkmcnt(0)
	s_ashr_i32 s12, s23, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshr_b32 s12, s12, 22
	s_add_u32 s12, s22, s12
	s_addc_u32 s13, s23, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_ashr_i64 s[14:15], s[12:13], 10
	s_lshl_b32 s12, s14, 10
	s_delay_alu instid0(SALU_CYCLE_1)
	s_sub_i32 s18, s22, s12
	s_bfe_i64 s[12:13], s[14:15], 0x200000
	s_mov_b32 s15, exec_lo
	v_cmpx_le_u32_e64 s18, v0
	s_xor_b32 s15, exec_lo, s15
; %bb.1:
	v_mad_u64_u32 v[1:2], null, v0, s14, s[18:19]
                                        ; implicit-def: $vgpr0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v3, s18, s12, v1
	v_add_co_ci_u32_e64 v4, null, s13, 0, s18
; %bb.2:                                ; %Flow80
	s_and_not1_saveexec_b32 s15, s15
; %bb.3:
	v_mad_u64_u32 v[1:2], null, v0, s14, v[0:1]
	s_add_u32 s12, s12, 1
	s_addc_u32 s13, s13, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v3, s12, s12, v1
	v_add_co_ci_u32_e64 v4, null, s13, 0, s12
; %bb.4:
	s_or_b32 exec_lo, exec_lo, s15
	v_ashrrev_i32_e32 v2, 31, v1
	s_mov_b32 s12, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i64_e64 v[3:4], v[1:2]
	s_cbranch_execz .LBB0_7
; %bb.5:                                ; %.lr.ph
	s_load_b64 s[18:19], s[2:3], 0x0
	s_load_b64 s[10:11], s[10:11], 0x0
	s_mov_b32 s3, 0x402921fb
	s_mov_b32 s2, 0x54442d18
	s_load_b64 s[4:5], s[4:5], 0x0
	s_load_b64 s[6:7], s[6:7], 0x0
	s_mov_b32 s15, 0
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[5:6], 0x40080000, s[18:19]
	v_mul_f64 v[7:8], s[18:19], s[2:3]
	s_load_b32 s3, s[8:9], 0x0
	s_lshl_b64 s[8:9], s[10:11], 3
	s_waitcnt lgkmcnt(0)
	s_add_i32 s10, s3, 2
	s_add_u32 s12, s0, s8
	s_addc_u32 s1, s1, s9
	s_ashr_i32 s13, s10, 31
	s_load_b64 s[8:9], s[16:17], 0x0
	s_add_i32 s10, s10, s13
	s_sub_i32 s16, -2, s3
	s_xor_b32 s14, s10, s13
	s_mov_b32 s10, 0x9abcaf48
	v_cvt_f32_u32_e32 v0, s14
	s_sub_i32 s0, 0, s14
	s_mov_b32 s11, 0x3e7ad7f2
	s_mov_b32 s3, 0x400921fb
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_rcp_iflag_f32_e32 v0, v0
	v_mul_f64 v[5:6], s[18:19], v[5:6]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_f64 v[7:8], s[18:19], v[7:8]
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, 0x4f7ffffe, v0
	v_cvt_u32_f32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v9, s0, v0
	v_mul_hi_u32 v9, v0, v9
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v0, v9
.LBB0_6:                                ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v9, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v10, v1, v9
	v_xor_b32_e32 v10, v10, v9
	v_xor_b32_e32 v9, s13, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v11, v10, v0
	v_mul_lo_u32 v12, v11, s14
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v10, v10, v12
	v_subrev_nc_u32_e32 v12, s14, v10
	v_cmp_le_u32_e32 vcc_lo, s14, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v10, v10, v12 :: v_dual_add_nc_u32 v13, 1, v11
	v_cndmask_b32_e32 v11, v11, v13, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_le_u32_e32 vcc_lo, s14, v10
	v_add_nc_u32_e32 v12, 1, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v10, v11, v12, vcc_lo
	v_xor_b32_e32 v10, v10, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v19, v10, v9
	v_mad_i64_i32 v[9:10], null, v19, v19, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_f64_i32_e32 v[10:11], v10
	v_cvt_f64_u32_e32 v[12:13], v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ldexp_f64 v[9:10], v[10:11], 32
	v_add_f64 v[9:10], v[9:10], v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[9:10], s[6:7], v[9:10]
	v_mul_f64 v[9:10], s[6:7], v[9:10]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f64 v[11:12], null, v[5:6], v[5:6], v[9:10]
	v_div_scale_f64 v[17:18], vcc_lo, v[9:10], v[5:6], v[9:10]
	v_rcp_f64_e32 v[13:14], v[11:12]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[15:16], -v[11:12], v[13:14], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[13:14], v[13:14], v[15:16], v[13:14]
	v_fma_f64 v[15:16], -v[11:12], v[13:14], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[13:14], v[13:14], v[15:16], v[13:14]
	v_mul_f64 v[15:16], v[17:18], v[13:14]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[11:12], -v[11:12], v[15:16], v[17:18]
	v_div_fmas_f64 v[11:12], v[11:12], v[13:14], v[15:16]
	v_cvt_f64_i32_e32 v[13:14], v19
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_div_fixup_f64 v[9:10], v[11:12], v[5:6], v[9:10]
	v_mad_u64_u32 v[11:12], null, s16, v19, v[1:2]
	v_cvt_f64_i32_e32 v[11:12], v11
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[9:10], -v[9:10], 1.0
	s_waitcnt lgkmcnt(0)
	v_fma_f64 v[11:12], s[4:5], v[11:12], s[8:9]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[9:10], 0x40080000, v[9:10]
	v_mul_f64 v[9:10], v[9:10], 4.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[9:10], v[9:10], s[2:3]
	v_mul_f64 v[9:10], v[9:10], s[10:11]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[9:10], v[11:12], v[9:10]
	v_mul_f64 v[9:10], v[9:10], v[13:14]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[9:10], s[6:7], v[9:10]
	v_div_scale_f64 v[11:12], null, v[7:8], v[7:8], v[9:10]
	v_div_scale_f64 v[17:18], vcc_lo, v[9:10], v[7:8], v[9:10]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[13:14], v[11:12]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[15:16], -v[11:12], v[13:14], 1.0
	v_fma_f64 v[13:14], v[13:14], v[15:16], v[13:14]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[15:16], -v[11:12], v[13:14], 1.0
	v_fma_f64 v[13:14], v[13:14], v[15:16], v[13:14]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[15:16], v[17:18], v[13:14]
	v_fma_f64 v[11:12], -v[11:12], v[15:16], v[17:18]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[11:12], v[11:12], v[13:14], v[15:16]
	v_div_fixup_f64 v[9:10], v[11:12], v[7:8], v[9:10]
	v_lshlrev_b64 v[11:12], 3, v[1:2]
	v_add_nc_u32_e32 v1, 1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v2, 31, v1
	v_add_co_u32 v11, s0, s12, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v12, s0, s1, v12, s0
	v_cmp_le_i64_e32 vcc_lo, v[3:4], v[1:2]
	s_or_b32 s15, vcc_lo, s15
	global_store_b64 v[11:12], v[9:10], off
	s_and_not1_b32 exec_lo, exec_lo, s15
	s_cbranch_execnz .LBB0_6
.LBB0_7:                                ; %Flow79
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z4initPdS_S_S_S_PlPiS0_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 64
		.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 20
		.amdhsa_next_free_sgpr 24
		.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	_Z4initPdS_S_S_S_PlPiS0_, .Lfunc_end0-_Z4initPdS_S_S_S_PlPiS0_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1064
; NumSgprs: 26
; NumVgprs: 20
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 26
; NumVGPRsForWavesPerEU: 20
; 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
	.protected	_Z3runPdS_S_PlS0_PiS1_S0_ ; -- Begin function _Z3runPdS_S_PlS0_PiS1_S0_
	.globl	_Z3runPdS_S_PlS0_PiS1_S0_
	.p2align	8
	.type	_Z3runPdS_S_PlS0_PiS1_S0_,@function
_Z3runPdS_S_PlS0_PiS1_S0_:              ; @_Z3runPdS_S_PlS0_PiS1_S0_
; %bb.0:
	s_mov_b32 s20, s15
	s_load_b512 s[4:19], s[0:1], 0x0
	s_mov_b32 s21, 0
                                        ; implicit-def: $vgpr6
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b64 s[22:23], s[20:21], 3
	s_waitcnt lgkmcnt(0)
	s_add_u32 s0, s6, s22
	s_addc_u32 s1, s7, s23
	s_add_u32 s2, s8, s22
	s_addc_u32 s3, s9, s23
	s_add_u32 s6, s10, s22
	s_addc_u32 s7, s11, s23
	s_add_u32 s8, s12, s22
	s_addc_u32 s9, s13, s23
	s_load_b64 s[10:11], s[8:9], 0x0
	s_lshl_b64 s[8:9], s[20:21], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s14, s14, s8
	s_addc_u32 s15, s15, s9
	s_add_u32 s8, s16, s8
	s_addc_u32 s9, s17, s9
	s_add_u32 s16, s18, s22
	s_addc_u32 s17, s19, s23
	s_waitcnt lgkmcnt(0)
	s_ashr_i32 s12, s11, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshr_b32 s12, s12, 22
	s_add_u32 s12, s10, s12
	s_addc_u32 s11, s11, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v4, s11, s12, 10
	v_lshlrev_b32_e32 v1, 10, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, s10, v1
	s_mov_b32 s10, exec_lo
                                        ; implicit-def: $vgpr1
	v_cmpx_ge_u32_e64 v0, v3
	s_xor_b32 s10, exec_lo, s10
; %bb.1:
	v_mad_u64_u32 v[1:2], null, v0, v4, v[3:4]
                                        ; implicit-def: $vgpr0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v1, v4
                                        ; implicit-def: $vgpr4
; %bb.2:                                ; %Flow199
	s_and_not1_saveexec_b32 s10, s10
; %bb.3:
	v_add_nc_u32_e32 v2, 1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v1, v2, v0
	v_add_nc_u32_e32 v6, v2, v1
; %bb.4:
	s_or_b32 exec_lo, exec_lo, s10
	s_load_b32 s20, s[8:9], 0x0
	s_load_b64 s[8:9], s[0:1], 0x0
	s_load_b64 s[10:11], s[2:3], 0x0
	s_load_b64 s[12:13], s[6:7], 0x0
	s_load_b32 s19, s[14:15], 0x0
	s_load_b64 s[6:7], s[16:17], 0x0
	v_cmp_lt_i32_e64 s0, v1, v6
	v_lshl_add_u32 v0, v1, 3, 0
	s_delay_alu instid0(VALU_DEP_2)
	s_and_saveexec_b32 s2, s0
	s_cbranch_execz .LBB1_7
; %bb.5:                                ; %.lr.ph
	v_ashrrev_i32_e32 v2, 31, v1
	s_waitcnt lgkmcnt(0)
	s_lshl_b64 s[14:15], s[6:7], 3
	v_lshl_add_u32 v4, v1, 3, 0
	s_add_u32 s1, s4, s14
	s_addc_u32 s3, s5, s15
	v_lshlrev_b64 v[2:3], 3, v[1:2]
	v_mov_b32_e32 v5, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v2, vcc_lo, s1, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s3, v3, vcc_lo
	s_mov_b32 s3, 0
.LBB1_6:                                ; =>This Inner Loop Header: Depth=1
	global_load_b64 v[7:8], v[2:3], off
	v_add_nc_u32_e32 v5, 1, v5
	v_add_co_u32 v2, vcc_lo, v2, 8
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e64 s1, v5, v6
	s_or_b32 s3, s1, s3
	s_waitcnt vmcnt(0)
	ds_store_b64 v4, v[7:8]
	v_add_nc_u32_e32 v4, 8, v4
	s_and_not1_b32 exec_lo, exec_lo, s3
	s_cbranch_execnz .LBB1_6
.LBB1_7:                                ; %Flow198
	s_or_b32 exec_lo, exec_lo, s2
	s_waitcnt lgkmcnt(0)
	v_cmp_lt_i64_e64 s1, s[12:13], 1
	s_barrier
	buffer_gl0_inv
	s_and_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccnz .LBB1_23
; %bb.8:                                ; %.preheader104.lr.ph
	s_add_i32 s3, s20, 2
	v_lshlrev_b32_e32 v4, 3, v1
	s_ashr_i32 s16, s3, 31
	v_add_lshl_u32 v5, s20, v1, 3
	s_add_i32 s1, s3, s16
	v_lshl_add_u32 v7, v1, 3, 0
	s_xor_b32 s17, s1, s16
	v_add3_u32 v8, 0, -8, v4
	v_cvt_f32_u32_e32 v2, s17
	s_sub_i32 s1, 0, s17
	v_add3_u32 v9, 0, 16, v5
	s_add_i32 s18, s20, 1
	s_add_i32 s19, s19, 1
	v_rcp_iflag_f32_e32 v2, v2
	s_mov_b64 s[14:15], 0
	s_waitcnt_depctr 0xfff
	v_dual_mul_f32 v2, 0x4f7ffffe, v2 :: v_dual_add_nc_u32 v11, 0, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v2, v2
	v_mul_lo_u32 v3, s1, v2
	s_lshl_b32 s1, s20, 3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_sub_i32 s20, -16, s1
	v_mul_hi_u32 v3, v2, v3
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v10, v2, v3
.LBB1_9:                                ; %.preheader104
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_11 Depth 2
                                        ;     Child Loop BB1_14 Depth 2
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB1_12
; %bb.10:                               ; %.lr.ph107.preheader
                                        ;   in Loop: Header=BB1_9 Depth=1
	v_mov_b32_e32 v2, v7
	v_dual_mov_b32 v3, v7 :: v_dual_mov_b32 v4, v1
	s_mov_b32 s2, 0
.LBB1_11:                               ; %.lr.ph107
                                        ;   Parent Loop BB1_9 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	ds_load_b64 v[12:13], v3
	v_add_nc_u32_e32 v4, 1, v4
	v_add_nc_u32_e32 v3, 8, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_cmp_ge_i32_e32 vcc_lo, v4, v6
	s_or_b32 s2, vcc_lo, s2
	s_waitcnt lgkmcnt(0)
	ds_store_b64 v2, v[12:13]
	v_add_nc_u32_e32 v2, 8, v2
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB1_11
.LBB1_12:                               ; %Flow194
                                        ;   in Loop: Header=BB1_9 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_and_saveexec_b32 s21, s0
	s_cbranch_execz .LBB1_22
; %bb.13:                               ; %.lr.ph111.preheader
                                        ;   in Loop: Header=BB1_9 Depth=1
	v_dual_mov_b32 v12, v11 :: v_dual_mov_b32 v13, v9
	v_dual_mov_b32 v14, v8 :: v_dual_mov_b32 v15, v1
	s_mov_b32 s22, 0
.LBB1_14:                               ; %.lr.ph111
                                        ;   Parent Loop BB1_9 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v2, 31, v15
	s_mov_b32 s23, exec_lo
	v_add_nc_u32_e32 v3, v15, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v3, v3, v2
	v_xor_b32_e32 v2, s16, v2
	v_mul_hi_u32 v4, v3, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v5, v4, s17
	v_sub_nc_u32_e32 v3, v3, v5
	v_add_nc_u32_e32 v5, 1, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v16, s17, v3
	v_cmp_le_u32_e32 vcc_lo, s17, v3
	v_dual_cndmask_b32 v4, v4, v5 :: v_dual_cndmask_b32 v3, v3, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v5, 1, v4
	v_cmp_le_u32_e32 vcc_lo, s17, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v3, v4, v5, vcc_lo
	v_xor_b32_e32 v3, v3, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v16, v3, v2
	v_mul_lo_u32 v2, v16, s3
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ne_u32_e64 v2, v15
	s_cbranch_execz .LBB1_21
; %bb.15:                               ;   in Loop: Header=BB1_14 Depth=2
	v_add_nc_u32_e32 v2, s18, v2
	v_cmp_ne_u32_e64 s1, 0, v16
	v_cmp_ne_u32_e64 s2, s19, v16
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_ne_u32_e32 vcc_lo, v2, v15
	s_and_b32 s1, vcc_lo, s1
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	s_and_b32 s1, s1, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 exec_lo, exec_lo, s1
	s_cbranch_execz .LBB1_21
; %bb.16:                               ;   in Loop: Header=BB1_14 Depth=2
	s_mov_b32 s1, exec_lo
                                        ; implicit-def: $vgpr2_vgpr3
                                        ; implicit-def: $vgpr4_vgpr5
	v_cmpx_ne_u32_e32 1, v16
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB1_18
; %bb.17:                               ;   in Loop: Header=BB1_14 Depth=2
	v_mul_lo_u32 v2, v16, v16
	v_lshlrev_b32_e32 v16, 1, v16
	ds_load_b64 v[18:19], v13
	v_or_b32_e32 v17, 1, v16
	v_sub_nc_u32_e32 v3, -2, v2
	v_cmp_gt_u32_e32 vcc_lo, 2, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v2, -2, v3, vcc_lo
	v_cmp_gt_u32_e32 vcc_lo, 3, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cvt_f64_i32_e32 v[4:5], v2
	ds_load_b64 v[2:3], v14 offset:8
	v_cndmask_b32_e32 v20, 0, v16, vcc_lo
	v_or_b32_e32 v16, 1, v20
	v_sub_nc_u32_e32 v20, 1, v20
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cvt_f64_i32_e32 v[16:17], v16
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[4:5], v[2:3], v[4:5]
	v_fma_f64 v[4:5], v[18:19], v[16:17], v[4:5]
	v_cvt_f64_i32_e32 v[16:17], v20
	v_add_nc_u32_e32 v18, s20, v14
	ds_load_b64 v[18:19], v18 offset:8
	s_waitcnt lgkmcnt(0)
	v_fma_f64 v[4:5], v[18:19], v[16:17], v[4:5]
.LBB1_18:                               ; %Flow189
                                        ;   in Loop: Header=BB1_14 Depth=2
	s_and_not1_saveexec_b32 s1, s1
	s_cbranch_execz .LBB1_20
; %bb.19:                               ;   in Loop: Header=BB1_14 Depth=2
	ds_load_b64 v[2:3], v14 offset:8
	ds_load_b64 v[16:17], v13
	s_waitcnt lgkmcnt(1)
	v_mul_f64 v[4:5], v[2:3], -4.0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[4:5], v[16:17], 2.0, v[4:5]
.LBB1_20:                               ; %.sink.split
                                        ;   in Loop: Header=BB1_14 Depth=2
	s_or_b32 exec_lo, exec_lo, s1
	ds_load_2addr_b64 v[16:19], v14 offset1:2
	s_waitcnt lgkmcnt(0)
	v_fma_f64 v[2:3], v[2:3], -2.0, v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[2:3], v[2:3], v[16:17]
	v_mul_f64 v[2:3], s[10:11], v[2:3]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[2:3], s[8:9], v[4:5], v[2:3]
	ds_load_b64 v[4:5], v12
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[2:3], v[4:5], v[2:3]
	ds_store_b64 v12, v[2:3]
.LBB1_21:                               ;   in Loop: Header=BB1_14 Depth=2
	s_or_b32 exec_lo, exec_lo, s23
	v_add_nc_u32_e32 v15, 1, v15
	v_add_nc_u32_e32 v14, 8, v14
	v_add_nc_u32_e32 v13, 8, v13
	v_add_nc_u32_e32 v12, 8, v12
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_ge_i32_e32 vcc_lo, v15, v6
	s_or_b32 s22, vcc_lo, s22
	s_and_not1_b32 exec_lo, exec_lo, s22
	s_cbranch_execnz .LBB1_14
.LBB1_22:                               ; %Flow192
                                        ;   in Loop: Header=BB1_9 Depth=1
	s_or_b32 exec_lo, exec_lo, s21
	s_add_u32 s14, s14, 1
	s_addc_u32 s15, s15, 0
	s_waitcnt lgkmcnt(0)
	s_cmp_eq_u64 s[14:15], s[12:13]
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB1_9
.LBB1_23:                               ; %.preheader
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB1_26
; %bb.24:                               ; %.lr.ph115
	v_ashrrev_i32_e32 v2, 31, v1
	s_lshl_b64 s[0:1], s[6:7], 3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	s_add_u32 s0, s4, s0
	s_addc_u32 s1, s5, s1
	v_lshlrev_b64 v[2:3], 3, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s0, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v3, vcc_lo
	s_mov_b32 s1, 0
.LBB1_25:                               ; =>This Inner Loop Header: Depth=1
	ds_load_b64 v[4:5], v0
	v_add_nc_u32_e32 v1, 1, v1
	v_add_nc_u32_e32 v0, 8, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_cmp_ge_i32_e32 vcc_lo, v1, v6
	s_or_b32 s1, vcc_lo, s1
	s_waitcnt lgkmcnt(0)
	global_store_b64 v[2:3], v[4:5], off
	v_add_co_u32 v2, s0, v2, 8
	v_add_co_ci_u32_e64 v3, s0, 0, v3, s0
	s_and_not1_b32 exec_lo, exec_lo, s1
	s_cbranch_execnz .LBB1_25
.LBB1_26:                               ; %Flow188
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z3runPdS_S_PlS0_PiS1_S0_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 64
		.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 21
		.amdhsa_next_free_sgpr 24
		.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_end1:
	.size	_Z3runPdS_S_PlS0_PiS1_S0_, .Lfunc_end1-_Z3runPdS_S_PlS0_PiS1_S0_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1376
; NumSgprs: 26
; NumVgprs: 21
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 26
; NumVGPRsForWavesPerEU: 21
; Occupancy: 16
; WaveLimiterHint : 0
; 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
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 64
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z4initPdS_S_S_S_PlPiS0_
    .private_segment_fixed_size: 0
    .sgpr_count:     26
    .sgpr_spill_count: 0
    .symbol:         _Z4initPdS_S_S_S_PlPiS0_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     20
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .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
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 64
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z3runPdS_S_PlS0_PiS1_S0_
    .private_segment_fixed_size: 0
    .sgpr_count:     26
    .sgpr_spill_count: 0
    .symbol:         _Z3runPdS_S_PlS0_PiS1_S0_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     21
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
