#include <stdio.h>
#include "math.h"
#include <pthread.h>
#include <unistd.h>
#include <string.h>
#include <sys/types.h>
#include <sys/socket.h>
#include <arpa/inet.h>
#include <netinet/in.h>
#define PORT    2084
#define MAXLINE 1024
/*
I will provide a word about tx,ty, and tz.
tx does not, in fact move the actual point (points are static), but defines how much mass will be donated.
After this donation, a density-based heat donation is calculated, but regardless of these trajectory values.
*/
__global__
void osmate(double tempRate, double massRate, double tempPush, double repulsionRate, int numPoints, double radius, double gConst, double * x, double * y, double * z, double * tx, double * ty, double * tz, double * newtx, double * newty, double * newtz, double * mass, double * temp, double * newmass, double * newtemp) {
  //The only part that really matters if I get the rest to work:
  int i=blockIdx.x + blockDim.x + threadIdx.x;
  for (int it=0;it<numPoints;it++) {//Iterate through all points
    // printf("%f",tx[it]);
    double dist=sqrt((x[it]-x[i])*(x[it]-x[i])+(y[it]-y[i])*(y[it]-y[i])+(z[it]-z[i])*(z[it]-z[i]));//This could be pre-calculated (these points are static). Whether the calculation happens should be a matter of how many points there are -- if there are too many points we will run out of memory, but if there are not as many it is a good idea. I'm leaving it for another day.
    if (dist!=0) {
      //Congratulations, you are not looking at the same point.
      double toChange=gConst*mass[i]*mass[it]/dist;
      newtx[i]=tx[i]+(toChange+tempPush*(temp[i]-temp[it])+repulsionRate*(mass[i]-mass[it]))*(x[it]-x[i]);
      newty[i]=ty[i]+(toChange+tempPush*(temp[i]-temp[it])+repulsionRate*(mass[i]-mass[it]))*(y[it]-y[i]);
      newtz[i]=tz[i]+(toChange+tempPush*(temp[i]-temp[it])+repulsionRate*(mass[i]-mass[it]))*(z[it]-z[i]);
    }
  }
  // if (x[i]!=0) {
  //   printf("%d\n",x[i]);
  // }
  for (int it=0;it<numPoints;it++) {
    //TODO make negative mass impossible.
    if (tx[i]!=0&&ty[i]!=0&&tz[i]!=0) {
      double massDisagreement=sqrt((x[it]-x[i])*(x[it]-x[i])/tx[i]+(y[it]-y[i])*(y[it]-y[i])/ty[i]+(z[it]-z[i])*(z[it]-z[i])/tz[i]);//TODO cases for rare 0 trajectory scenarios that would knacker it all.
      if (massDisagreement!=0) {
        newmass[it]=mass[it]+massRate/massDisagreement;//Hell if I know whether this is right, but it seems it should intuitively work. Also grants some cohesion. Note we don't subtract from the subject sample because the mass comes from behind it.
        double dist=sqrt((x[it]-x[i])*(x[it]-x[i])+(y[it]-y[i])*(y[it]-y[i])+(z[it]-z[i])*(z[it]-z[i]));//There are a lot of redundant calculations here...
        if (dist!=0) {
          double tempEx=tempRate*(temp[i]-temp[it])/dist;
          newtemp[it]=temp[it]+tempEx;
          newtemp[i]=temp[i]-tempEx;//Unless we want to go supernova. We should also have a mechanic where a mass increase "creates" heat, a decrease sucks it up, there is an external source of heat and heat can radiate away... Etc.
          // printf("%f",tx[i]);
        }
      }
    }
  }
}
//To tweak
double TEMP_RATE=0.1;
double MASS_RATE=0.4;
double TEMP_PUSH=0.03;
double REPULSION_RATE=1.0;
double SAMPLE_RADIUS=.1;
double GRAV_CONST=.0667;//Definitely not.

int ALTITUDE=5;//Radius of planet.
double LAYER_HEIGHT=.5;
int RES = 3;

//In generation, values may exceed maximums and minimums during simulation.
double leastAllowedMass=.001;
double mostAllowedMass=5;//We'll randomize between the two.

double leastAllowedTemp=.001;//Kelvin?
double mostAllowedTemp=50000;//A lot, hopefully not too much. Should check Planck heat.

