/*
 * CUDALEAPINT.CU: program to integrate hamiltonian system using leapfrog 
 *                 and CUDA
 * Based off of: 
 * https://courses.physics.ucsd.edu/2020/Winter/physics141/Assignments/leapint.c
 */

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "cudaLeapint.cuh"

int main(int argc, char **argv)
{
    /* Declaring Variables */
    int n, mstep, nout, nstep;
    float massBuffer[MAXPNT];
    float rxBuffer[MAXPNT];
    float ryBuffer[MAXPNT];
    float rzBuffer[MAXPNT];
    float vxBuffer[MAXPNT];
    float vyBuffer[MAXPNT];
    float vzBuffer[MAXPNT];
    float* mass = NULL;
    float* rx = NULL;
    float* ry = NULL;
    float* rz = NULL;
    float* vx = NULL;
    float* vy = NULL;
    float* vz = NULL;
    float* gm = NULL;
    char * nBodyFile;

    float tnow, dt;

    /* GM Constants in AU^3/Day^2 */
    float GMCONST[MAXPNT];
    for (int i = 0; i < MAXPNT; i++) {
        GMCONST[i] = 0.0;
    }

    /* Number of astronomical bodies */
    if (argc <= 1) {
        printf("Usage: %s [N-Body File]\n", argv[0]);
    }
    else {
        nBodyFile = argv[1];
    }

    /* Parsing through the nBodyFile for celestial bodies */
    printf("Reading values in from %s.\n", nBodyFile);
    FILE* fp = fopen((const char *) nBodyFile, "r");

    if (fp == NULL) {
        return 0;
    }

    int lineNumber = 0;
    char buffer[MAXBUFFER];
    char * delimiters = " \t";
    char * token;
    char * s;

    while (fgets(buffer, MAXBUFFER, fp)) {
        s = buffer;
        token = strtok(s, delimiters);
        if (token != NULL) {
            massBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            rxBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            ryBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            rzBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            vxBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            vyBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            vzBuffer[lineNumber] = atof(token);
        }
        lineNumber += 1;
    }
    n = lineNumber;
    fclose(fp);

    /* Allocating Unified Memory - accessible from CPU or GPU */
    cudaSetDevice(0);
    cudaMalloc(&mass, n*sizeof(float));
    cudaMalloc(&rx, n*sizeof(float));
    cudaMalloc(&ry, n*sizeof(float));
    cudaMalloc(&rz, n*sizeof(float));
    cudaMalloc(&vx, n*sizeof(float));
    cudaMalloc(&vy, n*sizeof(float));
    cudaMalloc(&vz, n*sizeof(float));
    cudaMalloc(&gm, n*sizeof(float));

    /* Copying memory to the device */
    cudaMemcpy(mass, massBuffer, n*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(rx, rxBuffer, n*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(ry, ryBuffer, n*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(rz, rzBuffer, n*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(vx, vxBuffer, n*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(vy, vyBuffer, n*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(vz, vzBuffer, n*sizeof(float), cudaMemcpyHostToDevice);


    /* Setting the gravitational constant for the peripheral bodies */
    for (int i = 0; i < n; i++) {
        GMCONST[i] = massBuffer[i];
    }
    cudaMemcpy(gm, GMCONST, n*sizeof(float), cudaMemcpyHostToDevice);


    /* Setting initial time */
    tnow = 0.0;

    /* next, set integration parameters */

    mstep = 700;                      /* number of steps to take  */
    nout = 1;                         /* steps between outputs    */
    dt = 0.02;                        /* timestep for integration */

    /* Checking to see if n is a multiple of nThreads*deviceCount 
       If not, then round down
       (As seen in David's code) */
    int nParticles = nThreads * int (float(n) / (nThreads));
    if (nParticles != n) {
        n = nParticles;
    }

    int numBlocks = n / nThreads;
    if (numBlocks == 0) {
        numBlocks = 1;
    }

    printf("numBlocks: %d\nnThreads: %d\n\nParticles: %d\n", numBlocks, nThreads, n);

    /* now, loop performing integration */

    /* Gravity Acceleration */
    {
        FILE* outFile;
        const char * filename;
        char * outBuffer = strtok(nBodyFile, ".");
        filename = strcat(outBuffer, "Sim.txt");
        outFile = fopen(filename, "w");

        /* Progress Bar Initialization */
        int maxBar = 30;
        char outputBar[maxBar];
        outputBar[0] = '|';
        outputBar[maxBar - 1]= '|';
        for (int i = 1; i < maxBar - 1; i++) {
            outputBar[i] = ' ';
        }

        printstate(rxBuffer, ryBuffer, rzBuffer, 
                   vxBuffer, vyBuffer, vzBuffer, n, tnow, outFile, filename);

        for (nstep = 0; nstep < mstep; nstep++) {	

            /* Progress Bar Handling/Management */
            int v = round((((double) nstep) / ((double) mstep)) * 100.0);
            int barIndex = (int) round(((double) v / 100.0) * maxBar) + 1;
            if (barIndex < maxBar - 1) {
                outputBar[barIndex] = '#';
            }
            printf("\r%s  %d%%", outputBar, v);
            fflush(stdout);

            /* then call output routine */

            leapstep <<<numBlocks, nThreads>>>(rx, ry, rz,
                                               vx, vy, vz,
                                               n, dt, gm, 0);
            cudaDeviceSynchronize();

            /* Copying memory from device to computer */
            cudaMemcpy(massBuffer, mass, n*sizeof(float), cudaMemcpyDeviceToHost);
            cudaMemcpy(rxBuffer, rx, n*sizeof(float), cudaMemcpyDeviceToHost);
            cudaMemcpy(ryBuffer, ry, n*sizeof(float), cudaMemcpyDeviceToHost);
            cudaMemcpy(rzBuffer, rz, n*sizeof(float), cudaMemcpyDeviceToHost);
            cudaMemcpy(vxBuffer, vx, n*sizeof(float), cudaMemcpyDeviceToHost);
            cudaMemcpy(vyBuffer, vy, n*sizeof(float), cudaMemcpyDeviceToHost);
            cudaMemcpy(vzBuffer, vz, n*sizeof(float), cudaMemcpyDeviceToHost);

            /* take integration step    */
            tnow = tnow + dt;			
            /* and update value of time */

            /* Printing out current positions and velocities */
            if (nstep % nout == 0) {
                printstate(rxBuffer, ryBuffer, rzBuffer, 
                           vxBuffer, vyBuffer, vzBuffer, n, tnow, outFile, filename);
            }
        }
        cudaDeviceSynchronize();

        /* Outputting to file */
        printstate(rxBuffer, ryBuffer, rzBuffer, 
                   vxBuffer, vyBuffer, vzBuffer, n, tnow, outFile, filename);
        /* Cleaning up progress bar */
        printf("\r%s  %d%%\n", outputBar, 100);
        fflush(stdout);
    }

    /* Freeing memory */
    cudaFree(mass);
    cudaFree(rx);
    cudaFree(ry);
    cudaFree(rz);
    cudaFree(vx);
    cudaFree(vy);
    cudaFree(vz);
    cudaFree(gm);
    cudaDeviceReset();
}

/*
 * LEAPSTEP: take one step using the leap-from integrator, formulated
 * as a mapping from t to t + dt.  WARNING: this integrator is not
 * accurate unless the timestep dt is fixed from one call to another.
 */

__global__ void leapstep(float rx[], float ry[], float rz[], 
                         float vx[], float vy[], float vz[], 
                         int n, float dt, float gmConst[], int deviceOffset)
{
    int index = deviceOffset + blockIdx.x * blockDim.x + threadIdx.x;
    float3 ac3;

    /* call acceleration code */
    ac3 = accel(rx, ry, rz, n, gmConst, deviceOffset, index);
    __syncthreads();

    /* Applying acceleration to velocity */
    vx[index] = vx[index] + 0.5 * dt * ac3.x;
    vy[index] = vy[index] + 0.5 * dt * ac3.y;
    vz[index] = vz[index] + 0.5 * dt * ac3.z;

    /* Applying velocity to position */
    rx[index] = rx[index] + dt * vx[index];
    ry[index] = ry[index] + dt * vy[index];
    rz[index] = rz[index] + dt * vz[index];

    /* call acceleration code */
    ac3 = accel(rx, ry, rz, n, gmConst, deviceOffset, index);
    __syncthreads();

    vx[index] = vx[index] + 0.5 * dt * ac3.x;
    vy[index] = vy[index] + 0.5 * dt * ac3.y;
    vz[index] = vz[index] + 0.5 * dt * ac3.z;
}

/*
 * ACCEL: compute accelerations for harmonic oscillator(s).
 */

__device__
float3 accel(float* rx, float* ry, float* rz, 
             int n, float gmConst[], int deviceOffset, int index)
{
    float3 ac3 = {0.0f, 0.0f, 0.0f};
    if (index != 0) {
        for (int j = 0; j < n; j++) {
            if (j != index) {
                float distVal = (rx[index]-rx[j])*(rx[index]-rx[j])
                                +(ry[index]-ry[j])*(ry[index]-ry[j])
                                +(rz[index]-rz[j])*(rz[index]-rz[j])
                                + SOFT_FACTOR;
                distVal = distVal * distVal * distVal;
                distVal = 1.0f / sqrtf(distVal);
        
                /* Summing up acceleration */
                ac3.x += -(rx[index]-rx[j])*gmConst[j]*distVal;
                ac3.y += -(ry[index]-ry[j])*gmConst[j]*distVal;
                ac3.z += -(rz[index]-rz[j])*gmConst[j]*distVal;
            }
            __syncthreads();
        }
    }
    return ac3;
}

/*
 * PRINTSTATE: output system state variables.
 */

void printstate(float rx[], float ry[], float rz[],
                float vx[], float vy[], float vz[], 
                int n, float tnow, FILE* outFile, const char * filename)
{
    int i;
    outFile = fopen(filename, "a+");
    for (i = 0; i < n; i++)	{		
        /* Printing out time, particle, position, and velocity */
        fprintf(outFile, 
                "%8.4f\t%4d\t%20.6f\t%20.6f\t%20.6f\t%20.6f\t%20.6f\t%20.6f\n", 
                tnow, i, rx[i], ry[i], rz[i], vx[i], vy[i], vz[i]);
    }
    fclose(outFile);
}


	.file	"tmpxft_00370d72_00000000-6_cudaLeapint.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2032:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2032:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0, @function
_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0:
.LFB2174:
	.cfi_startproc
	jmp	cudaMalloc@PLT
	.cfi_endproc
.LFE2174:
	.size	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0, .-_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	.globl	_Z5accelPfS_S_iS_ii
	.type	_Z5accelPfS_S_iS_ii, @function
_Z5accelPfS_S_iS_ii:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z5accelPfS_S_iS_ii, .-_Z5accelPfS_S_iS_ii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"a+"
.LC1:
	.string	"%8.4f\t%4d\t%20.6f\t%20.6f\t%20.6f\t%20.6f\t%20.6f\t%20.6f\n"
	.text
	.globl	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	.type	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc, @function
_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%r9, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%r8, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rcx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	leaq	.LC0(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdi, %rbx
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movl	96(%rsp), %eax
	movq	112(%rsp), %rdi
	movss	%xmm0, 12(%rsp)
	movl	%eax, 8(%rsp)
	call	fopen@PLT
	xorl	%r8d, %r8d
	movq	%rax, %rdi
.L6:
	movl	%r8d, %ecx
	cmpl	%r8d, 8(%rsp)
	jle	.L10
	leaq	.LC1(%rip), %rdx
	movl	$2, %esi
	movb	$7, %al
	movq	%r8, 24(%rsp)
	movq	%rdi, 16(%rsp)
	cvtss2sd	(%r15,%r8,4), %xmm6
	cvtss2sd	(%r14,%r8,4), %xmm5
	cvtss2sd	0(%r13,%r8,4), %xmm4
	cvtss2sd	(%r12,%r8,4), %xmm3
	cvtss2sd	0(%rbp,%r8,4), %xmm2
	cvtss2sd	(%rbx,%r8,4), %xmm1
	cvtss2sd	12(%rsp), %xmm0
	call	__fprintf_chk@PLT
	movq	24(%rsp), %r8
	movq	16(%rsp), %rdi
	incq	%r8
	jmp	.L6
.L10:
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	fclose@PLT
	.cfi_endproc
.LFE2029:
	.size	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc, .-_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	.globl	_Z42__device_stub__Z8leapstepPfS_S_S_S_S_ifS_iPfS_S_S_S_S_ifS_i
	.type	_Z42__device_stub__Z8leapstepPfS_S_S_S_S_ifS_iPfS_S_S_S_S_ifS_i, @function
_Z42__device_stub__Z8leapstepPfS_S_S_S_S_ifS_iPfS_S_S_S_S_ifS_i:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$232, %rsp
	.cfi_def_cfa_offset 240
	movq	248(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, (%rsp)
	movss	%xmm0, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 216(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	240(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 192(%rsp)
	movq	%rsp, %rax
	movq	%rax, 200(%rsp)
	leaq	256(%rsp), %rax
	movq	%rax, 208(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L11
	pushq	80(%rsp)
	.cfi_def_cfa_offset 248
	leaq	_Z8leapstepPfS_S_S_S_S_ifS_i(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 256
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 248
	popq	%rdx
	.cfi_def_cfa_offset 240
.L11:
	movq	216(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$232, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z42__device_stub__Z8leapstepPfS_S_S_S_S_ifS_iPfS_S_S_S_S_ifS_i, .-_Z42__device_stub__Z8leapstepPfS_S_S_S_S_ifS_iPfS_S_S_S_S_ifS_i
	.globl	_Z8leapstepPfS_S_S_S_S_ifS_i
	.type	_Z8leapstepPfS_S_S_S_S_ifS_i, @function
_Z8leapstepPfS_S_S_S_S_ifS_i:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z8leapstepPfS_S_S_S_S_ifS_iPfS_S_S_S_S_ifS_i
	.cfi_endproc
.LFE2055:
	.size	_Z8leapstepPfS_S_S_S_S_ifS_i, .-_Z8leapstepPfS_S_S_S_S_ifS_i
	.section	.rodata.str1.1
.LC3:
	.string	"Usage: %s [N-Body File]\n"
.LC4:
	.string	"Reading values in from %s.\n"
.LC5:
	.string	"r"
.LC6:
	.string	" \t"
.LC8:
	.string	"numBlocks: %d\nnThreads: %d\n\nParticles: %d\n"
.LC9:
	.string	"."
.LC10:
	.string	"Sim.txt"
.LC11:
	.string	"w"
.LC16:
	.string	"\r%s  %d%%"
.LC18:
	.string	"\r%s  %d%%\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	leaq	-483328(%rsp), %r11
.LPSRL0:
	subq	$4096, %rsp
	orq	$0, (%rsp)
	cmpq	%r11, %rsp
	jne	.LPSRL0
	subq	$952, %rsp
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	.cfi_offset 12, -48
	.cfi_offset 3, -56
	xorl	%ebx, %ebx
	movl	$15000, %ecx
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	movl	%edi, %edx
	movq	%rbx, -484240(%rbp)
	movq	%rbx, -484232(%rbp)
	decl	%edx
	movq	%rbx, -484224(%rbp)
	movq	%rbx, -484216(%rbp)
	movq	%rbx, -484208(%rbp)
	movq	%rbx, -484200(%rbp)
	movq	%rbx, -484192(%rbp)
	movq	%rbx, -484184(%rbp)
	leaq	-64152(%rbp), %rbx
	movq	%rbx, %rdi
	rep stosl
	jg	.L17
	movq	(%rsi), %rdx
	movl	$2, %edi
	leaq	.LC3(%rip), %rsi
	call	__printf_chk@PLT
	jmp	.L18
.L17:
	movq	8(%rsi), %rax
	movq	%rax, -484256(%rbp)
.L18:
	movq	-484256(%rbp), %rdx
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%r13d, %r13d
	leaq	.LC6(%rip), %r15
	call	__printf_chk@PLT
	movq	-484256(%rbp), %rdi
	leaq	.LC5(%rip), %rsi
	call	fopen@PLT
	movq	%rax, %r14
	testq	%rax, %rax
	je	.L19
.L20:
	leaq	-4152(%rbp), %r12
	movq	%r14, %rcx
	movl	$4096, %edx
	movl	$4096, %esi
	movq	%r12, %rdi
	call	__fgets_chk@PLT
	testq	%rax, %rax
	je	.L66
	movq	%r12, %rdi
	movq	%r15, %rsi
	call	strtok@PLT
	movq	%rax, %rdi
	testq	%rax, %rax
	je	.L21
	call	atof@PLT
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -484152(%rbp,%r13,4)
.L21:
	xorl	%edi, %edi
	movq	%r15, %rsi
	call	strtok@PLT
	movq	%rax, %rdi
	testq	%rax, %rax
	je	.L22
	call	atof@PLT
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -424152(%rbp,%r13,4)
.L22:
	xorl	%edi, %edi
	movq	%r15, %rsi
	call	strtok@PLT
	movq	%rax, %rdi
	testq	%rax, %rax
	je	.L23
	call	atof@PLT
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -364152(%rbp,%r13,4)
.L23:
	xorl	%edi, %edi
	movq	%r15, %rsi
	call	strtok@PLT
	movq	%rax, %rdi
	testq	%rax, %rax
	je	.L24
	call	atof@PLT
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -304152(%rbp,%r13,4)
.L24:
	xorl	%edi, %edi
	movq	%r15, %rsi
	call	strtok@PLT
	movq	%rax, %rdi
	testq	%rax, %rax
	je	.L25
	call	atof@PLT
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -244152(%rbp,%r13,4)
.L25:
	xorl	%edi, %edi
	movq	%r15, %rsi
	call	strtok@PLT
	movq	%rax, %rdi
	testq	%rax, %rax
	je	.L26
	call	atof@PLT
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -184152(%rbp,%r13,4)
.L26:
	xorl	%edi, %edi
	movq	%r15, %rsi
	call	strtok@PLT
	movq	%rax, %rdi
	testq	%rax, %rax
	je	.L27
	call	atof@PLT
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -124152(%rbp,%r13,4)
.L27:
	incq	%r13
	jmp	.L20
.L66:
	movq	%r14, %rdi
	movl	%r13d, %r15d
	movslq	%r13d, %r13
	call	fclose@PLT
	xorl	%edi, %edi
	leaq	-484152(%rbp), %r12
	leaq	-304152(%rbp), %r14
	call	cudaSetDevice@PLT
	leaq	0(,%r13,4), %rax
	leaq	-484240(%rbp), %rdi
	movq	%rax, %rsi
	movq	%rax, -484248(%rbp)
	leaq	-364152(%rbp), %r13
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	-484248(%rbp), %rsi
	leaq	-484232(%rbp), %rdi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	-484248(%rbp), %rsi
	leaq	-484224(%rbp), %rdi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	-484248(%rbp), %rsi
	leaq	-484216(%rbp), %rdi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	-484248(%rbp), %rsi
	leaq	-484208(%rbp), %rdi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	-484248(%rbp), %rsi
	leaq	-484200(%rbp), %rdi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	-484248(%rbp), %rsi
	leaq	-484192(%rbp), %rdi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	-484248(%rbp), %rsi
	leaq	-484184(%rbp), %rdi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	-484248(%rbp), %rdx
	movl	$1, %ecx
	movq	%r12, %rsi
	movq	-484240(%rbp), %rdi
	call	cudaMemcpy@PLT
	leaq	-424152(%rbp), %rax
	movl	$1, %ecx
	movq	-484248(%rbp), %rdx
	movq	-484232(%rbp), %rdi
	movq	%rax, %rsi
	movq	%rax, -484280(%rbp)
	call	cudaMemcpy@PLT
	movq	-484248(%rbp), %rdx
	movl	$1, %ecx
	movq	%r13, %rsi
	movq	-484224(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-484248(%rbp), %rdx
	movl	$1, %ecx
	movq	%r14, %rsi
	movq	-484216(%rbp), %rdi
	call	cudaMemcpy@PLT
	movq	-484248(%rbp), %rdx
	movl	$1, %ecx
	movq	-484208(%rbp), %rdi
	leaq	-244152(%rbp), %rsi
	call	cudaMemcpy@PLT
	movq	-484248(%rbp), %rdx
	movl	$1, %ecx
	movq	-484200(%rbp), %rdi
	leaq	-184152(%rbp), %rsi
	call	cudaMemcpy@PLT
	movq	-484248(%rbp), %rdx
	movl	$1, %ecx
	movq	-484192(%rbp), %rdi
	leaq	-124152(%rbp), %rsi
	call	cudaMemcpy@PLT
	xorl	%eax, %eax
.L29:
	cmpl	%eax, %r15d
	jle	.L67
	movss	(%r12,%rax,4), %xmm0
	movss	%xmm0, (%rbx,%rax,4)
	incq	%rax
	jmp	.L29
.L67:
	movq	-484248(%rbp), %rdx
	movq	-484184(%rbp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	cvtsi2ssl	%r15d, %xmm0
	mulss	.LC7(%rip), %xmm0
	cvttss2sil	%xmm0, %r15d
	movl	%r15d, %ebx
	sall	$10, %ebx
	testl	%r15d, %r15d
	jne	.L31
	movl	$1, %r15d
.L31:
	movl	%ebx, %r8d
	movl	$1024, %ecx
	movl	%r15d, %edx
	movl	$2, %edi
	leaq	.LC8(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	-484256(%rbp), %rdi
	leaq	.LC9(%rip), %rsi
	movq	%rsp, -484320(%rbp)
	call	strtok@PLT
	leaq	.LC10(%rip), %rsi
	movq	%rax, %rdi
	call	strcat@PLT
	leaq	.LC11(%rip), %rsi
	movq	%rax, %rdi
	movq	%rax, %r12
	call	fopen@PLT
	movq	%rax, -484264(%rbp)
.L32:
	movq	-484320(%rbp), %rax
	cmpq	%rax, %rsp
	je	.L33
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L32
.L33:
	subq	$32, %rsp
	orq	$0, 24(%rsp)
	movaps	.LC12(%rip), %xmm0
	movups	%xmm0, 1(%rsp)
	movq	%r13, %rsi
	movq	%r14, %rdx
	leaq	-244152(%rbp), %rcx
	movups	%xmm0, 13(%rsp)
	leaq	-124152(%rbp), %r9
	xorps	%xmm0, %xmm0
	xorl	%r13d, %r13d
	movq	%rsp, -484272(%rbp)
	leaq	-184152(%rbp), %r8
	movb	$124, (%rsp)
	movb	$124, 29(%rsp)
	pushq	%r11
	movq	-484280(%rbp), %rdi
	pushq	%r12
	pushq	-484264(%rbp)
	pushq	%rbx
	call	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	movslq	%ebx, %rax
	addq	$32, %rsp
	movl	$0x00000000, -484256(%rbp)
	salq	$2, %rax
	movq	%rax, -484248(%rbp)
.L37:
	cvtsi2sdl	%r13d, %xmm0
	divsd	.LC13(%rip), %xmm0
	mulsd	.LC14(%rip), %xmm0
	call	round@PLT
	cvttsd2sil	%xmm0, %r14d
	cvtsi2sdl	%r14d, %xmm0
	divsd	.LC14(%rip), %xmm0
	mulsd	.LC15(%rip), %xmm0
	call	round@PLT
	cvttsd2sil	%xmm0, %eax
	incl	%eax
	cmpl	$28, %eax
	jg	.L35
	movq	-484272(%rbp), %rcx
	cltq
	movb	$35, (%rcx,%rax)
.L35:
	movq	-484272(%rbp), %rdx
	movl	%r14d, %ecx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC16(%rip), %rsi
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	%r15d, -484176(%rbp)
	salq	$10, %rdx
	movl	$1, %ecx
	movabsq	$4294967297, %rax
	movq	%rax, -484172(%rbp)
	movl	-484168(%rbp), %esi
	movq	-484176(%rbp), %rdi
	movq	%rdx, -484164(%rbp)
	movl	$1, -484156(%rbp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L36
	movq	-484192(%rbp), %r9
	movq	-484200(%rbp), %r8
	movq	-484208(%rbp), %rcx
	movq	-484216(%rbp), %rdx
	movq	-484224(%rbp), %rsi
	movq	-484232(%rbp), %rdi
	pushq	%r10
	pushq	$0
	movss	.LC17(%rip), %xmm0
	pushq	-484184(%rbp)
	pushq	%rbx
	call	_Z42__device_stub__Z8leapstepPfS_S_S_S_S_ifS_iPfS_S_S_S_S_ifS_i
	addq	$32, %rsp
.L36:
	call	cudaDeviceSynchronize@PLT
	movq	-484248(%rbp), %rdx
	movl	$2, %ecx
	incl	%r13d
	movq	-484240(%rbp), %rsi
	leaq	-484152(%rbp), %rdi
	leaq	-124152(%rbp), %r14
	call	cudaMemcpy@PLT
	leaq	-424152(%rbp), %rax
	movl	$2, %ecx
	movq	-484248(%rbp), %rdx
	movq	-484232(%rbp), %rsi
	movq	%rax, %rdi
	movq	%rax, -484280(%rbp)
	call	cudaMemcpy@PLT
	leaq	-364152(%rbp), %rax
	movl	$2, %ecx
	movq	-484248(%rbp), %rdx
	movq	-484224(%rbp), %rsi
	movq	%rax, %rdi
	movq	%rax, -484288(%rbp)
	call	cudaMemcpy@PLT
	leaq	-304152(%rbp), %rax
	movl	$2, %ecx
	movq	-484248(%rbp), %rdx
	movq	-484216(%rbp), %rsi
	movq	%rax, %rdi
	movq	%rax, -484296(%rbp)
	call	cudaMemcpy@PLT
	leaq	-244152(%rbp), %rax
	movl	$2, %ecx
	movq	-484248(%rbp), %rdx
	movq	-484208(%rbp), %rsi
	movq	%rax, %rdi
	movq	%rax, -484304(%rbp)
	call	cudaMemcpy@PLT
	leaq	-184152(%rbp), %rax
	movl	$2, %ecx
	movq	-484248(%rbp), %rdx
	movq	-484200(%rbp), %rsi
	movq	%rax, %rdi
	movq	%rax, -484312(%rbp)
	call	cudaMemcpy@PLT
	movq	-484248(%rbp), %rdx
	movl	$2, %ecx
	movq	%r14, %rdi
	movq	-484192(%rbp), %rsi
	call	cudaMemcpy@PLT
	movss	.LC17(%rip), %xmm1
	movq	%r14, %r9
	addss	-484256(%rbp), %xmm1
	movss	%xmm1, -484256(%rbp)
	pushq	%rdx
	movss	-484256(%rbp), %xmm0
	movq	-484312(%rbp), %r8
	movq	-484304(%rbp), %rcx
	pushq	%r12
	pushq	-484264(%rbp)
	movq	-484296(%rbp), %rdx
	movq	-484288(%rbp), %rsi
	pushq	%rbx
	movq	-484280(%rbp), %rdi
	call	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	addq	$32, %rsp
	cmpl	$700, %r13d
	jne	.L37
	call	cudaDeviceSynchronize@PLT
	movq	%r14, %r9
	pushq	%rax
	movss	-484256(%rbp), %xmm0
	movq	-484312(%rbp), %r8
	movq	-484304(%rbp), %rcx
	pushq	%r12
	pushq	-484264(%rbp)
	movq	-484296(%rbp), %rdx
	movq	-484288(%rbp), %rsi
	pushq	%rbx
	movq	-484280(%rbp), %rdi
	call	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	movl	$100, %ecx
	addq	$32, %rsp
	xorl	%eax, %eax
	movq	-484272(%rbp), %rdx
	leaq	.LC18(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	stdout(%rip), %rdi
	call	fflush@PLT
	movq	-484320(%rbp), %rsp
	movq	-484240(%rbp), %rdi
	call	cudaFree@PLT
	movq	-484232(%rbp), %rdi
	call	cudaFree@PLT
	movq	-484224(%rbp), %rdi
	call	cudaFree@PLT
	movq	-484216(%rbp), %rdi
	call	cudaFree@PLT
	movq	-484208(%rbp), %rdi
	call	cudaFree@PLT
	movq	-484200(%rbp), %rdi
	call	cudaFree@PLT
	movq	-484192(%rbp), %rdi
	call	cudaFree@PLT
	movq	-484184(%rbp), %rdi
	call	cudaFree@PLT
	call	cudaDeviceReset@PLT
.L19:
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L38
	call	__stack_chk_fail@PLT
.L38:
	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
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC19:
	.string	"_Z8leapstepPfS_S_S_S_S_ifS_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2057:
	.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	_Z8leapstepPfS_S_S_S_S_ifS_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2057:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC7:
	.long	981467136
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC12:
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.byte	32
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC13:
	.long	0
	.long	1082515456
	.align 8
.LC14:
	.long	0
	.long	1079574528
	.align 8
.LC15:
	.long	0
	.long	1077805056
	.section	.rodata.cst4
	.align 4
.LC17:
	.long	1017370378
	.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:

/*
 * CUDALEAPINT.CU: program to integrate hamiltonian system using leapfrog 
 *                 and CUDA
 * Based off of: 
 * https://courses.physics.ucsd.edu/2020/Winter/physics141/Assignments/leapint.c
 */

#include "hip/hip_runtime.h"

#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "cudaLeapint.cuh"

int main(int argc, char **argv)
{
    /* Declaring Variables */
    int n, mstep, nout, nstep;
    float massBuffer[MAXPNT];
    float rxBuffer[MAXPNT];
    float ryBuffer[MAXPNT];
    float rzBuffer[MAXPNT];
    float vxBuffer[MAXPNT];
    float vyBuffer[MAXPNT];
    float vzBuffer[MAXPNT];
    float* mass = NULL;
    float* rx = NULL;
    float* ry = NULL;
    float* rz = NULL;
    float* vx = NULL;
    float* vy = NULL;
    float* vz = NULL;
    float* gm = NULL;
    char * nBodyFile;

    float tnow, dt;

    /* GM Constants in AU^3/Day^2 */
    float GMCONST[MAXPNT];
    for (int i = 0; i < MAXPNT; i++) {
        GMCONST[i] = 0.0;
    }

    /* Number of astronomical bodies */
    if (argc <= 1) {
        printf("Usage: %s [N-Body File]\n", argv[0]);
    }
    else {
        nBodyFile = argv[1];
    }

    /* Parsing through the nBodyFile for celestial bodies */
    printf("Reading values in from %s.\n", nBodyFile);
    FILE* fp = fopen((const char *) nBodyFile, "r");

    if (fp == NULL) {
        return 0;
    }

    int lineNumber = 0;
    char buffer[MAXBUFFER];
    char * delimiters = " \t";
    char * token;
    char * s;

    while (fgets(buffer, MAXBUFFER, fp)) {
        s = buffer;
        token = strtok(s, delimiters);
        if (token != NULL) {
            massBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            rxBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            ryBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            rzBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            vxBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            vyBuffer[lineNumber] = atof(token);
        }
        token = strtok(NULL, delimiters);
        if (token != NULL) {
            vzBuffer[lineNumber] = atof(token);
        }
        lineNumber += 1;
    }
    n = lineNumber;
    fclose(fp);

    /* Allocating Unified Memory - accessible from CPU or GPU */
    hipSetDevice(0);
    hipMalloc(&mass, n*sizeof(float));
    hipMalloc(&rx, n*sizeof(float));
    hipMalloc(&ry, n*sizeof(float));
    hipMalloc(&rz, n*sizeof(float));
    hipMalloc(&vx, n*sizeof(float));
    hipMalloc(&vy, n*sizeof(float));
    hipMalloc(&vz, n*sizeof(float));
    hipMalloc(&gm, n*sizeof(float));

    /* Copying memory to the device */
    hipMemcpy(mass, massBuffer, n*sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(rx, rxBuffer, n*sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(ry, ryBuffer, n*sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(rz, rzBuffer, n*sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(vx, vxBuffer, n*sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(vy, vyBuffer, n*sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(vz, vzBuffer, n*sizeof(float), hipMemcpyHostToDevice);


    /* Setting the gravitational constant for the peripheral bodies */
    for (int i = 0; i < n; i++) {
        GMCONST[i] = massBuffer[i];
    }
    hipMemcpy(gm, GMCONST, n*sizeof(float), hipMemcpyHostToDevice);


    /* Setting initial time */
    tnow = 0.0;

    /* next, set integration parameters */

    mstep = 700;                      /* number of steps to take  */
    nout = 1;                         /* steps between outputs    */
    dt = 0.02;                        /* timestep for integration */

    /* Checking to see if n is a multiple of nThreads*deviceCount 
       If not, then round down
       (As seen in David's code) */
    int nParticles = nThreads * int (float(n) / (nThreads));
    if (nParticles != n) {
        n = nParticles;
    }

    int numBlocks = n / nThreads;
    if (numBlocks == 0) {
        numBlocks = 1;
    }

    printf("numBlocks: %d\nnThreads: %d\n\nParticles: %d\n", numBlocks, nThreads, n);

    /* now, loop performing integration */

    /* Gravity Acceleration */
    {
        FILE* outFile;
        const char * filename;
        char * outBuffer = strtok(nBodyFile, ".");
        filename = strcat(outBuffer, "Sim.txt");
        outFile = fopen(filename, "w");

        /* Progress Bar Initialization */
        int maxBar = 30;
        char outputBar[maxBar];
        outputBar[0] = '|';
        outputBar[maxBar - 1]= '|';
        for (int i = 1; i < maxBar - 1; i++) {
            outputBar[i] = ' ';
        }

        printstate(rxBuffer, ryBuffer, rzBuffer, 
                   vxBuffer, vyBuffer, vzBuffer, n, tnow, outFile, filename);

        for (nstep = 0; nstep < mstep; nstep++) {	

            /* Progress Bar Handling/Management */
            int v = round((((double) nstep) / ((double) mstep)) * 100.0);
            int barIndex = (int) round(((double) v / 100.0) * maxBar) + 1;
            if (barIndex < maxBar - 1) {
                outputBar[barIndex] = '#';
            }
            printf("\r%s  %d%%", outputBar, v);
            fflush(stdout);

            /* then call output routine */

            leapstep <<<numBlocks, nThreads>>>(rx, ry, rz,
                                               vx, vy, vz,
                                               n, dt, gm, 0);
            hipDeviceSynchronize();

            /* Copying memory from device to computer */
            hipMemcpy(massBuffer, mass, n*sizeof(float), hipMemcpyDeviceToHost);
            hipMemcpy(rxBuffer, rx, n*sizeof(float), hipMemcpyDeviceToHost);
            hipMemcpy(ryBuffer, ry, n*sizeof(float), hipMemcpyDeviceToHost);
            hipMemcpy(rzBuffer, rz, n*sizeof(float), hipMemcpyDeviceToHost);
            hipMemcpy(vxBuffer, vx, n*sizeof(float), hipMemcpyDeviceToHost);
            hipMemcpy(vyBuffer, vy, n*sizeof(float), hipMemcpyDeviceToHost);
            hipMemcpy(vzBuffer, vz, n*sizeof(float), hipMemcpyDeviceToHost);

            /* take integration step    */
            tnow = tnow + dt;			
            /* and update value of time */

            /* Printing out current positions and velocities */
            if (nstep % nout == 0) {
                printstate(rxBuffer, ryBuffer, rzBuffer, 
                           vxBuffer, vyBuffer, vzBuffer, n, tnow, outFile, filename);
            }
        }
        hipDeviceSynchronize();

        /* Outputting to file */
        printstate(rxBuffer, ryBuffer, rzBuffer, 
                   vxBuffer, vyBuffer, vzBuffer, n, tnow, outFile, filename);
        /* Cleaning up progress bar */
        printf("\r%s  %d%%\n", outputBar, 100);
        fflush(stdout);
    }

    /* Freeing memory */
    hipFree(mass);
    hipFree(rx);
    hipFree(ry);
    hipFree(rz);
    hipFree(vx);
    hipFree(vy);
    hipFree(vz);
    hipFree(gm);
    hipDeviceReset();
}

/*
 * LEAPSTEP: take one step using the leap-from integrator, formulated
 * as a mapping from t to t + dt.  WARNING: this integrator is not
 * accurate unless the timestep dt is fixed from one call to another.
 */

__global__ void leapstep(float rx[], float ry[], float rz[], 
                         float vx[], float vy[], float vz[], 
                         int n, float dt, float gmConst[], int deviceOffset)
{
    int index = deviceOffset + blockIdx.x * blockDim.x + threadIdx.x;
    float3 ac3;

    /* call acceleration code */
    ac3 = accel(rx, ry, rz, n, gmConst, deviceOffset, index);
    __syncthreads();

    /* Applying acceleration to velocity */
    vx[index] = vx[index] + 0.5 * dt * ac3.x;
    vy[index] = vy[index] + 0.5 * dt * ac3.y;
    vz[index] = vz[index] + 0.5 * dt * ac3.z;

    /* Applying velocity to position */
    rx[index] = rx[index] + dt * vx[index];
    ry[index] = ry[index] + dt * vy[index];
    rz[index] = rz[index] + dt * vz[index];

    /* call acceleration code */
    ac3 = accel(rx, ry, rz, n, gmConst, deviceOffset, index);
    __syncthreads();

    vx[index] = vx[index] + 0.5 * dt * ac3.x;
    vy[index] = vy[index] + 0.5 * dt * ac3.y;
    vz[index] = vz[index] + 0.5 * dt * ac3.z;
}

/*
 * ACCEL: compute accelerations for harmonic oscillator(s).
 */

__device__
float3 accel(float* rx, float* ry, float* rz, 
             int n, float gmConst[], int deviceOffset, int index)
{
    float3 ac3 = {0.0f, 0.0f, 0.0f};
    if (index != 0) {
        for (int j = 0; j < n; j++) {
            if (j != index) {
                float distVal = (rx[index]-rx[j])*(rx[index]-rx[j])
                                +(ry[index]-ry[j])*(ry[index]-ry[j])
                                +(rz[index]-rz[j])*(rz[index]-rz[j])
                                + SOFT_FACTOR;
                distVal = distVal * distVal * distVal;
                distVal = 1.0f / sqrtf(distVal);
        
                /* Summing up acceleration */
                ac3.x += -(rx[index]-rx[j])*gmConst[j]*distVal;
                ac3.y += -(ry[index]-ry[j])*gmConst[j]*distVal;
                ac3.z += -(rz[index]-rz[j])*gmConst[j]*distVal;
            }
            __syncthreads();
        }
    }
    return ac3;
}

/*
 * PRINTSTATE: output system state variables.
 */

void printstate(float rx[], float ry[], float rz[],
                float vx[], float vy[], float vz[], 
                int n, float tnow, FILE* outFile, const char * filename)
{
    int i;
    outFile = fopen(filename, "a+");
    for (i = 0; i < n; i++)	{		
        /* Printing out time, particle, position, and velocity */
        fprintf(outFile, 
                "%8.4f\t%4d\t%20.6f\t%20.6f\t%20.6f\t%20.6f\t%20.6f\t%20.6f\n", 
                tnow, i, rx[i], ry[i], rz[i], vx[i], vy[i], vz[i]);
    }
    fclose(outFile);
}


	.text
	.file	"cudaLeapint.hip"
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI0_0:
	.long	0x3a800000                      # float 9.765625E-4
.LCPI0_5:
	.long	0x3ca3d70a                      # float 0.0199999996
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0
.LCPI0_1:
	.zero	16,32
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI0_2:
	.quad	0x4085e00000000000              # double 700
.LCPI0_3:
	.quad	0x4059000000000000              # double 100
.LCPI0_4:
	.quad	0x403e000000000000              # double 30
.LCPI0_6:
	.quad	0x3ff0000000000000              # double 1
	.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	$484248, %rsp                   # imm = 0x76398
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	movq	%rsi, %rbx
	movl	%edi, %r14d
	xorl	%eax, %eax
	movq	%rax, -120(%rbp)
	movq	%rax, -104(%rbp)
	movq	%rax, -96(%rbp)
	movq	%rax, -80(%rbp)
	movq	%rax, -72(%rbp)
	movq	%rax, -64(%rbp)
	movq	%rax, -56(%rbp)
	movq	%rax, -88(%rbp)
	leaq	-484288(%rbp), %rdi
	movl	$60000, %edx                    # imm = 0xEA60
	xorl	%esi, %esi
	callq	memset@PLT
	cmpl	$1, %r14d
	jg	.LBB0_2
# %bb.1:
	movq	(%rbx), %rsi
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	printf
                                        # implicit-def: $r15
	jmp	.LBB0_3
.LBB0_2:
	movq	8(%rbx), %r15
.LBB0_3:
	xorl	%ebx, %ebx
	movl	$.L.str.1, %edi
	movq	%r15, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.2, %esi
	movq	%r15, %rdi
	callq	fopen
	testq	%rax, %rax
	je	.LBB0_30
# %bb.4:
	movq	%rax, %r14
	movq	%r15, -48(%rbp)                 # 8-byte Spill
	leaq	-4288(%rbp), %rdi
	movl	$4096, %esi                     # imm = 0x1000
	movq	%rax, %rdx
	callq	fgets
	testq	%rax, %rax
	je	.LBB0_21
# %bb.5:                                # %.lr.ph.preheader
	xorl	%ebx, %ebx
	leaq	-4288(%rbp), %r15
.LBB0_6:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$.L.str.3, %esi
	movq	%r15, %rdi
	callq	strtok
	testq	%rax, %rax
	je	.LBB0_8
# %bb.7:                                #   in Loop: Header=BB0_6 Depth=1
	movq	%rax, %rdi
	callq	atof
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -424288(%rbp,%rbx,4)
.LBB0_8:                                #   in Loop: Header=BB0_6 Depth=1
	movl	$.L.str.3, %esi
	xorl	%edi, %edi
	callq	strtok
	testq	%rax, %rax
	je	.LBB0_10
# %bb.9:                                #   in Loop: Header=BB0_6 Depth=1
	movq	%rax, %rdi
	callq	atof
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -184288(%rbp,%rbx,4)
.LBB0_10:                               #   in Loop: Header=BB0_6 Depth=1
	movl	$.L.str.3, %esi
	xorl	%edi, %edi
	callq	strtok
	testq	%rax, %rax
	je	.LBB0_12
# %bb.11:                               #   in Loop: Header=BB0_6 Depth=1
	movq	%rax, %rdi
	callq	atof
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -124288(%rbp,%rbx,4)
.LBB0_12:                               #   in Loop: Header=BB0_6 Depth=1
	movl	$.L.str.3, %esi
	xorl	%edi, %edi
	callq	strtok
	testq	%rax, %rax
	je	.LBB0_14
# %bb.13:                               #   in Loop: Header=BB0_6 Depth=1
	movq	%rax, %rdi
	callq	atof
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -64288(%rbp,%rbx,4)
.LBB0_14:                               #   in Loop: Header=BB0_6 Depth=1
	movl	$.L.str.3, %esi
	xorl	%edi, %edi
	callq	strtok
	testq	%rax, %rax
	je	.LBB0_16
# %bb.15:                               #   in Loop: Header=BB0_6 Depth=1
	movq	%rax, %rdi
	callq	atof
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -364288(%rbp,%rbx,4)
.LBB0_16:                               #   in Loop: Header=BB0_6 Depth=1
	movl	$.L.str.3, %esi
	xorl	%edi, %edi
	callq	strtok
	testq	%rax, %rax
	je	.LBB0_18
# %bb.17:                               #   in Loop: Header=BB0_6 Depth=1
	movq	%rax, %rdi
	callq	atof
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -304288(%rbp,%rbx,4)
.LBB0_18:                               #   in Loop: Header=BB0_6 Depth=1
	movl	$.L.str.3, %esi
	xorl	%edi, %edi
	callq	strtok
	testq	%rax, %rax
	je	.LBB0_20
# %bb.19:                               #   in Loop: Header=BB0_6 Depth=1
	movq	%rax, %rdi
	callq	atof
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, -244288(%rbp,%rbx,4)
.LBB0_20:                               #   in Loop: Header=BB0_6 Depth=1
	incq	%rbx
	movq	%r15, %rdi
	movl	$4096, %esi                     # imm = 0x1000
	movq	%r14, %rdx
	callq	fgets
	testq	%rax, %rax
	jne	.LBB0_6
.LBB0_21:                               # %._crit_edge
	movq	%r14, %rdi
	callq	fclose
	xorl	%edi, %edi
	callq	hipSetDevice
	movl	%ebx, %r14d
	shlq	$2, %r14
	leaq	-120(%rbp), %r13
	movq	%r13, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-104(%rbp), %r12
	movq	%r12, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-96(%rbp), %r15
	movq	%r15, %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-80(%rbp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-72(%rbp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-64(%rbp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-56(%rbp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	-88(%rbp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	movq	(%r13), %rdi
	leaq	-424288(%rbp), %rsi
	movl	$1, %r13d
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r12), %rdi
	leaq	-184288(%rbp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r15), %rdi
	leaq	-124288(%rbp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	-80(%rbp), %rax
	movq	(%rax), %rdi
	leaq	-64288(%rbp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	-72(%rbp), %rax
	movq	(%rax), %rdi
	leaq	-364288(%rbp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	-64(%rbp), %rax
	movq	(%rax), %rdi
	leaq	-304288(%rbp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	-56(%rbp), %rax
	movq	(%rax), %rdi
	leaq	-244288(%rbp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%ebx, %ebx
	je	.LBB0_23
# %bb.22:                               # %.lr.ph154.preheader
	leaq	-484288(%rbp), %rdi
	leaq	-424288(%rbp), %rsi
	movq	%r14, %rdx
	callq	memcpy@PLT
.LBB0_23:                               # %._crit_edge155
	movabsq	$4294967296, %rax               # imm = 0x100000000
	movq	%rax, -128(%rbp)                # 8-byte Spill
	movq	-88(%rbp), %rdi
	leaq	-484288(%rbp), %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	cvtsi2ss	%ebx, %xmm0
	mulss	.LCPI0_0(%rip), %xmm0
	cvttss2si	%xmm0, %ebx
	movl	%ebx, %r15d
	shll	$10, %r15d
	leal	1023(%r15), %eax
	cmpl	$2047, %eax                     # imm = 0x7FF
	cmovbl	%r13d, %ebx
	movl	$.L.str.4, %edi
	movl	%ebx, %esi
	movl	$1024, %edx                     # imm = 0x400
	movl	%r15d, %ecx
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.5, %esi
	movq	-48(%rbp), %rdi                 # 8-byte Reload
	callq	strtok
	movq	%rax, %r14
	movq	%rax, %rdi
	callq	strlen@PLT
	movabsq	$32783537688897875, %rcx        # imm = 0x7478742E6D6953
	movq	%rcx, (%r14,%rax)
	movl	$.L.str.7, %esi
	movq	%r14, %rdi
	callq	fopen
	movq	%rsp, -152(%rbp)                # 8-byte Spill
	movq	%rsp, %rax
	leaq	-32(%rax), %rcx
	movq	%rcx, -144(%rbp)                # 8-byte Spill
	movq	%rcx, %rsp
	movb	$124, %cl
	movb	%cl, -32(%rax)
	movb	%cl, -3(%rax)
	leaq	-31(%rax), %rcx
	movq	%rcx, -160(%rbp)                # 8-byte Spill
	movapd	.LCPI0_1(%rip), %xmm0           # xmm0 = [32,32,32,32,32,32,32,32,32,32,32,32,32,32,32,32]
	movupd	%xmm0, -31(%rax)
	movupd	%xmm0, -19(%rax)
	subq	$32, %rsp
	movq	%r14, -136(%rbp)                # 8-byte Spill
	movq	%r14, 16(%rsp)
	movl	%r15d, (%rsp)
	leaq	-184288(%rbp), %rdi
	leaq	-124288(%rbp), %rsi
	leaq	-64288(%rbp), %rdx
	leaq	-364288(%rbp), %rcx
	leaq	-304288(%rbp), %r8
	leaq	-244288(%rbp), %r9
	xorpd	%xmm0, %xmm0
	movss	%xmm0, -48(%rbp)                # 4-byte Spill
	callq	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	addq	$32, %rsp
	btsq	$32, %rbx
	movq	%rbx, -168(%rbp)                # 8-byte Spill
	movq	%r15, -112(%rbp)                # 8-byte Spill
	movslq	%r15d, %r12
	shlq	$2, %r12
	xorpd	%xmm0, %xmm0
	movl	$700, %r13d                     # imm = 0x2BC
	addq	$1024, -128(%rbp)               # 8-byte Folded Spill
                                        # imm = 0x400
.LBB0_24:                               # =>This Inner Loop Header: Depth=1
	movsd	%xmm0, -176(%rbp)               # 8-byte Spill
	divsd	.LCPI0_2(%rip), %xmm0
	movsd	.LCPI0_3(%rip), %xmm1           # xmm1 = mem[0],zero
	mulsd	%xmm1, %xmm0
	callq	round@PLT
	movapd	%xmm0, -192(%rbp)               # 16-byte Spill
	cvttpd2dq	%xmm0, %xmm0
	cvtdq2pd	%xmm0, %xmm0
	divsd	.LCPI0_3(%rip), %xmm0
	mulsd	.LCPI0_4(%rip), %xmm0
	callq	round@PLT
	cvttsd2si	%xmm0, %eax
	cmpl	$27, %eax
	jg	.LBB0_26
# %bb.25:                               #   in Loop: Header=BB0_24 Depth=1
	cltq
	movq	-160(%rbp), %rcx                # 8-byte Reload
	movb	$35, (%rcx,%rax)
.LBB0_26:                               #   in Loop: Header=BB0_24 Depth=1
	cvttsd2si	-192(%rbp), %edx        # 16-byte Folded Reload
	movl	$.L.str.8, %edi
	movq	-144(%rbp), %rsi                # 8-byte Reload
	xorl	%eax, %eax
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	movq	-168(%rbp), %rdi                # 8-byte Reload
	movl	$1, %esi
	movq	-128(%rbp), %rdx                # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_28
# %bb.27:                               #   in Loop: Header=BB0_24 Depth=1
	movq	-104(%rbp), %rdi
	movq	-96(%rbp), %rsi
	movq	-80(%rbp), %rdx
	movq	-72(%rbp), %rcx
	movq	-64(%rbp), %r8
	movq	-56(%rbp), %r9
	subq	$8, %rsp
	movss	.LCPI0_5(%rip), %xmm0           # xmm0 = mem[0],zero,zero,zero
	pushq	$0
	pushq	-88(%rbp)
	pushq	-112(%rbp)                      # 8-byte Folded Reload
	callq	_Z23__device_stub__leapstepPfS_S_S_S_S_ifS_i
	addq	$32, %rsp
.LBB0_28:                               #   in Loop: Header=BB0_24 Depth=1
	callq	hipDeviceSynchronize
	movq	-120(%rbp), %rsi
	leaq	-424288(%rbp), %rdi
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-104(%rbp), %rsi
	leaq	-184288(%rbp), %rdi
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-96(%rbp), %rsi
	leaq	-124288(%rbp), %rdi
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-80(%rbp), %rsi
	leaq	-64288(%rbp), %rdi
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-72(%rbp), %rsi
	leaq	-364288(%rbp), %rbx
	movq	%rbx, %rdi
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-64(%rbp), %rsi
	leaq	-304288(%rbp), %r14
	movq	%r14, %rdi
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	-56(%rbp), %rsi
	leaq	-244288(%rbp), %r15
	movq	%r15, %rdi
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movss	-48(%rbp), %xmm0                # 4-byte Reload
                                        # xmm0 = mem[0],zero,zero,zero
	addss	.LCPI0_5(%rip), %xmm0
	subq	$32, %rsp
	movq	-136(%rbp), %rax                # 8-byte Reload
	movq	%rax, 16(%rsp)
	movq	-112(%rbp), %rax                # 8-byte Reload
	movl	%eax, (%rsp)
	leaq	-184288(%rbp), %rdi
	leaq	-124288(%rbp), %rsi
	leaq	-64288(%rbp), %rdx
	movq	%rbx, %rcx
	movq	%r14, %r8
	movq	%r15, %r9
	movss	%xmm0, -48(%rbp)                # 4-byte Spill
	callq	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	addq	$32, %rsp
	movsd	-176(%rbp), %xmm0               # 8-byte Reload
                                        # xmm0 = mem[0],zero
	addsd	.LCPI0_6(%rip), %xmm0
	decl	%r13d
	jne	.LBB0_24
# %bb.29:
	callq	hipDeviceSynchronize
	subq	$32, %rsp
	movq	-136(%rbp), %rax                # 8-byte Reload
	movq	%rax, 16(%rsp)
	movq	-112(%rbp), %rax                # 8-byte Reload
	movl	%eax, (%rsp)
	leaq	-184288(%rbp), %rdi
	leaq	-124288(%rbp), %rsi
	leaq	-64288(%rbp), %rdx
	leaq	-364288(%rbp), %rcx
	leaq	-304288(%rbp), %r8
	leaq	-244288(%rbp), %r9
	movss	-48(%rbp), %xmm0                # 4-byte Reload
                                        # xmm0 = mem[0],zero,zero,zero
	callq	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	addq	$32, %rsp
	movl	$.L.str.9, %edi
	movq	-144(%rbp), %rsi                # 8-byte Reload
	movl	$100, %edx
	xorl	%eax, %eax
	callq	printf
	movq	stdout(%rip), %rdi
	callq	fflush
	movq	-152(%rbp), %rsp                # 8-byte Reload
	movq	-120(%rbp), %rdi
	callq	hipFree
	movq	-104(%rbp), %rdi
	callq	hipFree
	movq	-96(%rbp), %rdi
	callq	hipFree
	movq	-80(%rbp), %rdi
	callq	hipFree
	movq	-72(%rbp), %rdi
	callq	hipFree
	movq	-64(%rbp), %rdi
	callq	hipFree
	movq	-56(%rbp), %rdi
	callq	hipFree
	movq	-88(%rbp), %rdi
	callq	hipFree
	callq	hipDeviceReset
.LBB0_30:
	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
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc # -- Begin function _Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	.type	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc,@function
_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc: # @_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movss	%xmm0, 8(%rsp)                  # 4-byte Spill
	movq	%r9, 32(%rsp)                   # 8-byte Spill
	movq	%r8, 24(%rsp)                   # 8-byte Spill
	movq	%rcx, 16(%rsp)                  # 8-byte Spill
	movq	%rdx, %r12
	movq	%rsi, %r13
	movq	%rdi, %rbp
	movl	96(%rsp), %r14d
	movq	112(%rsp), %rdi
	movl	$.L.str.10, %esi
	callq	fopen
	movq	%rax, %rbx
	testl	%r14d, %r14d
	jle	.LBB1_3
# %bb.1:                                # %.lr.ph
	xorps	%xmm0, %xmm0
	cvtss2sd	8(%rsp), %xmm0                  # 4-byte Folded Reload
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movl	%r14d, %r15d
	xorl	%r14d, %r14d
.LBB1_2:                                # =>This Inner Loop Header: Depth=1
	xorps	%xmm1, %xmm1
	cvtss2sd	(%rbp,%r14,4), %xmm1
	xorps	%xmm2, %xmm2
	cvtss2sd	(%r13,%r14,4), %xmm2
	xorps	%xmm3, %xmm3
	cvtss2sd	(%r12,%r14,4), %xmm3
	movq	16(%rsp), %rax                  # 8-byte Reload
	xorps	%xmm4, %xmm4
	cvtss2sd	(%rax,%r14,4), %xmm4
	movq	24(%rsp), %rax                  # 8-byte Reload
	xorps	%xmm5, %xmm5
	cvtss2sd	(%rax,%r14,4), %xmm5
	movq	32(%rsp), %rax                  # 8-byte Reload
	xorps	%xmm6, %xmm6
	cvtss2sd	(%rax,%r14,4), %xmm6
	movl	$.L.str.11, %esi
	movq	%rbx, %rdi
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movl	%r14d, %edx
	movb	$7, %al
	callq	fprintf
	incq	%r14
	cmpq	%r14, %r15
	jne	.LBB1_2
.LBB1_3:                                # %._crit_edge
	movq	%rbx, %rdi
	addq	$40, %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
	jmp	fclose                          # TAILCALL
.Lfunc_end1:
	.size	_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc, .Lfunc_end1-_Z10printstatePfS_S_S_S_S_ifP8_IO_FILEPKc
	.cfi_endproc
                                        # -- End function
	.globl	_Z23__device_stub__leapstepPfS_S_S_S_S_ifS_i # -- Begin function _Z23__device_stub__leapstepPfS_S_S_S_S_ifS_i
	.type	_Z23__device_stub__leapstepPfS_S_S_S_S_ifS_i,@function
_Z23__device_stub__leapstepPfS_S_S_S_S_ifS_i: # @_Z23__device_stub__leapstepPfS_S_S_S_S_ifS_i
	.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	$192, %rsp
	.cfi_def_cfa_offset 240
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	72(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	64(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	56(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	48(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	32(%rsp), %r8
	movq	%r9, (%r8)
	leaq	12(%rsp), %r9
	movss	%xmm0, (%r9)
	leaq	112(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	240(%rsp), %rax
	movq	%rax, 48(%rbx)
	movq	%r9, 56(%rbx)
	leaq	248(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	256(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	96(%rsp), %r14
	leaq	80(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z8leapstepPfS_S_S_S_S_ifS_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$208, %rsp
	.cfi_adjust_cfa_offset -208
	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_end2:
	.size	_Z23__device_stub__leapstepPfS_S_S_S_S_ifS_i, .Lfunc_end2-_Z23__device_stub__leapstepPfS_S_S_S_S_ifS_i
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8leapstepPfS_S_S_S_S_ifS_i, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB4_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Usage: %s [N-Body File]\n"
	.size	.L.str, 25

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Reading values in from %s.\n"
	.size	.L.str.1, 28

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

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

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"numBlocks: %d\nnThreads: %d\n\nParticles: %d\n"
	.size	.L.str.4, 43

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

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

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"\r%s  %d%%"
	.size	.L.str.8, 10

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

	.type	.L.str.9,@object                # @.str.9
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.9:
	.asciz	"\r%s  %d%%\n"
	.size	.L.str.9, 11

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"a+"
	.size	.L.str.10, 3

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"%8.4f\t%4d\t%20.6f\t%20.6f\t%20.6f\t%20.6f\t%20.6f\t%20.6f\n"
	.size	.L.str.11, 53

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z8leapstepPfS_S_S_S_S_ifS_i"
	.size	.L__unnamed_1, 29

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

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

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

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z8leapstepPfS_S_S_S_S_ifS_i ; -- Begin function _Z8leapstepPfS_S_S_S_S_ifS_i
	.globl	_Z8leapstepPfS_S_S_S_S_ifS_i
	.p2align	8
	.type	_Z8leapstepPfS_S_S_S_S_ifS_i,@function
_Z8leapstepPfS_S_S_S_S_ifS_i:           ; @_Z8leapstepPfS_S_S_S_S_ifS_i
; %bb.0:
	s_clause 0x3
	s_load_b32 s2, s[0:1], 0x54
	s_load_b32 s3, s[0:1], 0x40
	s_load_b256 s[4:11], s[0:1], 0x20
	s_load_b256 s[16:23], s[0:1], 0x0
	v_mov_b32_e32 v5, 0
	v_mov_b32_e32 v6, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	v_dual_mov_b32 v8, v6 :: v_dual_mov_b32 v7, v5
	v_dual_mov_b32 v12, v6 :: v_dual_mov_b32 v11, v5
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_mul_i32 s15, s15, s2
	s_cmp_gt_i32 s8, 0
	v_add3_u32 v0, s15, s3, v0
	s_cselect_b32 s0, -1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v1, 31, v0
	v_cmp_ne_u32_e32 vcc_lo, 0, v0
	v_lshlrev_b64 v[9:10], 2, v[0:1]
	s_and_b32 s1, s0, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s28, s1
	s_cbranch_execz .LBB0_6
; %bb.1:                                ; %.lr.ph.i
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, s16, v9
	v_add_co_ci_u32_e32 v2, vcc_lo, s17, v10, vcc_lo
	v_add_co_u32 v3, vcc_lo, s18, v9
	v_add_co_ci_u32_e32 v4, vcc_lo, s19, v10, vcc_lo
	v_add_co_u32 v5, vcc_lo, s20, v9
	v_add_co_ci_u32_e32 v6, vcc_lo, s21, v10, vcc_lo
	v_dual_mov_b32 v7, 0 :: v_dual_mov_b32 v8, v0
	v_dual_mov_b32 v12, 0 :: v_dual_mov_b32 v11, 0
	s_mov_b32 s2, 0x85f06f69
	s_mov_b32 s3, 0x3fd554c9
	s_mov_b64 s[12:13], s[10:11]
	s_mov_b64 s[14:15], s[20:21]
	s_mov_b64 s[24:25], s[18:19]
	s_mov_b64 s[26:27], s[16:17]
	s_mov_b32 s29, s8
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	s_mov_b32 s30, exec_lo
	v_cmpx_ne_u32_e32 0, v8
	s_cbranch_execz .LBB0_4
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	global_load_b32 v13, v[3:4], off
	global_load_b32 v14, v[1:2], off
	global_load_b32 v15, v[5:6], off
	s_load_b32 s0, s[24:25], 0x0
	s_load_b32 s31, s[26:27], 0x0
	s_load_b32 s33, s[14:15], 0x0
	s_waitcnt vmcnt(1) lgkmcnt(0)
	v_dual_subrev_f32 v16, s0, v13 :: v_dual_subrev_f32 v17, s31, v14
	s_waitcnt vmcnt(0)
	v_subrev_f32_e32 v15, s33, v15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v13, v16, v16
	v_fmac_f32_e32 v13, v17, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v13, v15, v15
	v_cvt_f64_f32_e32 v[13:14], v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[13:14], v[13:14], s[2:3]
	v_cvt_f32_f64_e32 v13, v[13:14]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v14, v13, v13
	v_mul_f32_e32 v13, v14, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v14, 0x4f800000, v13
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v13
	v_cndmask_b32_e32 v13, v13, v14, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v14, v13
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v18, -1, v14
	v_add_nc_u32_e32 v19, 1, v14
	v_fma_f32 v20, -v18, v14, v13
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v21, -v19, v14, v13
	v_cmp_ge_f32_e64 s0, 0, v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v14, v14, v18, s0
	v_cmp_lt_f32_e64 s0, 0, v21
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v14, v14, v19, s0
	s_load_b32 s0, s[12:13], 0x0
	v_mul_f32_e32 v18, 0x37800000, v14
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v14, v14, v18, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v13, 0x260
	v_cndmask_b32_e32 v13, v14, v13, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_div_scale_f32 v14, null, v13, v13, 1.0
	v_div_scale_f32 v20, vcc_lo, 1.0, v13, 1.0
	s_waitcnt lgkmcnt(0)
	v_mul_f32_e32 v17, s0, v17
	v_rcp_f32_e32 v18, v14
	v_mul_f32_e32 v15, s0, v15
	s_waitcnt_depctr 0xfff
	v_fma_f32 v19, -v14, v18, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v18, v19, v18
	v_mul_f32_e32 v19, v20, v18
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v21, -v14, v19, v20
	v_fmac_f32_e32 v19, v21, v18
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v14, -v14, v19, v20
	v_div_fmas_f32 v14, v14, v18, v19
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_fixup_f32 v13, v14, v13, 1.0
	v_mul_f32_e32 v14, s0, v16
	v_fma_f32 v7, -v17, v13, v7
	s_delay_alu instid0(VALU_DEP_2)
	v_fma_f32 v12, -v14, v13, v12
	v_fma_f32 v11, -v15, v13, v11
.LBB0_4:                                ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s30
	s_add_i32 s29, s29, -1
	s_add_u32 s26, s26, 4
	s_addc_u32 s27, s27, 0
	s_add_u32 s24, s24, 4
	s_addc_u32 s25, s25, 0
	s_add_u32 s14, s14, 4
	s_addc_u32 s15, s15, 0
	v_add_nc_u32_e32 v8, -1, v8
	s_add_u32 s12, s12, 4
	s_addc_u32 s13, s13, 0
	s_cmp_lg_u32 s29, 0
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB0_2
; %bb.5:                                ; %_Z5accelPfS_S_iS_ii.exit.loopexit
	v_cvt_f64_f32_e32 v[5:6], v7
	v_cvt_f64_f32_e32 v[7:8], v12
	v_cvt_f64_f32_e32 v[11:12], v11
.LBB0_6:                                ; %Flow149
	s_or_b32 exec_lo, exec_lo, s28
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_u32 v1, vcc_lo, s22, v9
	v_add_co_ci_u32_e32 v2, vcc_lo, s23, v10, vcc_lo
	s_barrier
	buffer_gl0_inv
	global_load_b32 v13, v[1:2], off
	v_cvt_f64_f32_e32 v[3:4], s9
	v_mov_b32_e32 v19, 0
	v_mov_b32_e32 v20, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_dual_mov_b32 v17, v19 :: v_dual_mov_b32 v18, v20
	v_mul_f64 v[3:4], v[3:4], 0.5
	s_waitcnt vmcnt(0)
	v_cvt_f64_f32_e32 v[13:14], v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[5:6], v[3:4], v[5:6], v[13:14]
	v_cvt_f32_f64_e32 v13, v[5:6]
	v_add_co_u32 v5, vcc_lo, s4, v9
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v10, vcc_lo
	global_store_b32 v[1:2], v13, off
	global_load_b32 v13, v[5:6], off
	s_waitcnt vmcnt(0)
	v_cvt_f64_f32_e32 v[13:14], v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[7:8], v[3:4], v[7:8], v[13:14]
	v_cvt_f32_f64_e32 v13, v[7:8]
	v_add_co_u32 v7, vcc_lo, s6, v9
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v10, vcc_lo
	global_store_b32 v[5:6], v13, off
	global_load_b32 v13, v[7:8], off
	s_waitcnt vmcnt(0)
	v_cvt_f64_f32_e32 v[13:14], v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[11:12], v[3:4], v[11:12], v[13:14]
	v_cvt_f32_f64_e32 v13, v[11:12]
	v_add_co_u32 v11, vcc_lo, s16, v9
	v_add_co_ci_u32_e32 v12, vcc_lo, s17, v10, vcc_lo
	global_store_b32 v[7:8], v13, off
	global_load_b32 v15, v[11:12], off
	global_load_b32 v13, v[1:2], off
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v15, s9, v13
	v_add_co_u32 v13, vcc_lo, s18, v9
	v_add_co_ci_u32_e32 v14, vcc_lo, s19, v10, vcc_lo
	global_store_b32 v[11:12], v15, off
	global_load_b32 v15, v[13:14], off
	global_load_b32 v16, v[5:6], off
	v_add_co_u32 v9, vcc_lo, s20, v9
	v_add_co_ci_u32_e32 v10, vcc_lo, s21, v10, vcc_lo
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v15, s9, v16
	global_store_b32 v[13:14], v15, off
	global_load_b32 v21, v[9:10], off
	global_load_b32 v15, v[7:8], off
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v21, s9, v15
	v_dual_mov_b32 v15, v19 :: v_dual_mov_b32 v16, v20
	global_store_b32 v[9:10], v21, off
	s_and_saveexec_b32 s4, s1
	s_cbranch_execz .LBB0_12
; %bb.7:                                ; %.lr.ph.i90.preheader
	v_dual_mov_b32 v15, 0 :: v_dual_mov_b32 v16, 0
	v_mov_b32_e32 v17, 0
	v_mov_b32_e32 v21, 0
	s_mov_b32 s2, 0x85f06f69
	s_mov_b32 s3, 0x3fd554c9
.LBB0_8:                                ; %.lr.ph.i90
                                        ; =>This Inner Loop Header: Depth=1
	s_mov_b32 s1, exec_lo
	v_cmpx_ne_u32_e32 0, v0
	s_cbranch_execz .LBB0_10
; %bb.9:                                ;   in Loop: Header=BB0_8 Depth=1
	global_load_b32 v18, v[13:14], off
	s_clause 0x1
	global_load_b32 v19, v16, s[18:19]
	global_load_b32 v20, v16, s[16:17]
	global_load_b32 v22, v[11:12], off
	global_load_b32 v23, v[9:10], off
	s_clause 0x1
	global_load_b32 v24, v16, s[20:21]
	global_load_b32 v25, v16, s[10:11]
	s_waitcnt vmcnt(5)
	v_sub_f32_e32 v26, v18, v19
	s_waitcnt vmcnt(3)
	v_sub_f32_e32 v20, v22, v20
	s_waitcnt vmcnt(1)
	v_sub_f32_e32 v22, v23, v24
	v_mul_f32_e32 v18, v26, v26
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v18, v20, v20
	v_fmac_f32_e32 v18, v22, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_f64_f32_e32 v[18:19], v18
	v_add_f64 v[18:19], v[18:19], s[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_f32_f64_e32 v18, v[18:19]
	v_mul_f32_e32 v19, v18, v18
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v22, v25, v22
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v18, v19, v18
	v_mul_f32_e32 v19, 0x4f800000, v18
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v18
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v18, v18, v19, vcc_lo
	v_sqrt_f32_e32 v19, v18
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v23, -1, v19
	v_add_nc_u32_e32 v24, 1, v19
	v_mul_f32_e32 v20, v25, v20
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v27, -v23, v19, v18
	v_fma_f32 v28, -v24, v19, v18
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s0, 0, v27
	v_cndmask_b32_e64 v19, v19, v23, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s0, 0, v28
	v_cndmask_b32_e64 v19, v19, v24, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v23, 0x37800000, v19
	v_cndmask_b32_e32 v19, v19, v23, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v18, 0x260
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v18, v19, v18, vcc_lo
	v_div_scale_f32 v19, null, v18, v18, 1.0
	v_div_scale_f32 v27, vcc_lo, 1.0, v18, 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v23, v19
	s_waitcnt_depctr 0xfff
	v_fma_f32 v24, -v19, v23, 1.0
	v_fmac_f32_e32 v23, v24, v23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v24, v27, v23
	v_fma_f32 v28, -v19, v24, v27
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v24, v28, v23
	v_fma_f32 v19, -v19, v24, v27
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v19, v19, v23, v24
	v_div_fixup_f32 v18, v19, v18, 1.0
	v_mul_f32_e32 v19, v25, v26
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v15, -v20, v18, v15
	v_fma_f32 v17, -v19, v18, v17
	v_fma_f32 v21, -v22, v18, v21
.LBB0_10:                               ;   in Loop: Header=BB0_8 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_add_i32 s8, s8, -1
	s_add_u32 s16, s16, 4
	s_addc_u32 s17, s17, 0
	s_add_u32 s18, s18, 4
	s_addc_u32 s19, s19, 0
	s_add_u32 s20, s20, 4
	s_addc_u32 s21, s21, 0
	v_add_nc_u32_e32 v0, -1, v0
	s_add_u32 s10, s10, 4
	s_addc_u32 s11, s11, 0
	s_cmp_lg_u32 s8, 0
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB0_8
; %bb.11:                               ; %_Z5accelPfS_S_iS_ii.exit100.loopexit
	v_cvt_f64_f32_e32 v[19:20], v15
	v_cvt_f64_f32_e32 v[17:18], v17
	v_cvt_f64_f32_e32 v[15:16], v21
.LBB0_12:                               ; %Flow
	s_or_b32 exec_lo, exec_lo, s4
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	global_load_b32 v0, v[1:2], off
	s_waitcnt vmcnt(0)
	v_cvt_f64_f32_e32 v[9:10], v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[9:10], v[3:4], v[19:20], v[9:10]
	v_cvt_f32_f64_e32 v0, v[9:10]
	global_store_b32 v[1:2], v0, off
	global_load_b32 v0, v[5:6], off
	s_waitcnt vmcnt(0)
	v_cvt_f64_f32_e32 v[0:1], v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[0:1], v[3:4], v[17:18], v[0:1]
	v_cvt_f32_f64_e32 v0, v[0:1]
	global_store_b32 v[5:6], v0, off
	global_load_b32 v0, v[7:8], off
	s_waitcnt vmcnt(0)
	v_cvt_f64_f32_e32 v[0:1], v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[0:1], v[3:4], v[15:16], v[0:1]
	v_cvt_f32_f64_e32 v0, v[0:1]
	global_store_b32 v[7:8], v0, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8leapstepPfS_S_S_S_S_ifS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 328
		.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 29
		.amdhsa_next_free_sgpr 34
		.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	_Z8leapstepPfS_S_S_S_S_ifS_i, .Lfunc_end0-_Z8leapstepPfS_S_S_S_S_ifS_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1760
; NumSgprs: 36
; NumVgprs: 29
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 4
; VGPRBlocks: 3
; NumSGPRsForWavesPerEU: 36
; NumVGPRsForWavesPerEU: 29
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .offset:         52
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .offset:         64
        .size:           4
        .value_kind:     by_value
      - .offset:         72
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         76
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         80
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         84
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         86
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         88
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         90
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         92
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         94
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         136
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 328
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8leapstepPfS_S_S_S_S_ifS_i
    .private_segment_fixed_size: 0
    .sgpr_count:     36
    .sgpr_spill_count: 0
    .symbol:         _Z8leapstepPfS_S_S_S_S_ifS_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     29
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