double leastAllowedTrajectory=.0001;//Along axis, I'm not a madman
double mostAllowedTrajectory=2;
void delay(int number_of_millis)
{
    // Converting time into milli_seconds
    // int milli_seconds = 1000 * number_of_seconds;

    // Storing start time
    clock_t start_time = clock();

    // looping till required time is not acheived
    while (clock() < start_time + number_of_millis)
        ;
}
void serialize_float(char *buffer, float value)
{
    unsigned int ivalue = *((unsigned int*)&value); // warning assumes 32-bit "unsigned int"
    buffer[0] = ivalue >> 24;
    buffer[1] = ivalue >> 16;
    buffer[2] = ivalue >> 8;
    buffer[3] = ivalue;
}
char* doublearraytochar(double * vals, int len) {
  char * result=(char *)malloc(len*sizeof(char)*4);
  for (int i=0;i<len;i++) {
    char * tempArr=(char *)malloc(sizeof(char)*4);
    serialize_float(tempArr,vals[i]);
    for (int i2=0;i2<4;i2++) {
      result[i*4+i2]=tempArr[i2];
    }
  }
  return result;
}
char * jsonArray(double * vals, int len, int bufPerNum, char * result, char * output) {
  result[0]='[';
  for (int i=0;i<len;i++) {
    if (i!=len-1) {
      snprintf(output,bufPerNum*2,"%f,",vals[i]);
    } else {
      snprintf(output,bufPerNum*2,"%f]",vals[i]);
    }
    strcat(result,output);
    // for (int i2=0;i2<(bufPerNum+1)*sizeof(char);i2++) {
    //   result[i*bufPerNum+i2+1]=output[i2];
    // }
  }
  return result;
}
int main() {
  srand ( time ( NULL));
  printf("Hello, you are hopefully simulating tectonics. If not, get out of here.\n");
  //Generate points...
  //The Declaration of Variable Allocation:
  int numSamples=0;
  for (double alt=LAYER_HEIGHT;alt<ALTITUDE;alt+=LAYER_HEIGHT) {
    //You may notice we project a cube into a sphere. Primarily I'm too lazy to generate something like an icosahedron, secondarily it would not have fine resolution (only subdivisions) if we used an icosahedron, tertiarily this allows us to check for seams, which are screaming alarms something is wrong, and fourth uv spheres get ridiculously fine detail only at the poles, which is a waste.
    numSamples+=(int)(pow(RES,3)-pow(RES-2,3));
  }
  printf("Allocating for %i samples (pass to your renderer your RES, %i, and your ALTITUDE, %i, like so: python display.py <RES> <ALTITUDE>)...\n", numSamples, RES,ALTITUDE);
  double * x;
  double * y;
  double * z;
  double * tx;
  double * ty;
  double * tz;
  double * temp;
  double * mass;
  double * newtx;
  double * newty;
  double * newtz;
  double * newtemp;
  double * newmass;
  cudaMallocManaged(&x,numSamples*sizeof(double));
  cudaMallocManaged(&y,numSamples*sizeof(double));
  cudaMallocManaged(&z,numSamples*sizeof(double));
  cudaMallocManaged(&tx,numSamples*sizeof(double));
  cudaMallocManaged(&ty,numSamples*sizeof(double));
  cudaMallocManaged(&tz,numSamples*sizeof(double));
  cudaMallocManaged(&temp,numSamples*sizeof(double));
  cudaMallocManaged(&mass,numSamples*sizeof(double));
  cudaMallocManaged(&newtx,numSamples*sizeof(double));
  cudaMallocManaged(&newty,numSamples*sizeof(double));
  cudaMallocManaged(&newtz,numSamples*sizeof(double));
  cudaMallocManaged(&newtemp,numSamples*sizeof(double));
  cudaMallocManaged(&newmass,numSamples*sizeof(double));
  printf("Filling...\n");
  //Fill values:
  int xi=0;
  int yi=0;
  int zi=0;
  int ai=0;
  for (double alt=LAYER_HEIGHT;alt<ALTITUDE;alt+=LAYER_HEIGHT) {
    xi=0;
    for (double xc=-1;xc<=1;xc+=2/(RES-1)) {
    yi=0;
      for (double yc=-1;yc<=1;yc+=2/(RES-1)) {
        zi=0;
        for (double zc=-1;zc<=1;zc+=2/(RES-1)) {
          if (xc==-1||xc==1||yc==-1||yc==1||zc==-1||zc==1) {
            //Cube coordinates projected:
            double altAdjust=alt/sqrt(pow(xc,2)+pow(yc,2.0)+pow(zc,2.0));
            x[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=xc*altAdjust;
            y[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=yc*altAdjust;
            z[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=zc*altAdjust;
            //Trajectories:
            tx[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=((double)rand()/RAND_MAX)*(mostAllowedTrajectory-leastAllowedTrajectory)+leastAllowedTrajectory;
            ty[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=((double)rand()/RAND_MAX)*(mostAllowedTrajectory-leastAllowedTrajectory)+leastAllowedTrajectory;
            tz[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=((double)rand()/RAND_MAX)*(mostAllowedTrajectory-leastAllowedTrajectory)+leastAllowedTrajectory;
            //Temp, mass:
            mass[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=((double)rand()/RAND_MAX)*(mostAllowedMass-leastAllowedMass)+leastAllowedMass;//Too lazy to change variable names despite grammatical innacuracy.
            temp[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=((double)rand()/RAND_MAX)*(mostAllowedTemp-leastAllowedTemp)+leastAllowedTemp;
          }
          zi++;
        }
        yi++;
      }
      xi++;
    }
    ai++;
  }
  printf("Generated, copying to GPU...\n");
  //Copy Values:
  printf("Copied. Starting kernel/server loop. Connect to port 2084 in you renderer.\n");
  //"Start server thread." Not anymore, I learned Kernels cannot write to host memory and the cpu cannot pull from kernels while they are running. So I guess we do it in a loop, which is slower but at least physically possible.
  //"pthread_join(thread_id,NULL); // Wait for server (test, don't actually do)." Also irrelevant.
  //Kernels:
  //More irrelevant comments for reference:
  // cudaDeviceProp deviceProp;
  // cudaGetDeviceProperties(&deviceProp,0);//The zero is for GPU zero, as I assume you have not bought extra GPUs for this task.
  // int rate=deviceProp.clockRate;
  //Dumping/Plagiarizing some server code:
  int sockfd;
  char buffer[numSamples*32];
  char *hello = "Hello from server";
  struct sockaddr_in servaddr, cliaddr;

  // Creating socket file descriptor
  if ( (sockfd = socket(AF_INET, SOCK_DGRAM, 0)) < 0 ) {
      perror("socket creation failed");
      exit(EXIT_FAILURE);
  }
  printf("wORkiNg...\n");
  memset(&servaddr, 0, sizeof(servaddr));
  memset(&cliaddr, 0, sizeof(cliaddr));

  // Filling server information

  sockfd=socket(AF_INET,SOCK_DGRAM,0);
  if (sockfd<0) {
    perror("Error on socket()");
  }
  int optval = 1;
  setsockopt(sockfd,SOL_SOCKET,SO_REUSEADDR,(const void *)&optval,sizeof(int));
  bzero((char *)&servaddr,sizeof(servaddr));
  servaddr.sin_family    = AF_INET; // IPv4
  // inet_pton(AF_INET, GetIP().c_str(), &(servaddr.sin_addr));
  servaddr.sin_addr.s_addr = INADDR_ANY;
  servaddr.sin_port = htons(PORT);
  printf("WorKInG...\n");
  // Bind the socket with the server address
  if ( bind(sockfd, (const struct sockaddr *)&servaddr,
          sizeof(servaddr)) < 0 )
  {
      perror("bind failed");
      exit(EXIT_FAILURE);
  }
  printf("Beginning Loop\n");
  int len,n;
  len=sizeof(cliaddr);
  n = recvfrom(sockfd, (char *)buffer, MAXLINE,
                MSG_WAITALL, ( struct sockaddr *) &cliaddr,
                (socklen_t *)&len);
  buffer[n]='\0';
  printf("%s,%i\n",buffer,n);
  char * one=(char*)calloc(numSamples*32,sizeof(char));
  one[0]='1';
  char * two=(char*)calloc(numSamples*32,sizeof(char));
  two[0]='2';
  char * result=(char*)calloc((numSamples*33+1),sizeof(char));
  char * output=(char *)calloc((numSamples*32+1),sizeof(char));
  while(n>0) {
    //Use numSamples for core count, one thread each. In the future we'll want to get multiple threads for when we run out of cores.
    osmate<<<numSamples,1>>>(TEMP_RATE,MASS_RATE,TEMP_PUSH,REPULSION_RATE,numSamples,ALTITUDE,GRAV_CONST,x,y,z,tx,ty,tz,newtx,newty,newtz,mass,temp,newmass,newtemp);
    cudaDeviceSynchronize();
    // printf("Joy to the world, the gpu has finished\n");
    // printf("Synchronized.\n");
    for (int i=0;i<numSamples;i++) {
      temp[i]=newtemp[i];
      mass[i]=newmass[i];
      tx[i]=newtx[i];
      ty[i]=newty[i];
      tz[i]=newtz[i];
    }
    //Send temp,mass
    // delay(10000);
    sendto(sockfd, one, numSamples*32, MSG_DONTWAIT, (const struct sockaddr *) &cliaddr,len);
    // delay(10000);
    // printf("\n");
    jsonArray(temp,numSamples,numSamples*32,result,output);
    sendto(sockfd, result, numSamples*32, MSG_DONTWAIT, (const struct sockaddr *) &cliaddr,len);
    // printf("%s\n",result);
    // delay(10000);
    sendto(sockfd, two, numSamples*32, MSG_DONTWAIT, (const struct sockaddr *) &cliaddr,len);
    // delay(10000);
    jsonArray(mass,numSamples,numSamples*32,result,output);
    sendto(sockfd, result, numSamples*32, MSG_DONTWAIT, (const struct sockaddr *) &cliaddr,len);
  }
}

	.file	"tmpxft_00384d1c_00000000-6_sim.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2058:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2058:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0, @function
_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0:
.LFB2203:
	.cfi_startproc
	movl	$1, %edx
	jmp	cudaMallocManaged@PLT
	.cfi_endproc
.LFE2203:
	.size	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0, .-_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	.globl	_Z5delayi
	.type	_Z5delayi, @function
_Z5delayi:
.LFB2051:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movslq	%edi, %rbx
	call	clock@PLT
	addq	%rax, %rbx
.L4:
	call	clock@PLT
	cmpq	%rbx, %rax
	jl	.L4
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2051:
	.size	_Z5delayi, .-_Z5delayi
	.globl	_Z15serialize_floatPcf
	.type	_Z15serialize_floatPcf, @function
_Z15serialize_floatPcf:
.LFB2052:
	.cfi_startproc
	endbr64
	movd	%xmm0, %eax
	bswap	%eax
	movl	%eax, (%rdi)
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z15serialize_floatPcf, .-_Z15serialize_floatPcf
	.globl	_Z17doublearraytocharPdi
	.type	_Z17doublearraytocharPdi, @function
_Z17doublearraytocharPdi:
.LFB2053:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movq	%rdi, %r13
	movslq	%esi, %rdi
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rdi, %r12
	salq	$2, %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	xorl	%ebx, %ebx
	pushq	%rcx
	.cfi_def_cfa_offset 48
	call	malloc@PLT
	movq	%rax, %rbp
.L10:
	cmpl	%ebx, %r12d
	jle	.L13
	movl	$4, %edi
	call	malloc@PLT
	cvtsd2ss	0(%r13,%rbx,8), %xmm0
	movq	%rax, %rdx
	movq	%rax, %rdi
	call	_Z15serialize_floatPcf
	movl	(%rdx), %eax
	movl	%eax, 0(%rbp,%rbx,4)
	incq	%rbx
	jmp	.L10
.L13:
	popq	%rdx
	.cfi_def_cfa_offset 40
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z17doublearraytocharPdi, .-_Z17doublearraytocharPdi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"%f,"
.LC1:
	.string	"%f]"
	.text
	.globl	_Z9jsonArrayPdiiPcS0_
	.type	_Z9jsonArrayPdiiPcS0_, @function
_Z9jsonArrayPdiiPcS0_:
.LFB2054:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	addl	%edx, %edx
	leal	-1(%rsi), %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%esi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rcx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movslq	%edx, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movb	$91, (%rcx)
	movq	%r8, 8(%rsp)
.L15:
	cmpl	%r13d, %r12d
	jge	.L21
	movsd	(%r14,%r12,8), %xmm0
	leaq	.LC0(%rip), %r8
	cmpl	%r12d, %r15d
	jne	.L20
	leaq	.LC1(%rip), %r8
.L20:
	movq	8(%rsp), %rdi
	orq	$-1, %rcx
	movl	$2, %edx
	movb	$1, %al
	movq	%rbx, %rsi
	incq	%r12
	call	__snprintf_chk@PLT
	movq	8(%rsp), %rsi
	movq	%rbp, %rdi
	call	strcat@PLT
	jmp	.L15
.L21:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbp, %rax
	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
.LFE2054:
	.size	_Z9jsonArrayPdiiPcS0_, .-_Z9jsonArrayPdiiPcS0_
	.globl	_Z56__device_stub__Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_ddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.type	_Z56__device_stub__Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_ddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_, @function
_Z56__device_stub__Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_ddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_:
.LFB2080:
	.cfi_startproc
	endbr64
	subq	$408, %rsp
	.cfi_def_cfa_offset 416
	movq	416(%rsp), %rax
	movl	%edi, 124(%rsp)
	leaq	184(%rsp), %rdi
	movq	%rsi, 96(%rsp)
	leaq	196(%rsp), %rsi
	movq	%rax, 56(%rsp)
	movq	424(%rsp), %rax
	movq	%rdx, 88(%rsp)
	leaq	168(%rsp), %rdx
	movq	%rax, 48(%rsp)
	movq	432(%rsp), %rax
	movq	%rcx, 80(%rsp)
	leaq	176(%rsp), %rcx
	movq	%rax, 40(%rsp)
	movq	440(%rsp), %rax
	movq	%r8, 72(%rsp)
	movq	%rax, 32(%rsp)
	movq	448(%rsp), %rax
	movq	%r9, 64(%rsp)
	movq	%rax, 24(%rsp)
	movq	456(%rsp), %rax
	movsd	%xmm0, 152(%rsp)
	movq	%rax, 16(%rsp)
	movq	464(%rsp), %rax
	movsd	%xmm1, 144(%rsp)
	movq	%rax, 8(%rsp)
	movq	472(%rsp), %rax
	movsd	%xmm2, 136(%rsp)
	movq	%rax, (%rsp)
	movsd	%xmm3, 128(%rsp)
	movsd	%xmm4, 112(%rsp)
	movsd	%xmm5, 104(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 392(%rsp)
	xorl	%eax, %eax
	leaq	152(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	144(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	136(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	128(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	124(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 336(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 344(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 352(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 360(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 368(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 376(%rsp)
	movq	%rsp, %rax
	movq	%rax, 384(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 184(%rsp)
	movl	$1, 192(%rsp)
	movq	%rax, 196(%rsp)
	movl	$1, 204(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L22
	pushq	176(%rsp)
	.cfi_def_cfa_offset 424
	leaq	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_(%rip), %rdi
	pushq	176(%rsp)
	.cfi_def_cfa_offset 432
	movq	212(%rsp), %rcx
	movl	220(%rsp), %r8d
	movq	200(%rsp), %rsi
	movl	208(%rsp), %edx
	leaq	248(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 424
	popq	%rdx
	.cfi_def_cfa_offset 416
.L22:
	movq	392(%rsp), %rax
	subq	%fs:40, %rax
	je	.L24
	call	__stack_chk_fail@PLT
.L24:
	addq	$408, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2080:
	.size	_Z56__device_stub__Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_ddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_, .-_Z56__device_stub__Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_ddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.globl	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.type	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_, @function
_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_:
.LFB2081:
	.cfi_startproc
	endbr64
	jmp	_Z56__device_stub__Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_ddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.cfi_endproc
.LFE2081:
	.size	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_, .-_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.section	.rodata.str1.1
.LC3:
	.string	"Hello, you are hopefully simulating tectonics. If not, get out of here.\n"
.LC5:
	.string	"Allocating for %i samples (pass to your renderer your RES, %i, and your ALTITUDE, %i, like so: python display.py <RES> <ALTITUDE>)...\n"
.LC6:
	.string	"Filling...\n"
.LC9:
	.string	"Generated, copying to GPU...\n"
.LC10:
	.string	"Copied. Starting kernel/server loop. Connect to port 2084 in you renderer.\n"
.LC11:
	.string	"socket creation failed"
.LC12:
	.string	"wORkiNg...\n"
.LC13:
	.string	"Error on socket()"
.LC14:
	.string	"WorKInG...\n"
.LC15:
	.string	"bind failed"
.LC16:
	.string	"Beginning Loop\n"
.LC17:
	.string	"%s,%i\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2055:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	xorl	%edi, %edi
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	.cfi_offset 12, -48
	.cfi_offset 3, -56
	xorl	%ebx, %ebx
	subq	$232, %rsp
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movsd	LAYER_HEIGHT(%rip), %xmm3
	movaps	%xmm3, %xmm2
.L28:
	movl	ALTITUDE(%rip), %r8d
	movl	RES(%rip), %r12d
	cvtsi2sdl	%r8d, %xmm0
	comisd	%xmm2, %xmm0
	jbe	.L66
	cvtsi2sdl	%r12d, %xmm0
	movsd	.LC4(%rip), %xmm1
	subl	$2, %r12d
	movsd	%xmm3, -240(%rbp)
	movsd	%xmm2, -232(%rbp)
	call	pow@PLT
	movsd	.LC4(%rip), %xmm1
	movq	%xmm0, %r13
	cvtsi2sdl	%r12d, %xmm0
	call	pow@PLT
	movsd	-232(%rbp), %xmm2
	movsd	-240(%rbp), %xmm3
	movaps	%xmm0, %xmm1
	movq	%r13, %xmm0
	subsd	%xmm1, %xmm0
	addsd	%xmm3, %xmm2
	cvttsd2sil	%xmm0, %eax
	addl	%eax, %ebx
	jmp	.L28
.L66:
	movl	%r12d, %ecx
	movslq	%ebx, %r12
	movl	%ebx, %edx
	movl	$2, %edi
	salq	$3, %r12
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r12, %rsi
	leaq	-216(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-208(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-200(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-192(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-184(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-176(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-168(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-160(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-152(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-144(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-136(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-128(%rbp), %rdi
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	%r12, %rsi
	leaq	-120(%rbp), %rdi
	xorl	%r12d, %r12d
	call	_Z17cudaMallocManagedIdE9cudaErrorPPT_mj.constprop.0.isra.0
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movsd	LAYER_HEIGHT(%rip), %xmm6
	movsd	.LC7(%rip), %xmm4
.L30:
	cvtsi2sdl	ALTITUDE(%rip), %xmm0
	comisd	%xmm6, %xmm0
	jbe	.L67
	movsd	.LC2(%rip), %xmm1
	xorl	%r14d, %r14d
	movaps	%xmm1, %xmm5
.L40:
	movaps	%xmm5, %xmm2
	xorl	%r13d, %r13d
.L39:
	movaps	%xmm5, %xmm3
	xorl	%r15d, %r15d
.L34:
	ucomisd	%xmm5, %xmm1
	movb	$0, %dl
	setnp	%al
	cmovne	%edx, %eax
	testb	%al, %al
	jne	.L31
	ucomisd	%xmm4, %xmm1
	setnp	%dl
	cmovne	%eax, %edx
	testb	%dl, %dl
	jne	.L31
	ucomisd	%xmm5, %xmm2
	setnp	%al
	cmovne	%edx, %eax
	testb	%al, %al
	jne	.L31
	ucomisd	%xmm4, %xmm2
	setnp	%dl
	cmovne	%eax, %edx
	testb	%dl, %dl
	jne	.L31
	ucomisd	%xmm5, %xmm3
	setnp	%al
	cmovne	%edx, %eax
	testb	%al, %al
	jne	.L31
	ucomisd	%xmm4, %xmm3
	setnp	%dl
	cmove	%edx, %eax
	testb	%al, %al
	je	.L32
.L31:
	movaps	%xmm2, %xmm7
	movaps	%xmm1, %xmm0
	movsd	%xmm6, -248(%rbp)
	mulsd	%xmm2, %xmm7
	movsd	%xmm2, -264(%rbp)
	mulsd	%xmm1, %xmm0
	movsd	%xmm1, -240(%rbp)
	movsd	%xmm3, -232(%rbp)
	addsd	%xmm7, %xmm0
	movaps	%xmm3, %xmm7
	mulsd	%xmm3, %xmm7
	addsd	%xmm7, %xmm0
	call	sqrt@PLT
	movsd	-248(%rbp), %xmm6
	movl	RES(%rip), %edx
	movaps	%xmm0, %xmm7
	movsd	-240(%rbp), %xmm1
	movsd	-264(%rbp), %xmm2
	movaps	%xmm6, %xmm0
	movl	%edx, %eax
	movsd	-232(%rbp), %xmm3
	movsd	%xmm6, -256(%rbp)
	divsd	%xmm7, %xmm0
	imull	%r14d, %eax
	movsd	%xmm1, -248(%rbp)
	movsd	%xmm2, -240(%rbp)
	addl	%r13d, %eax
	imull	%edx, %eax
	addl	%r15d, %eax
	imull	%edx, %eax
	movq	-216(%rbp), %rdx
	addl	%r12d, %eax
	cltq
	movaps	%xmm0, %xmm7
	mulsd	%xmm1, %xmm7
	movsd	%xmm7, (%rdx,%rax,8)
	movaps	%xmm0, %xmm7
	movq	-208(%rbp), %rdx
	mulsd	%xmm2, %xmm7
	mulsd	%xmm3, %xmm0
	movsd	%xmm7, (%rdx,%rax,8)
	movq	-200(%rbp), %rdx
	movsd	%xmm0, (%rdx,%rax,8)
	call	rand@PLT
	movl	RES(%rip), %edx
	movsd	leastAllowedTrajectory(%rip), %xmm8
	movl	%eax, %ecx
	movsd	mostAllowedTrajectory(%rip), %xmm7
	movl	%edx, %eax
	cvtsi2sdl	%ecx, %xmm0
	divsd	.LC8(%rip), %xmm0
	imull	%r14d, %eax
	subsd	%xmm8, %xmm7
	mulsd	%xmm7, %xmm0
	addl	%r13d, %eax
	imull	%edx, %eax
	addl	%r15d, %eax
	imull	%edx, %eax
	addsd	%xmm8, %xmm0
	movq	-192(%rbp), %rdx
	addl	%r12d, %eax
	cltq
	movsd	%xmm0, (%rdx,%rax,8)
	call	rand@PLT
	movl	RES(%rip), %edx
	movsd	leastAllowedTrajectory(%rip), %xmm8
	movl	%eax, %ecx
	movsd	mostAllowedTrajectory(%rip), %xmm7
	movl	%edx, %eax
	cvtsi2sdl	%ecx, %xmm0
	divsd	.LC8(%rip), %xmm0
	imull	%r14d, %eax
	subsd	%xmm8, %xmm7
	mulsd	%xmm7, %xmm0
	addl	%r13d, %eax
	imull	%edx, %eax
	addl	%r15d, %eax
	imull	%edx, %eax
	addsd	%xmm8, %xmm0
	movq	-184(%rbp), %rdx
	addl	%r12d, %eax
	cltq
	movsd	%xmm0, (%rdx,%rax,8)
	call	rand@PLT
	movl	RES(%rip), %edx
	movsd	leastAllowedTrajectory(%rip), %xmm8
	movl	%eax, %ecx
	movsd	mostAllowedTrajectory(%rip), %xmm7
	movl	%edx, %eax
	cvtsi2sdl	%ecx, %xmm0
	divsd	.LC8(%rip), %xmm0
	imull	%r14d, %eax
	subsd	%xmm8, %xmm7
	mulsd	%xmm7, %xmm0
	addl	%r13d, %eax
	imull	%edx, %eax
	addl	%r15d, %eax
	imull	%edx, %eax
	addsd	%xmm8, %xmm0
	movq	-176(%rbp), %rdx
	addl	%r12d, %eax
	cltq
	movsd	%xmm0, (%rdx,%rax,8)
	call	rand@PLT
	movl	RES(%rip), %edx
	movsd	leastAllowedMass(%rip), %xmm8
	movl	%eax, %ecx
	movsd	mostAllowedMass(%rip), %xmm7
	movl	%edx, %eax
	cvtsi2sdl	%ecx, %xmm0
	divsd	.LC8(%rip), %xmm0
	imull	%r14d, %eax
	subsd	%xmm8, %xmm7
	mulsd	%xmm7, %xmm0
	addl	%r13d, %eax
	imull	%edx, %eax
	addl	%r15d, %eax
	imull	%edx, %eax
	addsd	%xmm8, %xmm0
	movq	-160(%rbp), %rdx
	addl	%r12d, %eax
	cltq
	movsd	%xmm0, (%rdx,%rax,8)
	call	rand@PLT
	movl	RES(%rip), %edx
	movsd	leastAllowedTemp(%rip), %xmm8
	movl	%eax, %ecx
	movsd	mostAllowedTemp(%rip), %xmm7
	movsd	.LC7(%rip), %xmm4
	movl	%edx, %eax
	cvtsi2sdl	%ecx, %xmm0
	divsd	.LC8(%rip), %xmm0
	movsd	.LC2(%rip), %xmm5
	imull	%r14d, %eax
	subsd	%xmm8, %xmm7
	movsd	-256(%rbp), %xmm6
	movsd	-248(%rbp), %xmm1
	movsd	-240(%rbp), %xmm2
	movsd	-232(%rbp), %xmm3
	mulsd	%xmm7, %xmm0
	addl	%r13d, %eax
	imull	%edx, %eax
	addl	%r15d, %eax
	imull	%edx, %eax
	addsd	%xmm8, %xmm0
	movq	-168(%rbp), %rdx
	addl	%r12d, %eax
	cltq
	movsd	%xmm0, (%rdx,%rax,8)
.L32:
	movl	RES(%rip), %eax
	incl	%r15d
	leal	-1(%rax), %ecx
	movl	$2, %eax
	cltd
	idivl	%ecx
	cvtsi2sdl	%eax, %xmm0
	addsd	%xmm0, %xmm3
	comisd	%xmm3, %xmm4
	jnb	.L34
	addsd	%xmm0, %xmm2
	incl	%r13d
	comisd	%xmm2, %xmm4
	jnb	.L39
	addsd	%xmm0, %xmm1
	incl	%r14d
	comisd	%xmm1, %xmm4
	jnb	.L40
	addsd	LAYER_HEIGHT(%rip), %xmm6
	incl	%r12d
	jmp	.L30
.L67:
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	%ebx, %eax
	movq	%rsp, %rdx
	sall	$5, %eax
	movl	%eax, -248(%rbp)
	cltq
	movq	%rax, %rcx
	movq	%rax, -232(%rbp)
	andq	$-4096, %rcx
	subq	%rcx, %rdx
.L41:
	cmpq	%rdx, %rsp
	je	.L42
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L41
.L42:
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L43
	orq	$0, -8(%rsp,%rax)
.L43:
	movl	$2, %edi
	xorl	%edx, %edx
	movl	$2, %esi
	movq	%rsp, %r12
	call	socket@PLT
	leaq	.LC11(%rip), %rdi
	testl	%eax, %eax
	js	.L65
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorps	%xmm0, %xmm0
	xorl	%edx, %edx
	movl	$2, %esi
	movl	$2, %edi
	movups	%xmm0, -88(%rbp)
	movups	%xmm0, -72(%rbp)
	call	socket@PLT
	movl	%eax, -240(%rbp)
	testl	%eax, %eax
	jns	.L45
	leaq	.LC13(%rip), %rdi
	call	perror@PLT
.L45:
	movl	-240(%rbp), %edi
	movl	$2, %edx
	movl	$1, %esi
	leaq	-224(%rbp), %rcx
	movl	$4, %r8d
	leaq	-88(%rbp), %r13
	movl	$1, -224(%rbp)
	call	setsockopt@PLT
	xorl	%edi, %edi
	xorl	%r8d, %r8d
	xorl	%eax, %eax
	movq	%rdi, -84(%rbp)
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	movl	%r8d, -76(%rbp)
	movl	$604504066, -88(%rbp)
	call	__printf_chk@PLT
	movl	-240(%rbp), %edi
	movl	$16, %edx
	movq	%r13, %rsi
	call	bind@PLT
	testl	%eax, %eax
	jns	.L46
	leaq	.LC15(%rip), %rdi
.L65:
	call	perror@PLT
	movl	$1, %edi
	call	exit@PLT
.L46:
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	-72(%rbp), %r9
	movl	$256, %r8d
	movl	$16, -220(%rbp)
	pushq	%rax
	leaq	-220(%rbp), %rax
	movq	%r12, %rsi
	movl	$1024, %edx
	movq	-232(%rbp), %rcx
	movl	-240(%rbp), %edi
	pushq	%rax
	call	__recvfrom_chk@PLT
	movl	$2, %edi
	movl	%eax, %ecx
	movq	%rax, %r13
	cltq
	movb	$0, (%r12,%rax)
	xorl	%eax, %eax
	popq	%rdx
	movq	%r12, %rdx
	popq	%rsi
	leaq	.LC17(%rip), %rsi
	call	__printf_chk@PLT
	movq	-232(%rbp), %rdi
	movl	$1, %esi
	call	calloc@PLT
	movq	-232(%rbp), %rdi
	movl	$1, %esi
	movb	$49, (%rax)
	movq	%rax, %r12
	call	calloc@PLT
	imull	$33, %ebx, %edi
	movl	$1, %esi
	movb	$50, (%rax)
	movq	%rax, -264(%rbp)
	incl	%edi
	movslq	%edi, %rdi
	call	calloc@PLT
	movl	$1, %esi
	movq	%rax, -256(%rbp)
	movl	-248(%rbp), %eax
	leal	1(%rax), %edi
	movslq	%edi, %rdi
	call	calloc@PLT
	movq	%rax, -272(%rbp)
	testl	%r13d, %r13d
	jle	.L47
	movabsq	$4294967297, %r13
.L51:
	movq	%r13, -108(%rbp)
	movl	-104(%rbp), %esi
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	%ebx, -112(%rbp)
	movq	-112(%rbp), %rdi
	movq	%r13, %rdx
	movl	$1, %ecx
	movq	%r13, -100(%rbp)
	movl	$1, -92(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L48
	cvtsi2sdl	ALTITUDE(%rip), %xmm4
	pushq	-128(%rbp)
	movl	%ebx, %edi
	movsd	GRAV_CONST(%rip), %xmm5
	pushq	-120(%rbp)
	movsd	REPULSION_RATE(%rip), %xmm3
	pushq	-168(%rbp)
	pushq	-160(%rbp)
	pushq	-136(%rbp)
	pushq	-144(%rbp)
	pushq	-152(%rbp)
	pushq	-176(%rbp)
	movq	-184(%rbp), %r9
	movq	-192(%rbp), %r8
	movq	-200(%rbp), %rcx
	movq	-208(%rbp), %rdx
	movq	-216(%rbp), %rsi
	movsd	TEMP_PUSH(%rip), %xmm2
	movsd	MASS_RATE(%rip), %xmm1
	movsd	TEMP_RATE(%rip), %xmm0
	call	_Z56__device_stub__Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_ddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	addq	$64, %rsp
.L48:
	call	cudaDeviceSynchronize@PLT
	movq	-128(%rbp), %rdx
	movq	-120(%rbp), %rsi
	xorl	%eax, %eax
	movq	-168(%rbp), %rcx
	movq	-160(%rbp), %rdi
	movq	-152(%rbp), %r8
	movq	-192(%rbp), %r9
	movq	-144(%rbp), %r10
	movq	-184(%rbp), %r11
	movq	-136(%rbp), %r14
	movq	-176(%rbp), %r15
.L49:
	cmpl	%eax, %ebx
	jle	.L68
	movsd	(%rdx,%rax,8), %xmm0
	movsd	%xmm0, (%rcx,%rax,8)
	movsd	(%rsi,%rax,8), %xmm0
	movsd	%xmm0, (%rdi,%rax,8)
	movsd	(%r8,%rax,8), %xmm0
	movsd	%xmm0, (%r9,%rax,8)
	movsd	(%r10,%rax,8), %xmm0
	movsd	%xmm0, (%r11,%rax,8)
	movsd	(%r14,%rax,8), %xmm0
	movsd	%xmm0, (%r15,%rax,8)
	incq	%rax
	jmp	.L49
.L68:
	leaq	-72(%rbp), %r14
	movl	-220(%rbp), %r9d
	movq	-232(%rbp), %rdx
	movq	%r12, %rsi
	movl	-240(%rbp), %edi
	movq	%r14, %r8
	movl	$64, %ecx
	call	sendto@PLT
	movq	-272(%rbp), %r8
	movl	%ebx, %esi
	movq	-256(%rbp), %rcx
	movl	-248(%rbp), %edx
	movq	-168(%rbp), %rdi
	call	_Z9jsonArrayPdiiPcS0_
	movl	-220(%rbp), %r9d
	movq	%r14, %r8
	movq	-232(%rbp), %rdx
	movq	-256(%rbp), %rsi
	movl	-240(%rbp), %edi
	movl	$64, %ecx
	call	sendto@PLT
	movl	-220(%rbp), %r9d
	movq	%r14, %r8
	movq	-232(%rbp), %rdx
	movq	-264(%rbp), %rsi
	movl	-240(%rbp), %edi
	movl	$64, %ecx
	call	sendto@PLT
	movq	-272(%rbp), %r8
	movl	%ebx, %esi
	movq	-256(%rbp), %rcx
	movl	-248(%rbp), %edx
	movq	-160(%rbp), %rdi
	call	_Z9jsonArrayPdiiPcS0_
	movl	-220(%rbp), %r9d
	movq	%r14, %r8
	movq	-232(%rbp), %rdx
	movq	-256(%rbp), %rsi
	movl	-240(%rbp), %edi
	movl	$64, %ecx
	call	sendto@PLT
	jmp	.L51
.L47:
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L52
	call	__stack_chk_fail@PLT
.L52:
	leaq	-40(%rbp), %rsp
	xorl	%eax, %eax
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE2055:
	.size	main, .-main
	.section	.rodata.str1.1
.LC19:
	.string	"_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2083:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC19(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2083:
	.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
	.globl	mostAllowedTrajectory
	.data
	.align 8
	.type	mostAllowedTrajectory, @object
	.size	mostAllowedTrajectory, 8
mostAllowedTrajectory:
	.long	0
	.long	1073741824
	.globl	leastAllowedTrajectory
	.align 8
	.type	leastAllowedTrajectory, @object
	.size	leastAllowedTrajectory, 8
leastAllowedTrajectory:
	.long	-350469331
	.long	1058682594
	.globl	mostAllowedTemp
	.align 8
	.type	mostAllowedTemp, @object
	.size	mostAllowedTemp, 8
mostAllowedTemp:
	.long	0
	.long	1088973312
	.globl	leastAllowedTemp
	.align 8
	.type	leastAllowedTemp, @object
	.size	leastAllowedTemp, 8
leastAllowedTemp:
	.long	-755914244
	.long	1062232653
	.globl	mostAllowedMass
	.align 8
	.type	mostAllowedMass, @object
	.size	mostAllowedMass, 8
mostAllowedMass:
	.long	0
	.long	1075052544
	.globl	leastAllowedMass
	.align 8
	.type	leastAllowedMass, @object
	.size	leastAllowedMass, 8
leastAllowedMass:
	.long	-755914244
	.long	1062232653
	.globl	RES
	.align 4
	.type	RES, @object
	.size	RES, 4
RES:
	.long	3
	.globl	LAYER_HEIGHT
	.align 8
	.type	LAYER_HEIGHT, @object
	.size	LAYER_HEIGHT, 8
LAYER_HEIGHT:
	.long	0
	.long	1071644672
	.globl	ALTITUDE
	.align 4
	.type	ALTITUDE, @object
	.size	ALTITUDE, 4
ALTITUDE:
	.long	5
	.globl	GRAV_CONST
	.align 8
	.type	GRAV_CONST, @object
	.size	GRAV_CONST, 8
GRAV_CONST:
	.long	1319413953
	.long	1068569408
	.globl	SAMPLE_RADIUS
	.align 8
	.type	SAMPLE_RADIUS, @object
	.size	SAMPLE_RADIUS, 8
SAMPLE_RADIUS:
	.long	-1717986918
	.long	1069128089
	.globl	REPULSION_RATE
	.align 8
	.type	REPULSION_RATE, @object
	.size	REPULSION_RATE, 8
REPULSION_RATE:
	.long	0
	.long	1072693248
	.globl	TEMP_PUSH
	.align 8
	.type	TEMP_PUSH, @object
	.size	TEMP_PUSH, 8
TEMP_PUSH:
	.long	-343597384
	.long	1067366481
	.globl	MASS_RATE
	.align 8
	.type	MASS_RATE, @object
	.size	MASS_RATE, 8
MASS_RATE:
	.long	-1717986918
	.long	1071225241
	.globl	TEMP_RATE
	.align 8
	.type	TEMP_RATE, @object
	.size	TEMP_RATE, 8
TEMP_RATE:
	.long	-1717986918
	.long	1069128089
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC2:
	.long	0
	.long	-1074790400
	.align 8
.LC4:
	.long	0
	.long	1074266112
	.align 8
.LC7:
	.long	0
	.long	1072693248
	.align 8
.LC8:
	.long	-4194304
	.long	1105199103
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
#include <stdio.h>
#include "math.h"
#include <pthread.h>
#include <unistd.h>
#include <string.h>
#include <sys/types.h>
#include <sys/socket.h>
#include <arpa/inet.h>
#include <netinet/in.h>
#define PORT    2084
#define MAXLINE 1024
/*
I will provide a word about tx,ty, and tz.
tx does not, in fact move the actual point (points are static), but defines how much mass will be donated.
After this donation, a density-based heat donation is calculated, but regardless of these trajectory values.
*/
__global__
void osmate(double tempRate, double massRate, double tempPush, double repulsionRate, int numPoints, double radius, double gConst, double * x, double * y, double * z, double * tx, double * ty, double * tz, double * newtx, double * newty, double * newtz, double * mass, double * temp, double * newmass, double * newtemp) {
  //The only part that really matters if I get the rest to work:
  int i=blockIdx.x + blockDim.x + threadIdx.x;
  for (int it=0;it<numPoints;it++) {//Iterate through all points
    // printf("%f",tx[it]);
    double dist=sqrt((x[it]-x[i])*(x[it]-x[i])+(y[it]-y[i])*(y[it]-y[i])+(z[it]-z[i])*(z[it]-z[i]));//This could be pre-calculated (these points are static). Whether the calculation happens should be a matter of how many points there are -- if there are too many points we will run out of memory, but if there are not as many it is a good idea. I'm leaving it for another day.
    if (dist!=0) {
      //Congratulations, you are not looking at the same point.
      double toChange=gConst*mass[i]*mass[it]/dist;
      newtx[i]=tx[i]+(toChange+tempPush*(temp[i]-temp[it])+repulsionRate*(mass[i]-mass[it]))*(x[it]-x[i]);
      newty[i]=ty[i]+(toChange+tempPush*(temp[i]-temp[it])+repulsionRate*(mass[i]-mass[it]))*(y[it]-y[i]);
      newtz[i]=tz[i]+(toChange+tempPush*(temp[i]-temp[it])+repulsionRate*(mass[i]-mass[it]))*(z[it]-z[i]);
    }
  }
  // if (x[i]!=0) {
  //   printf("%d\n",x[i]);
  // }
  for (int it=0;it<numPoints;it++) {
    //TODO make negative mass impossible.
    if (tx[i]!=0&&ty[i]!=0&&tz[i]!=0) {
      double massDisagreement=sqrt((x[it]-x[i])*(x[it]-x[i])/tx[i]+(y[it]-y[i])*(y[it]-y[i])/ty[i]+(z[it]-z[i])*(z[it]-z[i])/tz[i]);//TODO cases for rare 0 trajectory scenarios that would knacker it all.
      if (massDisagreement!=0) {
        newmass[it]=mass[it]+massRate/massDisagreement;//Hell if I know whether this is right, but it seems it should intuitively work. Also grants some cohesion. Note we don't subtract from the subject sample because the mass comes from behind it.
        double dist=sqrt((x[it]-x[i])*(x[it]-x[i])+(y[it]-y[i])*(y[it]-y[i])+(z[it]-z[i])*(z[it]-z[i]));//There are a lot of redundant calculations here...
        if (dist!=0) {
          double tempEx=tempRate*(temp[i]-temp[it])/dist;
          newtemp[it]=temp[it]+tempEx;
          newtemp[i]=temp[i]-tempEx;//Unless we want to go supernova. We should also have a mechanic where a mass increase "creates" heat, a decrease sucks it up, there is an external source of heat and heat can radiate away... Etc.
          // printf("%f",tx[i]);
        }
      }
    }
  }
}
//To tweak
double TEMP_RATE=0.1;
double MASS_RATE=0.4;
double TEMP_PUSH=0.03;
double REPULSION_RATE=1.0;
double SAMPLE_RADIUS=.1;
double GRAV_CONST=.0667;//Definitely not.

int ALTITUDE=5;//Radius of planet.
double LAYER_HEIGHT=.5;
int RES = 3;

//In generation, values may exceed maximums and minimums during simulation.
double leastAllowedMass=.001;
double mostAllowedMass=5;//We'll randomize between the two.

double leastAllowedTemp=.001;//Kelvin?
double mostAllowedTemp=50000;//A lot, hopefully not too much. Should check Planck heat.

double leastAllowedTrajectory=.0001;//Along axis, I'm not a madman
double mostAllowedTrajectory=2;
void delay(int number_of_millis)
{
    // Converting time into milli_seconds
    // int milli_seconds = 1000 * number_of_seconds;

    // Storing start time
    clock_t start_time = clock();

    // looping till required time is not acheived
    while (clock() < start_time + number_of_millis)
        ;
}
void serialize_float(char *buffer, float value)
{
    unsigned int ivalue = *((unsigned int*)&value); // warning assumes 32-bit "unsigned int"
    buffer[0] = ivalue >> 24;
    buffer[1] = ivalue >> 16;
    buffer[2] = ivalue >> 8;
    buffer[3] = ivalue;
}
char* doublearraytochar(double * vals, int len) {
  char * result=(char *)malloc(len*sizeof(char)*4);
  for (int i=0;i<len;i++) {
    char * tempArr=(char *)malloc(sizeof(char)*4);
    serialize_float(tempArr,vals[i]);
    for (int i2=0;i2<4;i2++) {
      result[i*4+i2]=tempArr[i2];
    }
  }
  return result;
}
char * jsonArray(double * vals, int len, int bufPerNum, char * result, char * output) {
  result[0]='[';
  for (int i=0;i<len;i++) {
    if (i!=len-1) {
      snprintf(output,bufPerNum*2,"%f,",vals[i]);
    } else {
      snprintf(output,bufPerNum*2,"%f]",vals[i]);
    }
    strcat(result,output);
    // for (int i2=0;i2<(bufPerNum+1)*sizeof(char);i2++) {
    //   result[i*bufPerNum+i2+1]=output[i2];
    // }
  }
  return result;
}
int main() {
  srand ( time ( NULL));
  printf("Hello, you are hopefully simulating tectonics. If not, get out of here.\n");
  //Generate points...
  //The Declaration of Variable Allocation:
  int numSamples=0;
  for (double alt=LAYER_HEIGHT;alt<ALTITUDE;alt+=LAYER_HEIGHT) {
    //You may notice we project a cube into a sphere. Primarily I'm too lazy to generate something like an icosahedron, secondarily it would not have fine resolution (only subdivisions) if we used an icosahedron, tertiarily this allows us to check for seams, which are screaming alarms something is wrong, and fourth uv spheres get ridiculously fine detail only at the poles, which is a waste.
    numSamples+=(int)(pow(RES,3)-pow(RES-2,3));
  }
  printf("Allocating for %i samples (pass to your renderer your RES, %i, and your ALTITUDE, %i, like so: python display.py <RES> <ALTITUDE>)...\n", numSamples, RES,ALTITUDE);
  double * x;
  double * y;
  double * z;
  double * tx;
  double * ty;
  double * tz;
  double * temp;
  double * mass;
  double * newtx;
  double * newty;
  double * newtz;
  double * newtemp;
  double * newmass;
  hipMallocManaged(&x,numSamples*sizeof(double));
  hipMallocManaged(&y,numSamples*sizeof(double));
  hipMallocManaged(&z,numSamples*sizeof(double));
  hipMallocManaged(&tx,numSamples*sizeof(double));
  hipMallocManaged(&ty,numSamples*sizeof(double));
  hipMallocManaged(&tz,numSamples*sizeof(double));
  hipMallocManaged(&temp,numSamples*sizeof(double));
  hipMallocManaged(&mass,numSamples*sizeof(double));
  hipMallocManaged(&newtx,numSamples*sizeof(double));
  hipMallocManaged(&newty,numSamples*sizeof(double));
  hipMallocManaged(&newtz,numSamples*sizeof(double));
  hipMallocManaged(&newtemp,numSamples*sizeof(double));
  hipMallocManaged(&newmass,numSamples*sizeof(double));
  printf("Filling...\n");
  //Fill values:
  int xi=0;
  int yi=0;
  int zi=0;
  int ai=0;
  for (double alt=LAYER_HEIGHT;alt<ALTITUDE;alt+=LAYER_HEIGHT) {
    xi=0;
    for (double xc=-1;xc<=1;xc+=2/(RES-1)) {
    yi=0;
      for (double yc=-1;yc<=1;yc+=2/(RES-1)) {
        zi=0;
        for (double zc=-1;zc<=1;zc+=2/(RES-1)) {
          if (xc==-1||xc==1||yc==-1||yc==1||zc==-1||zc==1) {
            //Cube coordinates projected:
            double altAdjust=alt/sqrt(pow(xc,2)+pow(yc,2.0)+pow(zc,2.0));
            x[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=xc*altAdjust;
            y[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=yc*altAdjust;
            z[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=zc*altAdjust;
            //Trajectories:
            tx[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=((double)rand()/RAND_MAX)*(mostAllowedTrajectory-leastAllowedTrajectory)+leastAllowedTrajectory;
            ty[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=((double)rand()/RAND_MAX)*(mostAllowedTrajectory-leastAllowedTrajectory)+leastAllowedTrajectory;
            tz[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=((double)rand()/RAND_MAX)*(mostAllowedTrajectory-leastAllowedTrajectory)+leastAllowedTrajectory;
            //Temp, mass:
            mass[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=((double)rand()/RAND_MAX)*(mostAllowedMass-leastAllowedMass)+leastAllowedMass;//Too lazy to change variable names despite grammatical innacuracy.
            temp[xi*RES*RES*RES+yi*RES*RES+zi*RES+ai]=((double)rand()/RAND_MAX)*(mostAllowedTemp-leastAllowedTemp)+leastAllowedTemp;
          }
          zi++;
        }
        yi++;
      }
      xi++;
    }
    ai++;
  }
  printf("Generated, copying to GPU...\n");
  //Copy Values:
  printf("Copied. Starting kernel/server loop. Connect to port 2084 in you renderer.\n");
  //"Start server thread." Not anymore, I learned Kernels cannot write to host memory and the cpu cannot pull from kernels while they are running. So I guess we do it in a loop, which is slower but at least physically possible.
  //"pthread_join(thread_id,NULL); // Wait for server (test, don't actually do)." Also irrelevant.
  //Kernels:
  //More irrelevant comments for reference:
  // cudaDeviceProp deviceProp;
  // cudaGetDeviceProperties(&deviceProp,0);//The zero is for GPU zero, as I assume you have not bought extra GPUs for this task.
  // int rate=deviceProp.clockRate;
  //Dumping/Plagiarizing some server code:
  int sockfd;
  char buffer[numSamples*32];
  char *hello = "Hello from server";
  struct sockaddr_in servaddr, cliaddr;

  // Creating socket file descriptor
  if ( (sockfd = socket(AF_INET, SOCK_DGRAM, 0)) < 0 ) {
      perror("socket creation failed");
      exit(EXIT_FAILURE);
  }
  printf("wORkiNg...\n");
  memset(&servaddr, 0, sizeof(servaddr));
  memset(&cliaddr, 0, sizeof(cliaddr));

  // Filling server information

  sockfd=socket(AF_INET,SOCK_DGRAM,0);
  if (sockfd<0) {
    perror("Error on socket()");
  }
  int optval = 1;
  setsockopt(sockfd,SOL_SOCKET,SO_REUSEADDR,(const void *)&optval,sizeof(int));
  bzero((char *)&servaddr,sizeof(servaddr));
  servaddr.sin_family    = AF_INET; // IPv4
  // inet_pton(AF_INET, GetIP().c_str(), &(servaddr.sin_addr));
  servaddr.sin_addr.s_addr = INADDR_ANY;
  servaddr.sin_port = htons(PORT);
  printf("WorKInG...\n");
  // Bind the socket with the server address
  if ( bind(sockfd, (const struct sockaddr *)&servaddr,
          sizeof(servaddr)) < 0 )
  {
      perror("bind failed");
      exit(EXIT_FAILURE);
  }
  printf("Beginning Loop\n");
  int len,n;
  len=sizeof(cliaddr);
  n = recvfrom(sockfd, (char *)buffer, MAXLINE,
                MSG_WAITALL, ( struct sockaddr *) &cliaddr,
                (socklen_t *)&len);
  buffer[n]='\0';
  printf("%s,%i\n",buffer,n);
  char * one=(char*)calloc(numSamples*32,sizeof(char));
  one[0]='1';
  char * two=(char*)calloc(numSamples*32,sizeof(char));
  two[0]='2';
  char * result=(char*)calloc((numSamples*33+1),sizeof(char));
  char * output=(char *)calloc((numSamples*32+1),sizeof(char));
  while(n>0) {
    //Use numSamples for core count, one thread each. In the future we'll want to get multiple threads for when we run out of cores.
    osmate<<<numSamples,1>>>(TEMP_RATE,MASS_RATE,TEMP_PUSH,REPULSION_RATE,numSamples,ALTITUDE,GRAV_CONST,x,y,z,tx,ty,tz,newtx,newty,newtz,mass,temp,newmass,newtemp);
    hipDeviceSynchronize();
    // printf("Joy to the world, the gpu has finished\n");
    // printf("Synchronized.\n");
    for (int i=0;i<numSamples;i++) {
      temp[i]=newtemp[i];
      mass[i]=newmass[i];
      tx[i]=newtx[i];
      ty[i]=newty[i];
      tz[i]=newtz[i];
    }
    //Send temp,mass
    // delay(10000);
    sendto(sockfd, one, numSamples*32, MSG_DONTWAIT, (const struct sockaddr *) &cliaddr,len);
    // delay(10000);
    // printf("\n");
    jsonArray(temp,numSamples,numSamples*32,result,output);
    sendto(sockfd, result, numSamples*32, MSG_DONTWAIT, (const struct sockaddr *) &cliaddr,len);
    // printf("%s\n",result);
    // delay(10000);
    sendto(sockfd, two, numSamples*32, MSG_DONTWAIT, (const struct sockaddr *) &cliaddr,len);
    // delay(10000);
    jsonArray(mass,numSamples,numSamples*32,result,output);
    sendto(sockfd, result, numSamples*32, MSG_DONTWAIT, (const struct sockaddr *) &cliaddr,len);
  }
}

	.text
	.file	"sim.hip"
	.globl	_Z21__device_stub__osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_ # -- Begin function _Z21__device_stub__osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.type	_Z21__device_stub__osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_,@function
_Z21__device_stub__osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_: # @_Z21__device_stub__osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$304, %rsp                      # imm = 0x130
	.cfi_def_cfa_offset 352
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	104(%rsp), %rax
	movsd	%xmm0, (%rax)
	leaq	96(%rsp), %r10
	movsd	%xmm1, (%r10)
	leaq	88(%rsp), %r11
	movsd	%xmm2, (%r11)
	leaq	80(%rsp), %r14
	movsd	%xmm3, (%r14)
	leaq	4(%rsp), %r15
	movl	%edi, (%r15)
	leaq	72(%rsp), %rdi
	movsd	%xmm4, (%rdi)
	leaq	64(%rsp), %r12
	movsd	%xmm5, (%r12)
	leaq	56(%rsp), %r13
	movq	%rsi, (%r13)
	leaq	48(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	32(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	24(%rsp), %r8
	movq	%r9, (%r8)
	leaq	144(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%r10, 8(%rbx)
	movq	%r11, 16(%rbx)
	movq	%r14, 24(%rbx)
	movq	%r15, 32(%rbx)
	movq	%rdi, 40(%rbx)
	movq	%r12, 48(%rbx)
	movq	%r13, 56(%rbx)
	movq	%rsi, 64(%rbx)
	movq	%rdx, 72(%rbx)
	movq	%rcx, 80(%rbx)
	movq	%r8, 88(%rbx)
	leaq	352(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	360(%rsp), %rax
	movq	%rax, 104(%rbx)
	leaq	368(%rsp), %rax
	movq	%rax, 112(%rbx)
	leaq	376(%rsp), %rax
	movq	%rax, 120(%rbx)
	leaq	384(%rsp), %rax
	movq	%rax, 128(%rbx)
	leaq	392(%rsp), %rax
	movq	%rax, 136(%rbx)
	leaq	400(%rsp), %rax
	movq	%rax, 144(%rbx)
	leaq	408(%rsp), %rax
	movq	%rax, 152(%rbx)
	leaq	128(%rsp), %r14
	leaq	112(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$320, %rsp                      # imm = 0x140
	.cfi_adjust_cfa_offset -320
	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	_Z21__device_stub__osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_, .Lfunc_end0-_Z21__device_stub__osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z5delayi                       # -- Begin function _Z5delayi
	.type	_Z5delayi,@function
_Z5delayi:                              # @_Z5delayi
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movl	%edi, %ebx
	callq	clock
	movslq	%ebx, %rbx
	addq	%rax, %rbx
.LBB1_1:                                # =>This Inner Loop Header: Depth=1
	callq	clock
	cmpq	%rbx, %rax
	jl	.LBB1_1
# %bb.2:
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z5delayi, .Lfunc_end1-_Z5delayi
	.cfi_endproc
                                        # -- End function
	.globl	_Z15serialize_floatPcf          # -- Begin function _Z15serialize_floatPcf
	.type	_Z15serialize_floatPcf,@function
_Z15serialize_floatPcf:                 # @_Z15serialize_floatPcf
	.cfi_startproc
# %bb.0:
	movd	%xmm0, %eax
	bswapl	%eax
	movl	%eax, (%rdi)
	retq
.Lfunc_end2:
	.size	_Z15serialize_floatPcf, .Lfunc_end2-_Z15serialize_floatPcf
	.cfi_endproc
                                        # -- End function
	.globl	_Z17doublearraytocharPdi        # -- Begin function _Z17doublearraytocharPdi
	.type	_Z17doublearraytocharPdi,@function
_Z17doublearraytocharPdi:               # @_Z17doublearraytocharPdi
	.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	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%esi, %ebp
	movq	%rdi, %rbx
	movslq	%esi, %r15
	leaq	(,%r15,4), %rdi
	callq	malloc
	movq	%rax, %r14
	testl	%r15d, %r15d
	jle	.LBB3_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ebp, %r15d
	xorl	%r12d, %r12d
.LBB3_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$4, %edi
	callq	malloc
	xorps	%xmm0, %xmm0
	cvtsd2ss	(%rbx,%r12,8), %xmm0
	movd	%xmm0, %ecx
	bswapl	%ecx
	movl	%ecx, (%rax)
	movl	%ecx, (%r14,%r12,4)
	incq	%r12
	cmpq	%r12, %r15
	jne	.LBB3_2
.LBB3_3:                                # %._crit_edge
	movq	%r14, %rax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.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_end3:
	.size	_Z17doublearraytocharPdi, .Lfunc_end3-_Z17doublearraytocharPdi
	.cfi_endproc
                                        # -- End function
	.globl	_Z9jsonArrayPdiiPcS0_           # -- Begin function _Z9jsonArrayPdiiPcS0_
	.type	_Z9jsonArrayPdiiPcS0_,@function
_Z9jsonArrayPdiiPcS0_:                  # @_Z9jsonArrayPdiiPcS0_
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
                                        # kill: def $esi killed $esi def $rsi
	movq	%rcx, (%rsp)                    # 8-byte Spill
	movb	$91, (%rcx)
	testl	%esi, %esi
	jle	.LBB4_3
# %bb.1:                                # %.lr.ph
	movq	%r8, %r14
	movq	%rdi, %rbx
	leal	-1(%rsi), %r13d
	addl	%edx, %edx
	movslq	%edx, %r12
	movl	%esi, %ebp
	xorl	%r15d, %r15d
.LBB4_2:                                # =>This Inner Loop Header: Depth=1
	cmpq	%r15, %r13
	movl	$.L.str, %edx
	movl	$.L.str.1, %eax
	cmoveq	%rax, %rdx
	movsd	(%rbx,%r15,8), %xmm0            # xmm0 = mem[0],zero
	movq	%r14, %rdi
	movq	%r12, %rsi
	movb	$1, %al
	callq	snprintf
	movq	(%rsp), %rdi                    # 8-byte Reload
	movq	%r14, %rsi
	callq	strcat
	incq	%r15
	cmpq	%r15, %rbp
	jne	.LBB4_2
.LBB4_3:                                # %._crit_edge
	movq	(%rsp), %rax                    # 8-byte Reload
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	_Z9jsonArrayPdiiPcS0_, .Lfunc_end4-_Z9jsonArrayPdiiPcS0_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI5_0:
	.quad	0x4008000000000000              # double 3
.LCPI5_1:
	.quad	0xbff0000000000000              # double -1
.LCPI5_2:
	.quad	0x3ff0000000000000              # double 1
.LCPI5_3:
	.quad	0x41dfffffffc00000              # double 2147483647
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$248, %rsp
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	xorl	%ebx, %ebx
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	movl	$.Lstr, %edi
	callq	puts@PLT
	movsd	LAYER_HEIGHT(%rip), %xmm1       # xmm1 = mem[0],zero
	movl	ALTITUDE(%rip), %ecx
	cvtsi2sd	%ecx, %xmm0
	ucomisd	%xmm1, %xmm0
	jbe	.LBB5_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%ebx, %ebx
.LBB5_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movsd	%xmm1, -56(%rbp)                # 8-byte Spill
	xorps	%xmm0, %xmm0
	cvtsi2sdl	RES(%rip), %xmm0
	movsd	.LCPI5_0(%rip), %xmm1           # xmm1 = mem[0],zero
	callq	pow
	movsd	%xmm0, -72(%rbp)                # 8-byte Spill
	movl	RES(%rip), %eax
	addl	$-2, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	movsd	.LCPI5_0(%rip), %xmm1           # xmm1 = mem[0],zero
	callq	pow
	movsd	-56(%rbp), %xmm1                # 8-byte Reload
                                        # xmm1 = mem[0],zero
	movsd	-72(%rbp), %xmm2                # 8-byte Reload
                                        # xmm2 = mem[0],zero
	subsd	%xmm0, %xmm2
	cvttsd2si	%xmm2, %eax
	addl	%eax, %ebx
	addsd	LAYER_HEIGHT(%rip), %xmm1
	movl	ALTITUDE(%rip), %ecx
	xorps	%xmm0, %xmm0
	cvtsi2sd	%ecx, %xmm0
	ucomisd	%xmm1, %xmm0
	ja	.LBB5_2
.LBB5_3:                                # %._crit_edge
	movl	RES(%rip), %edx
	xorl	%r14d, %r14d
	movl	$.L.str.3, %edi
	movl	%ebx, %esi
	xorl	%eax, %eax
	callq	printf
	movq	%rbx, -168(%rbp)                # 8-byte Spill
	movslq	%ebx, %rbx
	shlq	$3, %rbx
	leaq	-232(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-224(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-216(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-128(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-120(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-112(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-88(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-80(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-208(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-200(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-192(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-184(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	-176(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movsd	LAYER_HEIGHT(%rip), %xmm1       # xmm1 = mem[0],zero
	cvtsi2sdl	ALTITUDE(%rip), %xmm0
	movsd	%xmm1, -64(%rbp)                # 8-byte Spill
	ucomisd	%xmm1, %xmm0
	jbe	.LBB5_15
# %bb.4:                                # %.preheader160.preheader
	movl	RES(%rip), %ecx
	movsd	.LCPI5_1(%rip), %xmm2           # xmm2 = mem[0],zero
	movsd	.LCPI5_2(%rip), %xmm3           # xmm3 = mem[0],zero
.LBB5_5:                                # %.preheader160
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_6 Depth 2
                                        #       Child Loop BB5_7 Depth 3
                                        #         Child Loop BB5_8 Depth 4
	xorl	%ebx, %ebx
	movapd	%xmm2, %xmm4
.LBB5_6:                                # %.preheader159
                                        #   Parent Loop BB5_5 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB5_7 Depth 3
                                        #         Child Loop BB5_8 Depth 4
	movapd	%xmm4, %xmm0
	movsd	.LCPI5_2(%rip), %xmm1           # xmm1 = mem[0],zero
	cmpeqpd	%xmm1, %xmm0
	movapd	%xmm4, %xmm1
	cmpeqpd	%xmm2, %xmm1
	orpd	%xmm0, %xmm1
	movss	%xmm1, -96(%rbp)                # 4-byte Spill
	movapd	%xmm4, %xmm0
	mulsd	%xmm4, %xmm0
	movapd	%xmm0, -144(%rbp)               # 16-byte Spill
	xorl	%r12d, %r12d
	movapd	%xmm2, %xmm5
	movapd	%xmm4, -256(%rbp)               # 16-byte Spill
.LBB5_7:                                # %.preheader
                                        #   Parent Loop BB5_5 Depth=1
                                        #     Parent Loop BB5_6 Depth=2
                                        # =>    This Loop Header: Depth=3
                                        #         Child Loop BB5_8 Depth 4
	ucomisd	%xmm2, %xmm5
	setnp	%al
	sete	%dl
	andb	%al, %dl
	orb	-96(%rbp), %dl                  # 1-byte Folded Reload
	ucomisd	%xmm3, %xmm5
	setnp	%al
	sete	%r13b
	andb	%al, %r13b
	orb	%dl, %r13b
	movapd	%xmm5, %xmm0
	mulsd	%xmm5, %xmm0
	addsd	-144(%rbp), %xmm0               # 16-byte Folded Reload
	movsd	%xmm0, -104(%rbp)               # 8-byte Spill
	xorl	%r15d, %r15d
	movapd	%xmm2, %xmm6
	movsd	%xmm5, -72(%rbp)                # 8-byte Spill
.LBB5_8:                                #   Parent Loop BB5_5 Depth=1
                                        #     Parent Loop BB5_6 Depth=2
                                        #       Parent Loop BB5_7 Depth=3
                                        # =>      This Inner Loop Header: Depth=4
	ucomisd	%xmm2, %xmm6
	setnp	%dl
	sete	%al
	ucomisd	%xmm3, %xmm6
	jne	.LBB5_9
	jnp	.LBB5_10
.LBB5_9:                                #   in Loop: Header=BB5_8 Depth=4
	andb	%dl, %al
	orb	%r13b, %al
	testb	$1, %al
	je	.LBB5_11
.LBB5_10:                               #   in Loop: Header=BB5_8 Depth=4
	movapd	%xmm6, %xmm0
	mulsd	%xmm6, %xmm0
	addsd	-104(%rbp), %xmm0               # 8-byte Folded Reload
	sqrtsd	%xmm0, %xmm0
	movsd	-64(%rbp), %xmm1                # 8-byte Reload
                                        # xmm1 = mem[0],zero
	divsd	%xmm0, %xmm1
	movapd	%xmm4, %xmm0
	mulsd	%xmm1, %xmm0
	movq	-232(%rbp), %rax
	movl	%ebx, %edx
	imull	%ecx, %edx
	addl	%r12d, %edx
	imull	%ecx, %edx
	addl	%r15d, %edx
	imull	%ecx, %edx
	addl	%r14d, %edx
	movslq	%edx, %rcx
	movsd	%xmm0, (%rax,%rcx,8)
	movapd	%xmm5, %xmm0
	mulsd	%xmm1, %xmm0
	movq	-224(%rbp), %rax
	movsd	%xmm0, (%rax,%rcx,8)
	mulsd	%xmm6, %xmm1
	movq	-216(%rbp), %rax
	movsd	%xmm1, (%rax,%rcx,8)
	movsd	%xmm6, -56(%rbp)                # 8-byte Spill
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	movsd	.LCPI5_3(%rip), %xmm1           # xmm1 = mem[0],zero
	divsd	%xmm1, %xmm0
	movsd	mostAllowedTrajectory(%rip), %xmm1 # xmm1 = mem[0],zero
	movsd	leastAllowedTrajectory(%rip), %xmm2 # xmm2 = mem[0],zero
	subsd	%xmm2, %xmm1
	mulsd	%xmm0, %xmm1
	addsd	%xmm2, %xmm1
	movq	-128(%rbp), %rax
	movl	RES(%rip), %ecx
	movl	%ebx, %edx
	imull	%ecx, %edx
	addl	%r12d, %edx
	imull	%ecx, %edx
	addl	%r15d, %edx
	imull	%ecx, %edx
	addl	%r14d, %edx
	movslq	%edx, %rcx
	movsd	%xmm1, (%rax,%rcx,8)
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	divsd	.LCPI5_3(%rip), %xmm0
	movsd	mostAllowedTrajectory(%rip), %xmm1 # xmm1 = mem[0],zero
	movsd	leastAllowedTrajectory(%rip), %xmm2 # xmm2 = mem[0],zero
	subsd	%xmm2, %xmm1
	mulsd	%xmm0, %xmm1
	addsd	%xmm2, %xmm1
	movq	-120(%rbp), %rax
	movl	RES(%rip), %ecx
	movl	%ebx, %edx
	imull	%ecx, %edx
	addl	%r12d, %edx
	imull	%ecx, %edx
	addl	%r15d, %edx
	imull	%ecx, %edx
	addl	%r14d, %edx
	movslq	%edx, %rcx
	movsd	%xmm1, (%rax,%rcx,8)
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	divsd	.LCPI5_3(%rip), %xmm0
	movsd	mostAllowedTrajectory(%rip), %xmm1 # xmm1 = mem[0],zero
	movsd	leastAllowedTrajectory(%rip), %xmm2 # xmm2 = mem[0],zero
	subsd	%xmm2, %xmm1
	mulsd	%xmm0, %xmm1
	addsd	%xmm2, %xmm1
	movq	-112(%rbp), %rax
	movl	RES(%rip), %ecx
	movl	%ebx, %edx
	imull	%ecx, %edx
	addl	%r12d, %edx
	imull	%ecx, %edx
	addl	%r15d, %edx
	imull	%ecx, %edx
	addl	%r14d, %edx
	movslq	%edx, %rcx
	movsd	%xmm1, (%rax,%rcx,8)
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	divsd	.LCPI5_3(%rip), %xmm0
	movsd	mostAllowedMass(%rip), %xmm1    # xmm1 = mem[0],zero
	movsd	leastAllowedMass(%rip), %xmm2   # xmm2 = mem[0],zero
	subsd	%xmm2, %xmm1
	mulsd	%xmm0, %xmm1
	addsd	%xmm2, %xmm1
	movq	-80(%rbp), %rax
	movl	RES(%rip), %ecx
	movl	%ebx, %edx
	imull	%ecx, %edx
	addl	%r12d, %edx
	imull	%ecx, %edx
	addl	%r15d, %edx
	imull	%ecx, %edx
	addl	%r14d, %edx
	movslq	%edx, %rcx
	movsd	%xmm1, (%rax,%rcx,8)
	callq	rand
	movsd	-56(%rbp), %xmm6                # 8-byte Reload
                                        # xmm6 = mem[0],zero
	movsd	-72(%rbp), %xmm5                # 8-byte Reload
                                        # xmm5 = mem[0],zero
	movapd	-256(%rbp), %xmm4               # 16-byte Reload
	movsd	.LCPI5_2(%rip), %xmm3           # xmm3 = mem[0],zero
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	divsd	.LCPI5_3(%rip), %xmm0
	movsd	mostAllowedTemp(%rip), %xmm1    # xmm1 = mem[0],zero
	movsd	leastAllowedTemp(%rip), %xmm2   # xmm2 = mem[0],zero
	subsd	%xmm2, %xmm1
	mulsd	%xmm0, %xmm1
	addsd	%xmm2, %xmm1
	movsd	.LCPI5_1(%rip), %xmm2           # xmm2 = mem[0],zero
	movq	-88(%rbp), %rax
	movl	RES(%rip), %ecx
	movl	%ebx, %edx
	imull	%ecx, %edx
	addl	%r12d, %edx
	imull	%ecx, %edx
	addl	%r15d, %edx
	imull	%ecx, %edx
	addl	%r14d, %edx
	movslq	%edx, %rdx
	movsd	%xmm1, (%rax,%rdx,8)
.LBB5_11:                               #   in Loop: Header=BB5_8 Depth=4
	leal	-1(%rcx), %esi
	movl	$2, %eax
	xorl	%edx, %edx
	idivl	%esi
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	incl	%r15d
	addsd	%xmm0, %xmm6
	ucomisd	%xmm6, %xmm3
	jae	.LBB5_8
# %bb.12:                               #   in Loop: Header=BB5_7 Depth=3
	incl	%r12d
	addsd	%xmm0, %xmm5
	ucomisd	%xmm5, %xmm3
	jae	.LBB5_7
# %bb.13:                               #   in Loop: Header=BB5_6 Depth=2
	incl	%ebx
	addsd	%xmm0, %xmm4
	ucomisd	%xmm4, %xmm3
	jae	.LBB5_6
# %bb.14:                               #   in Loop: Header=BB5_5 Depth=1
	incl	%r14d
	movsd	-64(%rbp), %xmm1                # 8-byte Reload
                                        # xmm1 = mem[0],zero
	addsd	LAYER_HEIGHT(%rip), %xmm1
	xorps	%xmm0, %xmm0
	cvtsi2sdl	ALTITUDE(%rip), %xmm0
	movsd	%xmm1, -64(%rbp)                # 8-byte Spill
	ucomisd	%xmm1, %xmm0
	ja	.LBB5_5
.LBB5_15:                               # %._crit_edge172
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movq	%rsp, -144(%rbp)                # 8-byte Spill
	movq	-168(%rbp), %r14                # 8-byte Reload
	movl	%r14d, %r12d
	shll	$5, %r12d
	movq	%rsp, %rbx
	subq	%r12, %rbx
	movq	%rbx, %rsp
	movl	$2, %edi
	movl	$2, %esi
	xorl	%edx, %edx
	callq	socket
	testl	%eax, %eax
	js	.LBB5_16
# %bb.18:
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	movq	$0, -276(%rbp)
	xorpd	%xmm0, %xmm0
	movapd	%xmm0, -160(%rbp)
	movl	$2, %edi
	movl	$2, %esi
	xorl	%edx, %edx
	callq	socket
	movl	%eax, %r13d
	testl	%eax, %eax
	js	.LBB5_19
.LBB5_20:
	leaq	-268(%rbp), %rcx
	movl	$1, (%rcx)
	movl	%r13d, %edi
	movl	$1, %esi
	movl	$2, %edx
	movl	$4, %r8d
	callq	setsockopt
	leaq	-284(%rbp), %r15
	movl	$16, %esi
	movq	%r15, %rdi
	callq	bzero
	movq	$604504066, (%r15)              # imm = 0x24080002
	movl	$.Lstr.5, %edi
	callq	puts@PLT
	movl	%r13d, %edi
	movq	%r15, %rsi
	movl	$16, %edx
	callq	bind
	testl	%eax, %eax
	js	.LBB5_21
# %bb.22:
	movl	$.Lstr.6, %edi
	callq	puts@PLT
	leaq	-44(%rbp), %r9
	movl	$16, (%r9)
	leaq	-160(%rbp), %r8
	movl	$1024, %edx                     # imm = 0x400
	movl	%r13d, -72(%rbp)                # 4-byte Spill
	movl	%r13d, %edi
	movq	%rbx, %rsi
	movl	$256, %ecx                      # imm = 0x100
	callq	recvfrom
	movslq	%eax, %r13
	movb	$0, (%rbx,%r13)
	movl	$.L.str.14, %edi
	movq	%rbx, %rsi
	movl	%r13d, %edx
	xorl	%eax, %eax
	callq	printf
	movslq	%r12d, %rbx
	movl	$1, %esi
	movq	%rbx, %rdi
	callq	calloc
	movq	%rax, -104(%rbp)                # 8-byte Spill
	movb	$49, (%rax)
	movl	$1, %esi
	movq	%rbx, %rdi
	callq	calloc
	movq	%rax, -256(%rbp)                # 8-byte Spill
	movb	$50, (%rax)
	leal	(%r14,%rbx), %eax
	incl	%eax
	movslq	%eax, %rdi
	movl	$1, %esi
	callq	calloc
	movq	%rax, -64(%rbp)                 # 8-byte Spill
	movq	%rbx, -96(%rbp)                 # 8-byte Spill
	movl	%ebx, %eax
	orl	$1, %eax
	movslq	%eax, %rdi
	movl	$1, %esi
	callq	calloc
	movq	%rax, -56(%rbp)                 # 8-byte Spill
	testl	%r13d, %r13d
	jle	.LBB5_30
# %bb.23:                               # %.lr.ph178
	movabsq	$4294967296, %rax               # imm = 0x100000000
	movl	%r14d, %r12d
	movq	%r12, %rcx
	btsq	$32, %rcx
	movq	%rcx, -264(%rbp)                # 8-byte Spill
	shlq	$3, %r12
	incq	%rax
	movq	%rax, -144(%rbp)                # 8-byte Spill
	movq	-96(%rbp), %rbx                 # 8-byte Reload
	movq	-64(%rbp), %r15                 # 8-byte Reload
.LBB5_24:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_28 Depth 2
	movq	-264(%rbp), %rdi                # 8-byte Reload
	movl	$1, %esi
	movq	-144(%rbp), %rdx                # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_26
# %bb.25:                               #   in Loop: Header=BB5_24 Depth=1
	movsd	TEMP_RATE(%rip), %xmm0          # xmm0 = mem[0],zero
	movsd	MASS_RATE(%rip), %xmm1          # xmm1 = mem[0],zero
	movsd	TEMP_PUSH(%rip), %xmm2          # xmm2 = mem[0],zero
	movsd	REPULSION_RATE(%rip), %xmm3     # xmm3 = mem[0],zero
	cvtsi2sdl	ALTITUDE(%rip), %xmm4
	movsd	GRAV_CONST(%rip), %xmm5         # xmm5 = mem[0],zero
	movq	-232(%rbp), %rsi
	movq	-224(%rbp), %rdx
	movq	-216(%rbp), %rcx
	movq	-128(%rbp), %r8
	movq	-120(%rbp), %r9
	movl	%r14d, %edi
	pushq	-184(%rbp)
	pushq	-176(%rbp)
	pushq	-88(%rbp)
	pushq	-80(%rbp)
	pushq	-192(%rbp)
	pushq	-200(%rbp)
	pushq	-208(%rbp)
	pushq	-112(%rbp)
	callq	_Z21__device_stub__osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	addq	$64, %rsp
.LBB5_26:                               #   in Loop: Header=BB5_24 Depth=1
	callq	hipDeviceSynchronize
	testl	%r14d, %r14d
	jle	.LBB5_29
# %bb.27:                               # %.lr.ph175
                                        #   in Loop: Header=BB5_24 Depth=1
	movq	-184(%rbp), %rax
	movq	-88(%rbp), %rcx
	movq	-176(%rbp), %rdx
	movq	-80(%rbp), %rsi
	movq	-208(%rbp), %rdi
	movq	-128(%rbp), %r8
	movq	-200(%rbp), %r9
	movq	-120(%rbp), %r10
	movq	-192(%rbp), %r11
	movq	-112(%rbp), %r13
	xorl	%r14d, %r14d
.LBB5_28:                               #   Parent Loop BB5_24 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rax,%r14), %xmm0              # xmm0 = mem[0],zero
	movsd	%xmm0, (%rcx,%r14)
	movsd	(%rdx,%r14), %xmm0              # xmm0 = mem[0],zero
	movsd	%xmm0, (%rsi,%r14)
	movsd	(%rdi,%r14), %xmm0              # xmm0 = mem[0],zero
	movsd	%xmm0, (%r8,%r14)
	movsd	(%r9,%r14), %xmm0               # xmm0 = mem[0],zero
	movsd	%xmm0, (%r10,%r14)
	movsd	(%r11,%r14), %xmm0              # xmm0 = mem[0],zero
	movsd	%xmm0, (%r13,%r14)
	addq	$8, %r14
	cmpq	%r14, %r12
	jne	.LBB5_28
.LBB5_29:                               # %._crit_edge176
                                        #   in Loop: Header=BB5_24 Depth=1
	movl	-44(%rbp), %r9d
	movl	-72(%rbp), %r13d                # 4-byte Reload
	movl	%r13d, %edi
	movq	-104(%rbp), %rsi                # 8-byte Reload
	movq	%rbx, %rdx
	movl	$64, %ecx
	leaq	-160(%rbp), %r8
	callq	sendto
	movq	-88(%rbp), %rdi
	movq	-168(%rbp), %r14                # 8-byte Reload
	movl	%r14d, %esi
	movl	%ebx, %edx
	movq	%r15, %rcx
	movq	-56(%rbp), %r8                  # 8-byte Reload
	callq	_Z9jsonArrayPdiiPcS0_
	movl	-44(%rbp), %r9d
	movl	%r13d, %edi
	movq	%r15, %rsi
	movq	%rbx, %rdx
	movl	$64, %ecx
	leaq	-160(%rbp), %r8
	callq	sendto
	movl	-44(%rbp), %r9d
	movl	%r13d, %edi
	movq	-256(%rbp), %rsi                # 8-byte Reload
	movq	%rbx, %rdx
	movl	$64, %ecx
	leaq	-160(%rbp), %r8
	callq	sendto
	movq	-80(%rbp), %rdi
	movl	%r14d, %esi
	movl	%ebx, %edx
	movq	%r15, %rcx
	movq	-56(%rbp), %r8                  # 8-byte Reload
	callq	_Z9jsonArrayPdiiPcS0_
	movl	-44(%rbp), %r9d
	movl	%r13d, %edi
	movq	%r15, %rsi
	movq	%rbx, %rdx
	movl	$64, %ecx
	leaq	-160(%rbp), %r8
	callq	sendto
	jmp	.LBB5_24
.LBB5_30:                               # %._crit_edge179
	movq	-144(%rbp), %rsp                # 8-byte Reload
	xorl	%eax, %eax
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.LBB5_19:
	.cfi_def_cfa %rbp, 16
	movl	$.L.str.10, %edi
	callq	perror
	jmp	.LBB5_20
.LBB5_16:
	movl	$.L.str.8, %edi
	jmp	.LBB5_17
.LBB5_21:
	movl	$.L.str.12, %edi
.LBB5_17:
	callq	perror
	movl	$1, %edi
	callq	exit
.Lfunc_end5:
	.size	main, .Lfunc_end5-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__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	.LBB7_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
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_,@object # @_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.section	.rodata,"a",@progbits
	.globl	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.p2align	3, 0x0
_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_:
	.quad	_Z21__device_stub__osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.size	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_, 8

	.type	TEMP_RATE,@object               # @TEMP_RATE
	.data
	.globl	TEMP_RATE
	.p2align	3, 0x0
TEMP_RATE:
	.quad	0x3fb999999999999a              # double 0.10000000000000001
	.size	TEMP_RATE, 8

	.type	MASS_RATE,@object               # @MASS_RATE
	.globl	MASS_RATE
	.p2align	3, 0x0
MASS_RATE:
	.quad	0x3fd999999999999a              # double 0.40000000000000002
	.size	MASS_RATE, 8

	.type	TEMP_PUSH,@object               # @TEMP_PUSH
	.globl	TEMP_PUSH
	.p2align	3, 0x0
TEMP_PUSH:
	.quad	0x3f9eb851eb851eb8              # double 0.029999999999999999
	.size	TEMP_PUSH, 8

	.type	REPULSION_RATE,@object          # @REPULSION_RATE
	.globl	REPULSION_RATE
	.p2align	3, 0x0
REPULSION_RATE:
	.quad	0x3ff0000000000000              # double 1
	.size	REPULSION_RATE, 8

	.type	SAMPLE_RADIUS,@object           # @SAMPLE_RADIUS
	.globl	SAMPLE_RADIUS
	.p2align	3, 0x0
SAMPLE_RADIUS:
	.quad	0x3fb999999999999a              # double 0.10000000000000001
	.size	SAMPLE_RADIUS, 8

	.type	GRAV_CONST,@object              # @GRAV_CONST
	.globl	GRAV_CONST
	.p2align	3, 0x0
GRAV_CONST:
	.quad	0x3fb113404ea4a8c1              # double 0.066699999999999995
	.size	GRAV_CONST, 8

	.type	ALTITUDE,@object                # @ALTITUDE
	.globl	ALTITUDE
	.p2align	2, 0x0
ALTITUDE:
	.long	5                               # 0x5
	.size	ALTITUDE, 4

	.type	LAYER_HEIGHT,@object            # @LAYER_HEIGHT
	.globl	LAYER_HEIGHT
	.p2align	3, 0x0
LAYER_HEIGHT:
	.quad	0x3fe0000000000000              # double 0.5
	.size	LAYER_HEIGHT, 8

	.type	RES,@object                     # @RES
	.globl	RES
	.p2align	2, 0x0
RES:
	.long	3                               # 0x3
	.size	RES, 4

	.type	leastAllowedMass,@object        # @leastAllowedMass
	.globl	leastAllowedMass
	.p2align	3, 0x0
leastAllowedMass:
	.quad	0x3f50624dd2f1a9fc              # double 0.001
	.size	leastAllowedMass, 8

	.type	mostAllowedMass,@object         # @mostAllowedMass
	.globl	mostAllowedMass
	.p2align	3, 0x0
mostAllowedMass:
	.quad	0x4014000000000000              # double 5
	.size	mostAllowedMass, 8

	.type	leastAllowedTemp,@object        # @leastAllowedTemp
	.globl	leastAllowedTemp
	.p2align	3, 0x0
leastAllowedTemp:
	.quad	0x3f50624dd2f1a9fc              # double 0.001
	.size	leastAllowedTemp, 8

	.type	mostAllowedTemp,@object         # @mostAllowedTemp
	.globl	mostAllowedTemp
	.p2align	3, 0x0
mostAllowedTemp:
	.quad	0x40e86a0000000000              # double 5.0E+4
	.size	mostAllowedTemp, 8

	.type	leastAllowedTrajectory,@object  # @leastAllowedTrajectory
	.globl	leastAllowedTrajectory
	.p2align	3, 0x0
leastAllowedTrajectory:
	.quad	0x3f1a36e2eb1c432d              # double 1.0E-4
	.size	leastAllowedTrajectory, 8

	.type	mostAllowedTrajectory,@object   # @mostAllowedTrajectory
	.globl	mostAllowedTrajectory
	.p2align	3, 0x0
mostAllowedTrajectory:
	.quad	0x4000000000000000              # double 2
	.size	mostAllowedTrajectory, 8

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

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"%f]"
	.size	.L.str.1, 4

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Allocating for %i samples (pass to your renderer your RES, %i, and your ALTITUDE, %i, like so: python display.py <RES> <ALTITUDE>)...\n"
	.size	.L.str.3, 135

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"socket creation failed"
	.size	.L.str.8, 23

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Error on socket()"
	.size	.L.str.10, 18

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"bind failed"
	.size	.L.str.12, 12

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"%s,%i\n"
	.size	.L.str.14, 7

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_"
	.size	.L__unnamed_1, 43

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

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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"Hello, you are hopefully simulating tectonics. If not, get out of here."
	.size	.Lstr, 72

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"Filling..."
	.size	.Lstr.1, 11

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"Generated, copying to GPU..."
	.size	.Lstr.2, 29

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"Copied. Starting kernel/server loop. Connect to port 2084 in you renderer."
	.size	.Lstr.3, 75

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"wORkiNg..."
	.size	.Lstr.4, 11

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"WorKInG..."
	.size	.Lstr.5, 11

	.type	.Lstr.6,@object                 # @str.6
.Lstr.6:
	.asciz	"Beginning Loop"
	.size	.Lstr.6, 15

	.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 _Z21__device_stub__osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_ ; -- Begin function _Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.globl	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
	.p2align	8
	.type	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_,@function
_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_: ; @_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
; %bb.0:
	s_load_b32 s33, s[0:1], 0x20
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_i32 s33, 1
	s_cbranch_scc1 .LBB0_13
; %bb.1:                                ; %.lr.ph
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0xac
	s_load_b512 s[16:31], s[0:1], 0x30
	s_load_b256 s[4:11], s[0:1], 0x70
	v_mov_b32_e32 v28, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	v_add3_u32 v0, s2, s15, v0
	s_clause 0x1
	s_load_b128 s[12:15], s[0:1], 0x90
	s_load_b256 s[36:43], s[0:1], 0x0
	s_mov_b64 s[0:1], s[10:11]
	s_mov_b64 s[2:3], s[8:9]
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[14:15], 3, v[0:1]
	v_add_co_u32 v0, vcc_lo, s18, v14
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s19, v15, vcc_lo
	v_add_co_u32 v2, vcc_lo, s20, v14
	v_add_co_ci_u32_e32 v3, vcc_lo, s21, v15, vcc_lo
	v_add_co_u32 v4, vcc_lo, s22, v14
	v_add_co_ci_u32_e32 v5, vcc_lo, s23, v15, vcc_lo
	v_add_co_u32 v16, vcc_lo, s8, v14
	v_add_co_ci_u32_e32 v17, vcc_lo, s9, v15, vcc_lo
	v_add_co_u32 v6, vcc_lo, s24, v14
	v_add_co_ci_u32_e32 v7, vcc_lo, s25, v15, vcc_lo
	v_add_co_u32 v8, vcc_lo, s10, v14
	v_add_co_ci_u32_e32 v9, vcc_lo, s11, v15, vcc_lo
	v_add_co_u32 v18, vcc_lo, s30, v14
	v_add_co_ci_u32_e32 v19, vcc_lo, s31, v15, vcc_lo
	v_add_co_u32 v10, vcc_lo, s26, v14
	v_add_co_ci_u32_e32 v11, vcc_lo, s27, v15, vcc_lo
	v_add_co_u32 v20, vcc_lo, s4, v14
	v_add_co_ci_u32_e32 v21, vcc_lo, s5, v15, vcc_lo
	v_add_co_u32 v12, vcc_lo, s28, v14
	v_add_co_ci_u32_e32 v13, vcc_lo, s29, v15, vcc_lo
	v_add_co_u32 v22, vcc_lo, s6, v14
	v_add_co_ci_u32_e32 v23, vcc_lo, s7, v15, vcc_lo
	s_mov_b64 s[4:5], s[22:23]
	s_mov_b64 s[6:7], s[20:21]
	s_mov_b64 s[24:25], s[18:19]
	s_mov_b32 s26, s33
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	global_load_b64 v[24:25], v[2:3], off
	s_clause 0x1
	global_load_b64 v[26:27], v28, s[6:7]
	global_load_b64 v[29:30], v28, s[24:25]
	global_load_b64 v[31:32], v[0:1], off
	global_load_b64 v[33:34], v28, s[4:5]
	global_load_b64 v[35:36], v[4:5], off
	s_mov_b32 s27, exec_lo
	s_waitcnt vmcnt(4)
	v_add_f64 v[26:27], v[26:27], -v[24:25]
	s_waitcnt vmcnt(2)
	v_add_f64 v[24:25], v[29:30], -v[31:32]
	s_waitcnt vmcnt(0)
	v_add_f64 v[29:30], v[33:34], -v[35:36]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[26:27], v[26:27], v[26:27]
	v_fma_f64 v[26:27], v[24:25], v[24:25], v[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[26:27], v[29:30], v[29:30], v[26:27]
	v_cmp_gt_f64_e32 vcc_lo, 0x10000000, v[26:27]
	v_cndmask_b32_e64 v29, 0, 1, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v29, 8, v29
	v_ldexp_f64 v[26:27], v[26:27], v29
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_rsq_f64_e32 v[29:30], v[26:27]
	s_waitcnt_depctr 0xfff
	v_mul_f64 v[31:32], v[26:27], v[29:30]
	v_mul_f64 v[29:30], v[29:30], 0.5
	v_fma_f64 v[33:34], -v[29:30], v[31:32], 0.5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[31:32], v[31:32], v[33:34], v[31:32]
	v_fma_f64 v[29:30], v[29:30], v[33:34], v[29:30]
	v_fma_f64 v[33:34], -v[31:32], v[31:32], v[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[31:32], v[33:34], v[29:30], v[31:32]
	v_fma_f64 v[33:34], -v[31:32], v[31:32], v[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f64 v[29:30], v[33:34], v[29:30], v[31:32]
	v_cndmask_b32_e64 v31, 0, 0xffffff80, vcc_lo
	v_cmp_class_f64_e64 vcc_lo, v[26:27], 0x260
	v_ldexp_f64 v[29:30], v[29:30], v31
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v27, v30, v27 :: v_dual_cndmask_b32 v26, v29, v26
	v_cmpx_neq_f64_e32 0, v[26:27]
	s_cbranch_execz .LBB0_4
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	global_load_b64 v[29:30], v[16:17], off
	s_clause 0x1
	global_load_b64 v[31:32], v28, s[2:3]
	global_load_b64 v[35:36], v28, s[0:1]
	global_load_b64 v[39:40], v[8:9], off
	s_waitcnt vmcnt(3)
	v_mul_f64 v[33:34], v[29:30], s[16:17]
	s_waitcnt vmcnt(2)
	v_add_f64 v[29:30], v[29:30], -v[31:32]
	s_waitcnt vmcnt(0)
	v_add_f64 v[35:36], v[39:40], -v[35:36]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[33:34], v[33:34], v[31:32]
	v_div_scale_f64 v[37:38], null, v[26:27], v[26:27], v[33:34]
	v_div_scale_f64 v[47:48], vcc_lo, v[33:34], v[26:27], v[33:34]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[41:42], v[37:38]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[43:44], -v[37:38], v[41:42], 1.0
	v_fma_f64 v[41:42], v[41:42], v[43:44], v[41:42]
	global_load_b64 v[43:44], v[6:7], off
	v_fma_f64 v[45:46], -v[37:38], v[41:42], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[41:42], v[41:42], v[45:46], v[41:42]
	v_mul_f64 v[45:46], v[47:48], v[41:42]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[37:38], -v[37:38], v[45:46], v[47:48]
	v_div_fmas_f64 v[37:38], v[37:38], v[41:42], v[45:46]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_div_fixup_f64 v[26:27], v[37:38], v[26:27], v[33:34]
	s_waitcnt lgkmcnt(0)
	v_fma_f64 v[31:32], v[35:36], s[40:41], v[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[29:30], v[29:30], s[42:43], v[31:32]
	s_waitcnt vmcnt(0)
	v_fma_f64 v[24:25], v[24:25], v[29:30], v[43:44]
	global_store_b64 v[18:19], v[24:25], off
	global_load_b64 v[24:25], v[8:9], off
	global_load_b64 v[29:30], v28, s[0:1]
	global_load_b64 v[31:32], v[16:17], off
	s_clause 0x1
	global_load_b64 v[33:34], v28, s[2:3]
	global_load_b64 v[35:36], v28, s[6:7]
	global_load_b64 v[37:38], v[2:3], off
	global_load_b64 v[39:40], v[10:11], off
	s_waitcnt vmcnt(5)
	v_add_f64 v[24:25], v[24:25], -v[29:30]
	s_waitcnt vmcnt(3)
	v_add_f64 v[29:30], v[31:32], -v[33:34]
	s_waitcnt vmcnt(1)
	v_add_f64 v[31:32], v[35:36], -v[37:38]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[24:25], v[24:25], s[40:41], v[26:27]
	v_fma_f64 v[24:25], v[29:30], s[42:43], v[24:25]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[24:25], v[24:25], v[31:32], v[39:40]
	global_store_b64 v[20:21], v[24:25], off
	global_load_b64 v[24:25], v[8:9], off
	global_load_b64 v[29:30], v28, s[0:1]
	global_load_b64 v[31:32], v[16:17], off
	s_clause 0x1
	global_load_b64 v[33:34], v28, s[2:3]
	global_load_b64 v[35:36], v28, s[4:5]
	global_load_b64 v[37:38], v[4:5], off
	global_load_b64 v[39:40], v[12:13], off
	s_waitcnt vmcnt(5)
	v_add_f64 v[24:25], v[24:25], -v[29:30]
	s_waitcnt vmcnt(3)
	v_add_f64 v[29:30], v[31:32], -v[33:34]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f64 v[24:25], v[24:25], s[40:41], v[26:27]
	s_waitcnt vmcnt(1)
	v_add_f64 v[26:27], v[35:36], -v[37:38]
	v_fma_f64 v[24:25], v[29:30], s[42:43], v[24:25]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[24:25], v[24:25], v[26:27], v[39:40]
	global_store_b64 v[22:23], v[24:25], off
.LBB0_4:                                ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s27
	s_add_i32 s26, s26, -1
	s_add_u32 s24, s24, 8
	s_addc_u32 s25, s25, 0
	s_add_u32 s6, s6, 8
	s_addc_u32 s7, s7, 0
	s_add_u32 s4, s4, 8
	s_addc_u32 s5, s5, 0
	s_add_u32 s2, s2, 8
	s_addc_u32 s3, s3, 0
	s_add_u32 s0, s0, 8
	s_addc_u32 s1, s1, 0
	s_cmp_eq_u32 s26, 0
	s_cbranch_scc0 .LBB0_2
; %bb.5:                                ; %.preheader
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v14, vcc_lo, s14, v14
	v_add_co_ci_u32_e32 v15, vcc_lo, s15, v15, vcc_lo
	v_mov_b32_e32 v22, 0
.LBB0_6:                                ; =>This Inner Loop Header: Depth=1
	global_load_b64 v[16:17], v[6:7], off
	s_mov_b32 s2, exec_lo
	s_waitcnt vmcnt(0)
	v_cmpx_neq_f64_e32 0, v[16:17]
	s_cbranch_execz .LBB0_12
; %bb.7:                                ;   in Loop: Header=BB0_6 Depth=1
	global_load_b64 v[18:19], v[10:11], off
	s_waitcnt vmcnt(0)
	v_cmp_neq_f64_e32 vcc_lo, 0, v[18:19]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_12
; %bb.8:                                ;   in Loop: Header=BB0_6 Depth=1
	global_load_b64 v[20:21], v[12:13], off
	s_waitcnt vmcnt(0)
	v_cmp_neq_f64_e32 vcc_lo, 0, v[20:21]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_12
; %bb.9:                                ;   in Loop: Header=BB0_6 Depth=1
	global_load_b64 v[23:24], v[0:1], off
	global_load_b64 v[25:26], v22, s[18:19]
	global_load_b64 v[27:28], v[2:3], off
	s_clause 0x1
	global_load_b64 v[29:30], v22, s[20:21]
	global_load_b64 v[31:32], v22, s[22:23]
	global_load_b64 v[33:34], v[4:5], off
	s_waitcnt vmcnt(4)
	v_add_f64 v[23:24], v[25:26], -v[23:24]
	s_waitcnt vmcnt(2)
	v_add_f64 v[25:26], v[29:30], -v[27:28]
	s_waitcnt vmcnt(0)
	v_add_f64 v[27:28], v[31:32], -v[33:34]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f64 v[23:24], v[23:24], v[23:24]
	v_mul_f64 v[25:26], v[25:26], v[25:26]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f64 v[27:28], v[27:28], v[27:28]
	v_div_scale_f64 v[29:30], null, v[16:17], v[16:17], v[23:24]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_div_scale_f64 v[31:32], null, v[18:19], v[18:19], v[25:26]
	v_div_scale_f64 v[33:34], null, v[20:21], v[20:21], v[27:28]
	v_div_scale_f64 v[47:48], vcc_lo, v[23:24], v[16:17], v[23:24]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_rcp_f64_e32 v[35:36], v[29:30]
	v_rcp_f64_e32 v[37:38], v[31:32]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(TRANS32_DEP_3)
	v_rcp_f64_e32 v[39:40], v[33:34]
	v_fma_f64 v[41:42], -v[29:30], v[35:36], 1.0
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[43:44], -v[31:32], v[37:38], 1.0
	v_fma_f64 v[45:46], -v[33:34], v[39:40], 1.0
	v_fma_f64 v[35:36], v[35:36], v[41:42], v[35:36]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[37:38], v[37:38], v[43:44], v[37:38]
	v_fma_f64 v[39:40], v[39:40], v[45:46], v[39:40]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[41:42], -v[29:30], v[35:36], 1.0
	v_fma_f64 v[43:44], -v[31:32], v[37:38], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[45:46], -v[33:34], v[39:40], 1.0
	v_fma_f64 v[35:36], v[35:36], v[41:42], v[35:36]
	v_div_scale_f64 v[41:42], s0, v[25:26], v[18:19], v[25:26]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[37:38], v[37:38], v[43:44], v[37:38]
	v_fma_f64 v[39:40], v[39:40], v[45:46], v[39:40]
	v_div_scale_f64 v[43:44], s1, v[27:28], v[20:21], v[27:28]
	v_mul_f64 v[45:46], v[47:48], v[35:36]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f64 v[49:50], v[41:42], v[37:38]
	v_mul_f64 v[51:52], v[43:44], v[39:40]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[29:30], -v[29:30], v[45:46], v[47:48]
	v_fma_f64 v[31:32], -v[31:32], v[49:50], v[41:42]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[33:34], -v[33:34], v[51:52], v[43:44]
	v_div_fmas_f64 v[29:30], v[29:30], v[35:36], v[45:46]
	s_mov_b32 vcc_lo, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_fmas_f64 v[31:32], v[31:32], v[37:38], v[49:50]
	s_mov_b32 vcc_lo, s1
	v_div_fmas_f64 v[33:34], v[33:34], v[39:40], v[51:52]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_div_fixup_f64 v[16:17], v[29:30], v[16:17], v[23:24]
	v_div_fixup_f64 v[18:19], v[31:32], v[18:19], v[25:26]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_fixup_f64 v[20:21], v[33:34], v[20:21], v[27:28]
	v_add_f64 v[16:17], v[16:17], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[16:17], v[16:17], v[20:21]
	v_cmp_gt_f64_e32 vcc_lo, 0x10000000, v[16:17]
	v_cndmask_b32_e64 v18, 0, 1, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v18, 8, v18
	v_ldexp_f64 v[16:17], v[16:17], v18
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_rsq_f64_e32 v[18:19], v[16:17]
	s_waitcnt_depctr 0xfff
	v_mul_f64 v[20:21], v[16:17], v[18:19]
	v_mul_f64 v[18:19], v[18:19], 0.5
	v_fma_f64 v[23:24], -v[18:19], v[20:21], 0.5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[20:21], v[20:21], v[23:24], v[20:21]
	v_fma_f64 v[18:19], v[18:19], v[23:24], v[18:19]
	v_fma_f64 v[23:24], -v[20:21], v[20:21], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], v[23:24], v[18:19], v[20:21]
	v_fma_f64 v[23:24], -v[20:21], v[20:21], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f64 v[18:19], v[23:24], v[18:19], v[20:21]
	v_cndmask_b32_e64 v20, 0, 0xffffff80, vcc_lo
	v_cmp_class_f64_e64 vcc_lo, v[16:17], 0x260
	v_ldexp_f64 v[18:19], v[18:19], v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v17, v19, v17 :: v_dual_cndmask_b32 v16, v18, v16
	v_cmp_neq_f64_e32 vcc_lo, 0, v[16:17]
	s_delay_alu instid0(VALU_DEP_4)
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_12
; %bb.10:                               ;   in Loop: Header=BB0_6 Depth=1
	v_div_scale_f64 v[18:19], null, v[16:17], v[16:17], s[38:39]
	global_load_b64 v[23:24], v22, s[8:9]
	v_div_scale_f64 v[27:28], vcc_lo, s[38:39], v[16:17], s[38:39]
	v_rcp_f64_e32 v[20:21], v[18:19]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[25:26], -v[18:19], v[20:21], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], v[20:21], v[25:26], v[20:21]
	v_fma_f64 v[25:26], -v[18:19], v[20:21], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], v[20:21], v[25:26], v[20:21]
	v_mul_f64 v[25:26], v[27:28], v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[18:19], -v[18:19], v[25:26], v[27:28]
	v_div_fmas_f64 v[18:19], v[18:19], v[20:21], v[25:26]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_div_fixup_f64 v[16:17], v[18:19], v[16:17], s[38:39]
	s_waitcnt vmcnt(0)
	v_add_f64 v[16:17], v[16:17], v[23:24]
	global_store_b64 v22, v[16:17], s[12:13]
	global_load_b64 v[16:17], v[2:3], off
	s_clause 0x1
	global_load_b64 v[18:19], v22, s[20:21]
	global_load_b64 v[20:21], v22, s[18:19]
	global_load_b64 v[23:24], v[0:1], off
	global_load_b64 v[25:26], v22, s[22:23]
	global_load_b64 v[27:28], v[4:5], off
	s_waitcnt vmcnt(4)
	v_add_f64 v[16:17], v[18:19], -v[16:17]
	s_waitcnt vmcnt(2)
	v_add_f64 v[18:19], v[20:21], -v[23:24]
	s_waitcnt vmcnt(0)
	v_add_f64 v[20:21], v[25:26], -v[27:28]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[16:17], v[16:17], v[16:17]
	v_fma_f64 v[16:17], v[18:19], v[18:19], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[16:17], v[20:21], v[20:21], v[16:17]
	v_cmp_gt_f64_e32 vcc_lo, 0x10000000, v[16:17]
	v_cndmask_b32_e64 v18, 0, 1, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v18, 8, v18
	v_ldexp_f64 v[16:17], v[16:17], v18
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_rsq_f64_e32 v[18:19], v[16:17]
	s_waitcnt_depctr 0xfff
	v_mul_f64 v[20:21], v[16:17], v[18:19]
	v_mul_f64 v[18:19], v[18:19], 0.5
	v_fma_f64 v[23:24], -v[18:19], v[20:21], 0.5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[20:21], v[20:21], v[23:24], v[20:21]
	v_fma_f64 v[18:19], v[18:19], v[23:24], v[18:19]
	v_fma_f64 v[23:24], -v[20:21], v[20:21], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], v[23:24], v[18:19], v[20:21]
	v_fma_f64 v[23:24], -v[20:21], v[20:21], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f64 v[18:19], v[23:24], v[18:19], v[20:21]
	v_cndmask_b32_e64 v20, 0, 0xffffff80, vcc_lo
	v_cmp_class_f64_e64 vcc_lo, v[16:17], 0x260
	v_ldexp_f64 v[18:19], v[18:19], v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v17, v19, v17 :: v_dual_cndmask_b32 v16, v18, v16
	v_cmp_neq_f64_e32 vcc_lo, 0, v[16:17]
	s_delay_alu instid0(VALU_DEP_4)
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_12
; %bb.11:                               ;   in Loop: Header=BB0_6 Depth=1
	global_load_b64 v[18:19], v[8:9], off
	global_load_b64 v[20:21], v22, s[10:11]
	s_waitcnt vmcnt(0)
	v_add_f64 v[18:19], v[18:19], -v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[18:19], v[18:19], s[36:37]
	v_div_scale_f64 v[23:24], null, v[16:17], v[16:17], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[25:26], v[23:24]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[27:28], -v[23:24], v[25:26], 1.0
	v_fma_f64 v[25:26], v[25:26], v[27:28], v[25:26]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[27:28], -v[23:24], v[25:26], 1.0
	v_fma_f64 v[25:26], v[25:26], v[27:28], v[25:26]
	v_div_scale_f64 v[27:28], vcc_lo, v[18:19], v[16:17], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[29:30], v[27:28], v[25:26]
	v_fma_f64 v[23:24], -v[23:24], v[29:30], v[27:28]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[23:24], v[23:24], v[25:26], v[29:30]
	v_div_fixup_f64 v[16:17], v[23:24], v[16:17], v[18:19]
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[18:19], v[20:21], v[16:17]
	global_store_b64 v22, v[18:19], s[14:15]
	global_load_b64 v[18:19], v[8:9], off
	s_waitcnt vmcnt(0)
	v_add_f64 v[16:17], v[18:19], -v[16:17]
	global_store_b64 v[14:15], v[16:17], off
.LBB0_12:                               ;   in Loop: Header=BB0_6 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_add_i32 s33, s33, -1
	s_add_u32 s18, s18, 8
	s_addc_u32 s19, s19, 0
	s_add_u32 s20, s20, 8
	s_addc_u32 s21, s21, 0
	s_add_u32 s22, s22, 8
	s_addc_u32 s23, s23, 0
	s_add_u32 s8, s8, 8
	s_addc_u32 s9, s9, 0
	s_add_u32 s12, s12, 8
	s_addc_u32 s13, s13, 0
	s_add_u32 s10, s10, 8
	s_addc_u32 s11, s11, 0
	s_add_u32 s14, s14, 8
	s_addc_u32 s15, s15, 0
	s_cmp_eq_u32 s33, 0
	s_cbranch_scc0 .LBB0_6
.LBB0_13:                               ; %._crit_edge
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 416
		.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 53
		.amdhsa_next_free_sgpr 44
		.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	_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_, .Lfunc_end0-_Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 2604
; NumSgprs: 46
; NumVgprs: 53
; ScratchSize: 0
; MemoryBound: 1
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 5
; VGPRBlocks: 6
; NumSGPRsForWavesPerEU: 46
; NumVGPRsForWavesPerEU: 53
; 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:
      - .offset:         0
        .size:           8
        .value_kind:     by_value
      - .offset:         8
        .size:           8
        .value_kind:     by_value
      - .offset:         16
        .size:           8
        .value_kind:     by_value
      - .offset:         24
        .size:           8
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           8
        .value_kind:     by_value
      - .offset:         48
        .size:           8
        .value_kind:     by_value
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         80
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         88
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         96
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         104
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         112
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         120
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         128
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         136
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         144
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         152
        .size:           8
        .value_kind:     global_buffer
      - .offset:         160
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         164
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         168
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         172
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         174
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         176
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         178
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         180
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         182
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         200
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         208
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         216
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         224
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 416
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     46
    .sgpr_spill_count: 0
    .symbol:         _Z6osmateddddiddPdS_S_S_S_S_S_S_S_S_S_S_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     53
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
