/*
Author: Jason He
Version: 1.0 20210605 Serial version.
Version: 2.0 20210607 CUDA version.
*/

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

/*
This program solve the Stokes Flow problem- 36.1.2 in the textbook.

Inputs: 
	N, number of grid points in each direction.
	mu, double precision viscosity.
	P, double precision pressure drop.
	w, double precision relaxation parameter.
    tol, double precision error tolerance.
	K, maximum number of iterations.

Outputs:
	print arguments
	output final grid value into "StokesU.out", "StokesV.out" and "StokesP.out"
	print run time and output it to runtime.dat

E.g.
$ ./stokes 128 1 1 0.4 1e-9 100000

*/

//declare global variables
const int threadsPerBlock = 1024;//for NVIDIA TESLA K20c maximum is 1024
/*the findmaxres kernel demand that threadsPerBlock should be a power of 2.
Also in the three update kernels, each block compare and store the residuals into
the global array dev_res[threadsPerBlock], which only work with maximum
threadsPerBlock=1024, since in this case only one block runs at a time.*/
__device__ __constant__ int dev_N;
__device__ __constant__ double dev_mu;
__device__ __constant__ double dev_P;
__device__ __constant__ double dev_w;
__device__ __constant__ double dev_dx;

//this kernel set values of 1d length-n double array to 0
__global__ void init(double* a, int n){
    int k = threadIdx.x + blockIdx.x*blockDim.x;
    if(k < n){
        a[k] = 0.;
    }
}

//the three update functions update the grid values and return the maximum residual
__global__ void updateU(double* dev_u, double* dev_p, double* dev_res, int color){
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < dev_N*(dev_N-1)){
        int j = i/(dev_N-1);
        int k = i%(dev_N-1);
        if((j+k)%2==color){
            double r;
            int jp = i + dev_N-1;//(j+1)*(dev_N-1) + k;
            int jm = i - dev_N+1;//(j-1)*(dev_N-1) + k;
            int kp = i + 1;//j*(dev_N-1) + k+1;
            int km = i - 1;//j*(dev_N-1) + k-1;
            int pi = j*(dev_N-1) + k;
            int pm = pi - dev_N+1;//(j-1)*(dev_N-1) + k;
            if(j==0){
                if(k==0){
                    r = dev_mu*(dev_u[jp]+dev_u[kp]-4*dev_u[i])
                        -dev_dx*2*(dev_p[pi]-dev_P);
                }
                else if(k<dev_N-2){
                    r = dev_mu*(dev_u[jp]+dev_u[kp]+dev_u[km]-3*dev_u[i])
                        -dev_dx*2*(dev_p[pi]-dev_P);
                }
                else{
                    r = dev_mu*(dev_u[jp]+dev_u[km]-4*dev_u[i])
                        -dev_dx*2*(dev_p[pi]-dev_P);
                }
            }
            else if(j<dev_N-1){
                if(k==0){
                    r = dev_mu*(dev_u[jp]+dev_u[jm]+dev_u[kp]-5*dev_u[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
                else if(k<dev_N-2){
                    r = dev_mu*(dev_u[jp]+dev_u[jm]+dev_u[kp]+dev_u[km]-4*dev_u[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
                else{
                    r = dev_mu*(dev_u[jp]+dev_u[jm]+dev_u[km]-5*dev_u[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
            }
            else{
                if(k==0){
                    r = dev_mu*(dev_u[jm]+dev_u[kp]-4*dev_u[i])
                        -dev_dx*2*(0-dev_p[pm]);
                }
                else if(k<dev_N-2){
                    r = dev_mu*(dev_u[jm]+dev_u[kp]+dev_u[km]-3*dev_u[i])
                        -dev_dx*2*(0-dev_p[pm]);
                }
                else{
                    r = dev_mu*(dev_u[jm]+dev_u[km]-4*dev_u[i])
                        -dev_dx*2*(0-dev_p[pm]);
                }
            }
            dev_u[i] += dev_w*r;
            dev_res[threadIdx.x] = max(fabs(r), dev_res[threadIdx.x]);
        }
    }
}

__global__ void updateV(double* dev_v, double* dev_p, double* dev_res, int color){
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < (dev_N-1)*dev_N){
        int k = i%dev_N;
        if(0<k && k<dev_N-1){
            int j = i/dev_N;
            if((j+k)%2==color){
                double r;
                int jp = i + dev_N;//(j+1)*dev_N + k;
                int jm = i - dev_N;//(j-1)*dev_N + k;
                int kp = i + 1;//j*dev_N + k+1;
                int km = i - 1;//j*dev_N + k-1;
                int pi = j*(dev_N-1) + k;
                int pm = pi - 1;//j*(dev_N-1) + k-1;
                if(j==0){
                    r = dev_mu*(dev_v[kp]+dev_v[km]+dev_v[jp]-3*dev_v[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
                else if(j<dev_N-2){
                    r = dev_mu*(dev_v[kp]+dev_v[km]+dev_v[jp]+dev_v[jm]-4*dev_v[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
                else{
                    r = dev_mu*(dev_v[kp]+dev_v[km]+dev_v[jm]-3*dev_v[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
                dev_v[i] += dev_w*r;
                dev_res[threadIdx.x] = max(fabs(r), dev_res[threadIdx.x]);
            }
        }
    }
}

__global__ void updateP(double* dev_u, double* dev_v, double* dev_p, double* dev_res){
    
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < (dev_N-1)*(dev_N-1)){
        int j = i/(dev_N-1);
        int k = i%(dev_N-1);
        int up = (j+1)*(dev_N-1) + k;
        int um = j*(dev_N-1) + k;
        int vp = j*dev_N + k+1;
        int vm = j*dev_N + k;
        double r;
        r = -(dev_u[up]-dev_u[um])-(dev_v[vp]-dev_v[vm]);
        dev_p[i] += dev_w*r;//update all p's
        dev_res[threadIdx.x] = max(fabs(r), dev_res[threadIdx.x]);//not comfortable with this
    }
}

//this kernel find the max value
__global__ void findmaxres(double* dev_res, int i){
    int k = threadIdx.x + blockIdx.x*blockDim.x;
    if(k%(2*i)==0){
        dev_res[k] = max(dev_res[k], dev_res[k+i]);
    }
}

int main(int argc, char* argv[]){
    //Choose gpu device
    cudaDeviceProp prop;
    int dev;
    memset(&prop, 0, sizeof(cudaDeviceProp));
    prop.multiProcessorCount = 13;
    cudaChooseDevice(&dev, &prop);
    cudaSetDevice(dev);

    //load input parameters
    int argi = 0;
    int N = atoi(argv[++argi]);
    double mu = atof(argv[++argi]);
    double P = atof(argv[++argi]);
    double w = atof(argv[++argi]);
    double tol = atof(argv[++argi]);
    int K = atoi(argv[++argi]);
    printf( "N = %d\nmu = %lf\nP = %lf\nw = %lf\ntau = %e\nK = %d\n",
            N, mu, P, w, tol, K);
    
    //calculate parameters for the iteration equation
    double dx = 1./(N-1);
    //calculate parameters for blocks and threads
    const int blocksPerGrid = (N*(N-1) + threadsPerBlock - 1)/threadsPerBlock;
    printf("BlocksPerGrid = %d\nThreadsPerBlock = %d\n", blocksPerGrid, threadsPerBlock);

    //copy parameters to device
    cudaMemcpyToSymbol(dev_N, &N, sizeof(double));
    cudaMemcpyToSymbol(dev_mu, &mu, sizeof(double));
    cudaMemcpyToSymbol(dev_P, &P, sizeof(double));
    cudaMemcpyToSymbol(dev_w, &w, sizeof(double));
    cudaMemcpyToSymbol(dev_dx, &dx, sizeof(double));

    //initialize value of the grid on device
    double *dev_u, *dev_v, *dev_p;
    cudaMalloc((void**)&dev_u, N*(N-1)*sizeof(double));
    cudaMalloc((void**)&dev_v, (N-1)*N*sizeof(double));
    cudaMalloc((void**)&dev_p, (N-1)*(N-1)*sizeof(double));
    init<<<blocksPerGrid, threadsPerBlock>>>(dev_u, N*(N-1));
    init<<<blocksPerGrid, threadsPerBlock>>>(dev_v, (N-1)*N);
    init<<<blocksPerGrid, threadsPerBlock>>>(dev_p, (N-1)*(N-1));
    
    //initialize grids of residual
    double *dev_res;
    cudaMalloc((void**)&dev_res, threadsPerBlock*sizeof(double));

    //initialize parameters for iterations
    int iter=0;//number of iteration
    double maxres=1.0;//maximum residual of u,v,p on every grid points
    float runtime;//record runtime
    clock_t t;
    t = clock();
    //main loop
    while(iter<K && maxres>tol){
        init<<<1, threadsPerBlock>>>(dev_res, threadsPerBlock);
        updateU<<<blocksPerGrid, threadsPerBlock>>>(dev_u, dev_p, dev_res, 0);
        updateU<<<blocksPerGrid, threadsPerBlock>>>(dev_u, dev_p, dev_res, 1);
        updateV<<<blocksPerGrid, threadsPerBlock>>>(dev_v, dev_p, dev_res, 0);
        updateV<<<blocksPerGrid, threadsPerBlock>>>(dev_v, dev_p, dev_res, 1);
        updateP<<<blocksPerGrid, threadsPerBlock>>>(dev_u, dev_v, dev_p, dev_res);
        for(int i=1; i<threadsPerBlock; i*=2)
            findmaxres<<<1, threadsPerBlock>>>(dev_res, i);
        cudaMemcpy(&maxres, dev_res, sizeof(double), cudaMemcpyDeviceToHost);
        iter++;
    }
    runtime = (float)(clock()-t)/CLOCKS_PER_SEC;
    FILE* tfile = fopen("runtime.dat", "a");
    fprintf(tfile, "%d %f\n", N, runtime/iter);
    fclose(tfile);
    printf("Runtime = %f seconds\nOne iteration runtime = %f seconds\n",
            runtime, runtime/iter);

    //print iterations and residuals
    printf("number of iterations = %d\nresidual = %e\n", iter, maxres);
    
    //initialize the grids on the host
    double *u = (double*)malloc(N*(N-1)*sizeof(double));
    double *v = (double*)malloc((N-1)*N*sizeof(double));
    double *p = (double*)malloc((N-1)*(N-1)*sizeof(double));

    //output to files
    cudaMemcpy(u, dev_u, N*(N-1)*sizeof(double), cudaMemcpyDeviceToHost);
    cudaMemcpy(v, dev_v, (N-1)*N*sizeof(double), cudaMemcpyDeviceToHost);
    cudaMemcpy(p, dev_p, (N-1)*(N-1)*sizeof(double), cudaMemcpyDeviceToHost);
    FILE* ufile = fopen("stokesU.out","w");
    fwrite(u, sizeof(double), N*(N-1), ufile);
    fclose(ufile);
    FILE* vfile = fopen("stokesV.out","w");
    fwrite(v, sizeof(double), (N-1)*N, vfile);
    fclose(vfile);
    FILE* pfile = fopen("stokesP.out","w");
    fwrite(p, sizeof(double), (N-1)*(N-1), pfile);
    fclose(pfile);

    //free memories
    free(u);
    free(v);
    free(p);
    cudaFree(dev_u);
    cudaFree(dev_v);
    cudaFree(dev_p);
    cudaFree(dev_res);

    return 0;
}
	.file	"tmpxft_00201955_00000000-6_stokes.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z24__device_stub__Z4initPdiPdi
	.type	_Z24__device_stub__Z4initPdiPdi, @function
_Z24__device_stub__Z4initPdiPdi:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movl	%esi, 4(%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z4initPdi(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L2:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z24__device_stub__Z4initPdiPdi, .-_Z24__device_stub__Z4initPdiPdi
	.globl	_Z4initPdi
	.type	_Z4initPdi, @function
_Z4initPdi:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z24__device_stub__Z4initPdiPdi
	.cfi_endproc
.LFE2053:
	.size	_Z4initPdi, .-_Z4initPdi
	.globl	_Z31__device_stub__Z7updateUPdS_S_iPdS_S_i
	.type	_Z31__device_stub__Z7updateUPdS_S_iPdS_S_i, @function
_Z31__device_stub__Z7updateUPdS_S_iPdS_S_i:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z7updateUPdS_S_i(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L8:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z31__device_stub__Z7updateUPdS_S_iPdS_S_i, .-_Z31__device_stub__Z7updateUPdS_S_iPdS_S_i
	.globl	_Z7updateUPdS_S_i
	.type	_Z7updateUPdS_S_i, @function
_Z7updateUPdS_S_i:
.LFB2055:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z7updateUPdS_S_iPdS_S_i
	.cfi_endproc
.LFE2055:
	.size	_Z7updateUPdS_S_i, .-_Z7updateUPdS_S_i
	.globl	_Z31__device_stub__Z7updateVPdS_S_iPdS_S_i
	.type	_Z31__device_stub__Z7updateVPdS_S_iPdS_S_i, @function
_Z31__device_stub__Z7updateVPdS_S_iPdS_S_i:
.LFB2056:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z7updateVPdS_S_i(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L13:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2056:
	.size	_Z31__device_stub__Z7updateVPdS_S_iPdS_S_i, .-_Z31__device_stub__Z7updateVPdS_S_iPdS_S_i
	.globl	_Z7updateVPdS_S_i
	.type	_Z7updateVPdS_S_i, @function
_Z7updateVPdS_S_i:
.LFB2057:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z7updateVPdS_S_iPdS_S_i
	.cfi_endproc
.LFE2057:
	.size	_Z7updateVPdS_S_i, .-_Z7updateVPdS_S_i
	.globl	_Z32__device_stub__Z7updatePPdS_S_S_PdS_S_S_
	.type	_Z32__device_stub__Z7updatePPdS_S_S_PdS_S_S_, @function
_Z32__device_stub__Z7updatePPdS_S_S_PdS_S_S_:
.LFB2058:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z7updatePPdS_S_S_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L18:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2058:
	.size	_Z32__device_stub__Z7updatePPdS_S_S_PdS_S_S_, .-_Z32__device_stub__Z7updatePPdS_S_S_PdS_S_S_
	.globl	_Z7updatePPdS_S_S_
	.type	_Z7updatePPdS_S_S_, @function
_Z7updatePPdS_S_S_:
.LFB2059:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z7updatePPdS_S_S_PdS_S_S_
	.cfi_endproc
.LFE2059:
	.size	_Z7updatePPdS_S_S_, .-_Z7updatePPdS_S_S_
	.globl	_Z31__device_stub__Z10findmaxresPdiPdi
	.type	_Z31__device_stub__Z10findmaxresPdiPdi, @function
_Z31__device_stub__Z10findmaxresPdiPdi:
.LFB2060:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movl	%esi, 4(%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z10findmaxresPdi(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L23:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L25
	call	__stack_chk_fail@PLT
.L25:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2060:
	.size	_Z31__device_stub__Z10findmaxresPdiPdi, .-_Z31__device_stub__Z10findmaxresPdiPdi
	.globl	_Z10findmaxresPdi
	.type	_Z10findmaxresPdi, @function
_Z10findmaxresPdi:
.LFB2061:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z10findmaxresPdiPdi
	.cfi_endproc
.LFE2061:
	.size	_Z10findmaxresPdi, .-_Z10findmaxresPdi
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"N = %d\nmu = %lf\nP = %lf\nw = %lf\ntau = %e\nK = %d\n"
.LC2:
	.string	"BlocksPerGrid = %d\nThreadsPerBlock = %d\n"
.LC4:
	.string	"a"
.LC5:
	.string	"runtime.dat"
.LC6:
	.string	"%d %f\n"
.LC7:
	.string	"Runtime = %f seconds\nOne iteration runtime = %f seconds\n"
.LC8:
	.string	"number of iterations = %d\nresidual = %e\n"
.LC9:
	.string	"w"
.LC10:
	.string	"stokesU.out"
.LC11:
	.string	"stokesV.out"
.LC12:
	.string	"stokesP.out"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$258, %ecx
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	subq	$1192, %rsp
	.cfi_def_cfa_offset 1248
	movq	%fs:40, %rax
	movq	%rax, 1176(%rsp)
	xorl	%eax, %eax
	leaq	144(%rsp), %rsi
	movq	%rsi, %rdi
	rep stosl
	leaq	40(%rsp), %rdi
	movl	$13, 532(%rsp)
	call	cudaChooseDevice@PLT
	movl	40(%rsp), %edi
	call	cudaSetDevice@PLT
	movq	8(%rbx), %rdi
	call	atoi@PLT
	movq	16(%rbx), %rdi
	movl	%eax, %ebp
	movl	%eax, 44(%rsp)
	call	atof@PLT
	movq	24(%rbx), %rdi
	movq	%xmm0, %r12
	movsd	%xmm0, 48(%rsp)
	call	atof@PLT
	movq	32(%rbx), %rdi
	movsd	%xmm0, 56(%rsp)
	movsd	%xmm0, 24(%rsp)
	call	atof@PLT
	movq	40(%rbx), %rdi
	movsd	%xmm0, 64(%rsp)
	movsd	%xmm0, 16(%rsp)
	call	atof@PLT
	movq	48(%rbx), %rdi
	movsd	%xmm0, 8(%rsp)
	call	atoi@PLT
	movsd	8(%rsp), %xmm3
	movsd	16(%rsp), %xmm2
	movl	%ebp, %edx
	movsd	24(%rsp), %xmm1
	movl	%eax, %ecx
	movq	%r12, %xmm0
	movl	%eax, %r13d
	leaq	.LC0(%rip), %rsi
	movl	$2, %edi
	movb	$4, %al
	call	__printf_chk@PLT
	movl	44(%rsp), %eax
	movsd	.LC1(%rip), %xmm0
	movl	$1024, %ecx
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	leal	-1(%rax), %edx
	cvtsi2sdl	%edx, %xmm1
	imull	%edx, %eax
	divsd	%xmm1, %xmm0
	addl	$1023, %eax
	cltd
	idivl	%ecx
	movl	%eax, %edx
	movl	%eax, %ebx
	xorl	%eax, %eax
	movsd	%xmm0, 72(%rsp)
	call	__printf_chk@PLT
	xorl	%ecx, %ecx
	leaq	44(%rsp), %rsi
	movl	$1, %r8d
	movl	$8, %edx
	leaq	_ZL5dev_N(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	xorl	%ecx, %ecx
	leaq	48(%rsp), %rsi
	movl	$1, %r8d
	movl	$8, %edx
	leaq	_ZL6dev_mu(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	xorl	%ecx, %ecx
	leaq	56(%rsp), %rsi
	movl	$1, %r8d
	movl	$8, %edx
	leaq	_ZL5dev_P(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	xorl	%ecx, %ecx
	leaq	64(%rsp), %rsi
	movl	$1, %r8d
	movl	$8, %edx
	leaq	_ZL5dev_w(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movl	$1, %r8d
	xorl	%ecx, %ecx
	movl	$8, %edx
	leaq	72(%rsp), %rsi
	leaq	_ZL6dev_dx(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movl	44(%rsp), %eax
	leaq	80(%rsp), %rdi
	leal	-1(%rax), %esi
	imull	%eax, %esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	cudaMalloc@PLT
	movl	44(%rsp), %eax
	leaq	88(%rsp), %rdi
	leal	-1(%rax), %esi
	imull	%eax, %esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	cudaMalloc@PLT
	movl	44(%rsp), %eax
	leaq	96(%rsp), %rdi
	leal	-1(%rax), %esi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	cudaMalloc@PLT
	movl	%ebx, 120(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	movl	$4194305, %edx
	movl	$1, %ecx
	movl	$1, 140(%rsp)
	movq	%rax, 124(%rsp)
	movq	120(%rsp), %rdi
	salq	$10, %rdx
	movl	128(%rsp), %esi
	movq	%rdx, 132(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L29
	movl	44(%rsp), %eax
	movq	80(%rsp), %rdi
	leal	-1(%rax), %esi
	imull	%eax, %esi
	call	_Z24__device_stub__Z4initPdiPdi
.L29:
	movl	%ebx, 120(%rsp)
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	salq	$10, %rdx
	movl	$1, %ecx
	movl	$1, 140(%rsp)
	movq	%rax, 124(%rsp)
	movq	120(%rsp), %rdi
	movl	128(%rsp), %esi
	movq	%rdx, 132(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L30
	movl	44(%rsp), %eax
	movq	88(%rsp), %rdi
	leal	-1(%rax), %esi
	imull	%eax, %esi
	call	_Z24__device_stub__Z4initPdiPdi
.L30:
	movl	%ebx, 120(%rsp)
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rax
	salq	$10, %rdx
	movl	$1, %ecx
	movl	$1, 140(%rsp)
	movq	%rax, 124(%rsp)
	movq	120(%rsp), %rdi
	movl	128(%rsp), %esi
	movq	%rdx, 132(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L31
	movl	44(%rsp), %eax
	movq	96(%rsp), %rdi
	leal	-1(%rax), %esi
	imull	%esi, %esi
	call	_Z24__device_stub__Z4initPdiPdi
.L31:
	leaq	104(%rsp), %rdi
	movl	$8192, %esi
	movl	$4194305, %ebp
	xorl	%r12d, %r12d
	call	cudaMalloc@PLT
	movq	.LC1(%rip), %rax
	salq	$10, %rbp
	movq	%rax, 112(%rsp)
	call	clock@PLT
	movq	%rax, 16(%rsp)
.L32:
	cmpl	%r13d, %r12d
	jl	.L41
.L43:
	call	clock@PLT
	movq	16(%rsp), %rbx
	leaq	.LC4(%rip), %rsi
	leaq	.LC5(%rip), %rdi
	leaq	.LC9(%rip), %r14
	subq	%rbx, %rax
	cvtsi2ssq	%rax, %xmm0
	divss	.LC3(%rip), %xmm0
	movd	%xmm0, %ebx
	call	fopen@PLT
	cvtsi2ssl	%r12d, %xmm0
	movd	%ebx, %xmm1
	movl	44(%rsp), %ecx
	leaq	.LC6(%rip), %rdx
	movq	%rax, %rbp
	movl	$2, %esi
	movq	%rax, %rdi
	movb	$1, %al
	divss	%xmm0, %xmm1
	cvtss2sd	%xmm1, %xmm1
	movaps	%xmm1, %xmm0
	movsd	%xmm1, 8(%rsp)
	call	__fprintf_chk@PLT
	movq	%rbp, %rdi
	call	fclose@PLT
	movd	%ebx, %xmm4
	movl	$2, %edi
	movb	$2, %al
	movsd	8(%rsp), %xmm1
	leaq	.LC7(%rip), %rsi
	cvtss2sd	%xmm4, %xmm0
	call	__printf_chk@PLT
	movsd	112(%rsp), %xmm0
	movl	%r12d, %edx
	movb	$1, %al
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	44(%rsp), %r13d
	leal	-1(%r13), %ebx
	imull	%ebx, %r13d
	imull	%ebx, %ebx
	movslq	%r13d, %r13
	salq	$3, %r13
	movq	%r13, %rdi
	call	malloc@PLT
	movq	%r13, %rdi
	movq	%rax, %r12
	call	malloc@PLT
	movslq	%ebx, %rdi
	salq	$3, %rdi
	movq	%rax, %rbp
	call	malloc@PLT
	movq	80(%rsp), %rsi
	movq	%r13, %rdx
	movq	%r12, %rdi
	movl	$2, %ecx
	movq	%rax, %rbx
	call	cudaMemcpy@PLT
	movl	44(%rsp), %eax
	movq	88(%rsp), %rsi
	movq	%rbp, %rdi
	movl	$2, %ecx
	leal	-1(%rax), %edx
	imull	%eax, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	44(%rsp), %eax
	movq	96(%rsp), %rsi
	movq	%rbx, %rdi
	movl	$2, %ecx
	leal	-1(%rax), %edx
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movq	%r14, %rsi
	leaq	.LC10(%rip), %rdi
	call	fopen@PLT
	movl	$8, %esi
	movq	%r12, %rdi
	movq	%rax, %r13
	movl	44(%rsp), %eax
	movq	%r13, %rcx
	leal	-1(%rax), %edx
	imull	%eax, %edx
	movslq	%edx, %rdx
	call	fwrite@PLT
	movq	%r13, %rdi
	call	fclose@PLT
	movq	%r14, %rsi
	leaq	.LC11(%rip), %rdi
	call	fopen@PLT
	movl	$8, %esi
	movq	%rbp, %rdi
	movq	%rax, %r13
	movl	44(%rsp), %eax
	movq	%r13, %rcx
	leal	-1(%rax), %edx
	imull	%eax, %edx
	movslq	%edx, %rdx
	call	fwrite@PLT
	movq	%r13, %rdi
	call	fclose@PLT
	movq	%r14, %rsi
	leaq	.LC12(%rip), %rdi
	call	fopen@PLT
	movl	$8, %esi
	movq	%rbx, %rdi
	movq	%rax, %r13
	movl	44(%rsp), %eax
	movq	%r13, %rcx
	leal	-1(%rax), %edx
	imull	%edx, %edx
	movslq	%edx, %rdx
	call	fwrite@PLT
	movq	%r13, %rdi
	call	fclose@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	80(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	movq	1176(%rsp), %rax
	subq	%fs:40, %rax
	je	.L44
	call	__stack_chk_fail@PLT
.L41:
	movsd	112(%rsp), %xmm0
	comisd	8(%rsp), %xmm0
	jbe	.L43
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbp, %rdx
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movq	%rbp, 132(%rsp)
	movl	$1, 140(%rsp)
	movq	%rdi, 120(%rsp)
	movl	$1, 128(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	movq	104(%rsp), %rdi
	movl	$1024, %esi
	call	_Z24__device_stub__Z4initPdiPdi
.L33:
	movl	%ebx, 120(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbp, %rdx
	movabsq	$4294967297, %rax
	movl	$1, %ecx
	movq	%rbp, 132(%rsp)
	movq	%rax, 124(%rsp)
	movq	120(%rsp), %rdi
	movl	128(%rsp), %esi
	movl	$1, 140(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L34
	movq	104(%rsp), %rdx
	movq	96(%rsp), %rsi
	xorl	%ecx, %ecx
	movq	80(%rsp), %rdi
	call	_Z31__device_stub__Z7updateUPdS_S_iPdS_S_i
.L34:
	movl	%ebx, 120(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbp, %rdx
	movabsq	$4294967297, %rax
	movl	$1, %ecx
	movq	%rbp, 132(%rsp)
	movq	%rax, 124(%rsp)
	movq	120(%rsp), %rdi
	movl	128(%rsp), %esi
	movl	$1, 140(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L35
	movq	104(%rsp), %rdx
	movq	96(%rsp), %rsi
	movl	$1, %ecx
	movq	80(%rsp), %rdi
	call	_Z31__device_stub__Z7updateUPdS_S_iPdS_S_i
.L35:
	movl	%ebx, 120(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbp, %rdx
	movabsq	$4294967297, %rax
	movl	$1, %ecx
	movq	%rbp, 132(%rsp)
	movq	%rax, 124(%rsp)
	movq	120(%rsp), %rdi
	movl	128(%rsp), %esi
	movl	$1, 140(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L36
	movq	104(%rsp), %rdx
	movq	96(%rsp), %rsi
	xorl	%ecx, %ecx
	movq	88(%rsp), %rdi
	call	_Z31__device_stub__Z7updateVPdS_S_iPdS_S_i
.L36:
	movl	%ebx, 120(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbp, %rdx
	movabsq	$4294967297, %rax
	movl	$1, %ecx
	movq	%rbp, 132(%rsp)
	movq	%rax, 124(%rsp)
	movq	120(%rsp), %rdi
	movl	128(%rsp), %esi
	movl	$1, 140(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L37
	movq	104(%rsp), %rdx
	movq	96(%rsp), %rsi
	movl	$1, %ecx
	movq	88(%rsp), %rdi
	call	_Z31__device_stub__Z7updateVPdS_S_iPdS_S_i
.L37:
	movl	%ebx, 120(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbp, %rdx
	movabsq	$4294967297, %rax
	movl	$1, %ecx
	movq	%rbp, 132(%rsp)
	movq	%rax, 124(%rsp)
	movq	120(%rsp), %rdi
	movl	128(%rsp), %esi
	movl	$1, 140(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L38
	movq	104(%rsp), %rcx
	movq	96(%rsp), %rdx
	movq	88(%rsp), %rsi
	movq	80(%rsp), %rdi
	call	_Z32__device_stub__Z7updatePPdS_S_S_PdS_S_S_
.L38:
	movl	$10, %r14d
	movl	$1, %r15d
.L40:
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbp, %rdx
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movq	%rbp, 132(%rsp)
	movl	$1, 140(%rsp)
	movq	%rdi, 120(%rsp)
	movl	$1, 128(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L39
	movq	104(%rsp), %rdi
	movl	%r15d, %esi
	call	_Z31__device_stub__Z10findmaxresPdiPdi
.L39:
	addl	%r15d, %r15d
	decl	%r14d
	jne	.L40
	movq	104(%rsp), %rsi
	leaq	112(%rsp), %rdi
	movl	$2, %ecx
	incl	%r12d
	movl	$8, %edx
	call	cudaMemcpy@PLT
	jmp	.L32
.L44:
	addq	$1192, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC13:
	.string	"_Z10findmaxresPdi"
.LC14:
	.string	"_Z7updatePPdS_S_S_"
.LC15:
	.string	"_Z7updateVPdS_S_i"
.LC16:
	.string	"_Z7updateUPdS_S_i"
.LC17:
	.string	"_Z4initPdi"
.LC18:
	.string	"dev_N"
.LC19:
	.string	"dev_mu"
.LC20:
	.string	"dev_P"
.LC21:
	.string	"dev_w"
.LC22:
	.string	"dev_dx"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2063:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC13(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z10findmaxresPdi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC14(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z7updatePPdS_S_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC15(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z7updateVPdS_S_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC16(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z7updateUPdS_S_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC17(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z4initPdi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC18(%rip), %rdx
	movl	$4, %r9d
	leaq	_ZL5dev_N(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC19(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL6dev_mu(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC20(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL5dev_P(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC21(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL5dev_w(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC22(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$8, %r9d
	movq	%rdx, %rcx
	leaq	_ZL6dev_dx(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2063:
	.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
	.local	_ZL6dev_dx
	.comm	_ZL6dev_dx,8,8
	.local	_ZL5dev_w
	.comm	_ZL5dev_w,8,8
	.local	_ZL5dev_P
	.comm	_ZL5dev_P,8,8
	.local	_ZL6dev_mu
	.comm	_ZL6dev_mu,8,8
	.local	_ZL5dev_N
	.comm	_ZL5dev_N,4,4
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC1:
	.long	0
	.long	1072693248
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC3:
	.long	1232348160
	.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:

/*
Author: Jason He
Version: 1.0 20210605 Serial version.
Version: 2.0 20210607 CUDA version.
*/


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

/*
This program solve the Stokes Flow problem- 36.1.2 in the textbook.

Inputs: 
	N, number of grid points in each direction.
	mu, double precision viscosity.
	P, double precision pressure drop.
	w, double precision relaxation parameter.
    tol, double precision error tolerance.
	K, maximum number of iterations.

Outputs:
	print arguments
	output final grid value into "StokesU.out", "StokesV.out" and "StokesP.out"
	print run time and output it to runtime.dat

E.g.
$ ./stokes 128 1 1 0.4 1e-9 100000

*/

//declare global variables
const int threadsPerBlock = 1024;//for NVIDIA TESLA K20c maximum is 1024
/*the findmaxres kernel demand that threadsPerBlock should be a power of 2.
Also in the three update kernels, each block compare and store the residuals into
the global array dev_res[threadsPerBlock], which only work with maximum
threadsPerBlock=1024, since in this case only one block runs at a time.*/
__device__ __constant__ int dev_N;
__device__ __constant__ double dev_mu;
__device__ __constant__ double dev_P;
__device__ __constant__ double dev_w;
__device__ __constant__ double dev_dx;

//this kernel set values of 1d length-n double array to 0
__global__ void init(double* a, int n){
    int k = threadIdx.x + blockIdx.x*blockDim.x;
    if(k < n){
        a[k] = 0.;
    }
}

//the three update functions update the grid values and return the maximum residual
__global__ void updateU(double* dev_u, double* dev_p, double* dev_res, int color){
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < dev_N*(dev_N-1)){
        int j = i/(dev_N-1);
        int k = i%(dev_N-1);
        if((j+k)%2==color){
            double r;
            int jp = i + dev_N-1;//(j+1)*(dev_N-1) + k;
            int jm = i - dev_N+1;//(j-1)*(dev_N-1) + k;
            int kp = i + 1;//j*(dev_N-1) + k+1;
            int km = i - 1;//j*(dev_N-1) + k-1;
            int pi = j*(dev_N-1) + k;
            int pm = pi - dev_N+1;//(j-1)*(dev_N-1) + k;
            if(j==0){
                if(k==0){
                    r = dev_mu*(dev_u[jp]+dev_u[kp]-4*dev_u[i])
                        -dev_dx*2*(dev_p[pi]-dev_P);
                }
                else if(k<dev_N-2){
                    r = dev_mu*(dev_u[jp]+dev_u[kp]+dev_u[km]-3*dev_u[i])
                        -dev_dx*2*(dev_p[pi]-dev_P);
                }
                else{
                    r = dev_mu*(dev_u[jp]+dev_u[km]-4*dev_u[i])
                        -dev_dx*2*(dev_p[pi]-dev_P);
                }
            }
            else if(j<dev_N-1){
                if(k==0){
                    r = dev_mu*(dev_u[jp]+dev_u[jm]+dev_u[kp]-5*dev_u[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
                else if(k<dev_N-2){
                    r = dev_mu*(dev_u[jp]+dev_u[jm]+dev_u[kp]+dev_u[km]-4*dev_u[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
                else{
                    r = dev_mu*(dev_u[jp]+dev_u[jm]+dev_u[km]-5*dev_u[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
            }
            else{
                if(k==0){
                    r = dev_mu*(dev_u[jm]+dev_u[kp]-4*dev_u[i])
                        -dev_dx*2*(0-dev_p[pm]);
                }
                else if(k<dev_N-2){
                    r = dev_mu*(dev_u[jm]+dev_u[kp]+dev_u[km]-3*dev_u[i])
                        -dev_dx*2*(0-dev_p[pm]);
                }
                else{
                    r = dev_mu*(dev_u[jm]+dev_u[km]-4*dev_u[i])
                        -dev_dx*2*(0-dev_p[pm]);
                }
            }
            dev_u[i] += dev_w*r;
            dev_res[threadIdx.x] = max(fabs(r), dev_res[threadIdx.x]);
        }
    }
}

__global__ void updateV(double* dev_v, double* dev_p, double* dev_res, int color){
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < (dev_N-1)*dev_N){
        int k = i%dev_N;
        if(0<k && k<dev_N-1){
            int j = i/dev_N;
            if((j+k)%2==color){
                double r;
                int jp = i + dev_N;//(j+1)*dev_N + k;
                int jm = i - dev_N;//(j-1)*dev_N + k;
                int kp = i + 1;//j*dev_N + k+1;
                int km = i - 1;//j*dev_N + k-1;
                int pi = j*(dev_N-1) + k;
                int pm = pi - 1;//j*(dev_N-1) + k-1;
                if(j==0){
                    r = dev_mu*(dev_v[kp]+dev_v[km]+dev_v[jp]-3*dev_v[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
                else if(j<dev_N-2){
                    r = dev_mu*(dev_v[kp]+dev_v[km]+dev_v[jp]+dev_v[jm]-4*dev_v[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
                else{
                    r = dev_mu*(dev_v[kp]+dev_v[km]+dev_v[jm]-3*dev_v[i])
                        -dev_dx*(dev_p[pi]-dev_p[pm]);
                }
                dev_v[i] += dev_w*r;
                dev_res[threadIdx.x] = max(fabs(r), dev_res[threadIdx.x]);
            }
        }
    }
}

__global__ void updateP(double* dev_u, double* dev_v, double* dev_p, double* dev_res){
    
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    if(i < (dev_N-1)*(dev_N-1)){
        int j = i/(dev_N-1);
        int k = i%(dev_N-1);
        int up = (j+1)*(dev_N-1) + k;
        int um = j*(dev_N-1) + k;
        int vp = j*dev_N + k+1;
        int vm = j*dev_N + k;
        double r;
        r = -(dev_u[up]-dev_u[um])-(dev_v[vp]-dev_v[vm]);
        dev_p[i] += dev_w*r;//update all p's
        dev_res[threadIdx.x] = max(fabs(r), dev_res[threadIdx.x]);//not comfortable with this
    }
}

//this kernel find the max value
__global__ void findmaxres(double* dev_res, int i){
    int k = threadIdx.x + blockIdx.x*blockDim.x;
    if(k%(2*i)==0){
        dev_res[k] = max(dev_res[k], dev_res[k+i]);
    }
}

int main(int argc, char* argv[]){
    //Choose gpu device
    hipDeviceProp_t prop;
    int dev;
    memset(&prop, 0, sizeof(hipDeviceProp_t));
    prop.multiProcessorCount = 13;
    hipChooseDevice(&dev, &prop);
    hipSetDevice(dev);

    //load input parameters
    int argi = 0;
    int N = atoi(argv[++argi]);
    double mu = atof(argv[++argi]);
    double P = atof(argv[++argi]);
    double w = atof(argv[++argi]);
    double tol = atof(argv[++argi]);
    int K = atoi(argv[++argi]);
    printf( "N = %d\nmu = %lf\nP = %lf\nw = %lf\ntau = %e\nK = %d\n",
            N, mu, P, w, tol, K);
    
    //calculate parameters for the iteration equation
    double dx = 1./(N-1);
    //calculate parameters for blocks and threads
    const int blocksPerGrid = (N*(N-1) + threadsPerBlock - 1)/threadsPerBlock;
    printf("BlocksPerGrid = %d\nThreadsPerBlock = %d\n", blocksPerGrid, threadsPerBlock);

    //copy parameters to device
    hipMemcpyToSymbol(HIP_SYMBOL(dev_N), &N, sizeof(double));
    hipMemcpyToSymbol(HIP_SYMBOL(dev_mu), &mu, sizeof(double));
    hipMemcpyToSymbol(HIP_SYMBOL(dev_P), &P, sizeof(double));
    hipMemcpyToSymbol(HIP_SYMBOL(dev_w), &w, sizeof(double));
    hipMemcpyToSymbol(HIP_SYMBOL(dev_dx), &dx, sizeof(double));

    //initialize value of the grid on device
    double *dev_u, *dev_v, *dev_p;
    hipMalloc((void**)&dev_u, N*(N-1)*sizeof(double));
    hipMalloc((void**)&dev_v, (N-1)*N*sizeof(double));
    hipMalloc((void**)&dev_p, (N-1)*(N-1)*sizeof(double));
    init<<<blocksPerGrid, threadsPerBlock>>>(dev_u, N*(N-1));
    init<<<blocksPerGrid, threadsPerBlock>>>(dev_v, (N-1)*N);
    init<<<blocksPerGrid, threadsPerBlock>>>(dev_p, (N-1)*(N-1));
    
    //initialize grids of residual
    double *dev_res;
    hipMalloc((void**)&dev_res, threadsPerBlock*sizeof(double));

    //initialize parameters for iterations
    int iter=0;//number of iteration
    double maxres=1.0;//maximum residual of u,v,p on every grid points
    float runtime;//record runtime
    clock_t t;
    t = clock();
    //main loop
    while(iter<K && maxres>tol){
        init<<<1, threadsPerBlock>>>(dev_res, threadsPerBlock);
        updateU<<<blocksPerGrid, threadsPerBlock>>>(dev_u, dev_p, dev_res, 0);
        updateU<<<blocksPerGrid, threadsPerBlock>>>(dev_u, dev_p, dev_res, 1);
        updateV<<<blocksPerGrid, threadsPerBlock>>>(dev_v, dev_p, dev_res, 0);
        updateV<<<blocksPerGrid, threadsPerBlock>>>(dev_v, dev_p, dev_res, 1);
        updateP<<<blocksPerGrid, threadsPerBlock>>>(dev_u, dev_v, dev_p, dev_res);
        for(int i=1; i<threadsPerBlock; i*=2)
            findmaxres<<<1, threadsPerBlock>>>(dev_res, i);
        hipMemcpy(&maxres, dev_res, sizeof(double), hipMemcpyDeviceToHost);
        iter++;
    }
    runtime = (float)(clock()-t)/CLOCKS_PER_SEC;
    FILE* tfile = fopen("runtime.dat", "a");
    fprintf(tfile, "%d %f\n", N, runtime/iter);
    fclose(tfile);
    printf("Runtime = %f seconds\nOne iteration runtime = %f seconds\n",
            runtime, runtime/iter);

    //print iterations and residuals
    printf("number of iterations = %d\nresidual = %e\n", iter, maxres);
    
    //initialize the grids on the host
    double *u = (double*)malloc(N*(N-1)*sizeof(double));
    double *v = (double*)malloc((N-1)*N*sizeof(double));
    double *p = (double*)malloc((N-1)*(N-1)*sizeof(double));

    //output to files
    hipMemcpy(u, dev_u, N*(N-1)*sizeof(double), hipMemcpyDeviceToHost);
    hipMemcpy(v, dev_v, (N-1)*N*sizeof(double), hipMemcpyDeviceToHost);
    hipMemcpy(p, dev_p, (N-1)*(N-1)*sizeof(double), hipMemcpyDeviceToHost);
    FILE* ufile = fopen("stokesU.out","w");
    fwrite(u, sizeof(double), N*(N-1), ufile);
    fclose(ufile);
    FILE* vfile = fopen("stokesV.out","w");
    fwrite(v, sizeof(double), (N-1)*N, vfile);
    fclose(vfile);
    FILE* pfile = fopen("stokesP.out","w");
    fwrite(p, sizeof(double), (N-1)*(N-1), pfile);
    fclose(pfile);

    //free memories
    free(u);
    free(v);
    free(p);
    hipFree(dev_u);
    hipFree(dev_v);
    hipFree(dev_p);
    hipFree(dev_res);

    return 0;
}
	.text
	.file	"stokes.hip"
	.globl	_Z19__device_stub__initPdi      # -- Begin function _Z19__device_stub__initPdi
	.type	_Z19__device_stub__initPdi,@function
_Z19__device_stub__initPdi:             # @_Z19__device_stub__initPdi
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	4(%rsp), %rcx
	movl	%esi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%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	$_Z4initPdi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z19__device_stub__initPdi, .Lfunc_end0-_Z19__device_stub__initPdi
	.cfi_endproc
                                        # -- End function
	.globl	_Z22__device_stub__updateUPdS_S_i # -- Begin function _Z22__device_stub__updateUPdS_S_i
	.type	_Z22__device_stub__updateUPdS_S_i,@function
_Z22__device_stub__updateUPdS_S_i:      # @_Z22__device_stub__updateUPdS_S_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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%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	$_Z7updateUPdS_S_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z22__device_stub__updateUPdS_S_i, .Lfunc_end1-_Z22__device_stub__updateUPdS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z22__device_stub__updateVPdS_S_i # -- Begin function _Z22__device_stub__updateVPdS_S_i
	.type	_Z22__device_stub__updateVPdS_S_i,@function
_Z22__device_stub__updateVPdS_S_i:      # @_Z22__device_stub__updateVPdS_S_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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%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	$_Z7updateVPdS_S_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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	_Z22__device_stub__updateVPdS_S_i, .Lfunc_end2-_Z22__device_stub__updateVPdS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z22__device_stub__updatePPdS_S_S_ # -- Begin function _Z22__device_stub__updatePPdS_S_S_
	.type	_Z22__device_stub__updatePPdS_S_S_,@function
_Z22__device_stub__updatePPdS_S_S_:     # @_Z22__device_stub__updatePPdS_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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	16(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rsp, %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z7updatePPdS_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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_end3:
	.size	_Z22__device_stub__updatePPdS_S_S_, .Lfunc_end3-_Z22__device_stub__updatePPdS_S_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z25__device_stub__findmaxresPdi # -- Begin function _Z25__device_stub__findmaxresPdi
	.type	_Z25__device_stub__findmaxresPdi,@function
_Z25__device_stub__findmaxresPdi:       # @_Z25__device_stub__findmaxresPdi
	.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	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	4(%rsp), %rcx
	movl	%esi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%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	$_Z10findmaxresPdi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	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_end4:
	.size	_Z25__device_stub__findmaxresPdi, .Lfunc_end4-_Z25__device_stub__findmaxresPdi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI5_0:
	.quad	0x3ff0000000000000              # double 1
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI5_1:
	.long	0x49742400                      # float 1.0E+6
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$1592, %rsp                     # imm = 0x638
	.cfi_def_cfa_offset 1648
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %r15
	movabsq	$4294968320, %rbx               # imm = 0x100000400
	leaq	120(%rsp), %r14
	movl	$1472, %edx                     # imm = 0x5C0
	movq	%r14, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movl	$13, 388(%r14)
	leaq	108(%rsp), %r12
	movq	%r12, %rdi
	movq	%r14, %rsi
	callq	hipChooseDeviceR0600
	movl	(%r12), %edi
	callq	hipSetDevice
	movq	8(%r15), %rdi
	callq	atoi
	movl	%eax, %r12d
	leaq	4(%rsp), %r14
	movl	%eax, (%r14)
	movq	16(%r15), %rdi
	callq	atof
	movsd	%xmm0, 48(%rsp)                 # 8-byte Spill
	leaq	96(%rsp), %rax
	movsd	%xmm0, (%rax)
	movq	24(%r15), %rdi
	callq	atof
	movsd	%xmm0, 72(%rsp)                 # 8-byte Spill
	leaq	88(%rsp), %rax
	movsd	%xmm0, (%rax)
	movq	32(%r15), %rdi
	callq	atof
	movsd	%xmm0, 64(%rsp)                 # 8-byte Spill
	leaq	80(%rsp), %rax
	movsd	%xmm0, (%rax)
	movq	40(%r15), %rdi
	callq	atof
	movsd	%xmm0, 24(%rsp)                 # 8-byte Spill
	movq	48(%r15), %rdi
	callq	atoi
	movl	%eax, %ebp
	movl	$.L.str, %edi
	movl	%r12d, %esi
	movsd	48(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movsd	72(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	movsd	64(%rsp), %xmm2                 # 8-byte Reload
                                        # xmm2 = mem[0],zero
	movsd	24(%rsp), %xmm3                 # 8-byte Reload
                                        # xmm3 = mem[0],zero
	movl	%eax, %edx
	movb	$4, %al
	callq	printf
	leal	-1(%r12), %r15d
	xorps	%xmm0, %xmm0
	cvtsi2sd	%r15d, %xmm0
	movsd	.LCPI5_0(%rip), %xmm1           # xmm1 = mem[0],zero
	divsd	%xmm0, %xmm1
	leaq	112(%rsp), %r13
	movsd	%xmm1, (%r13)
	imull	%r12d, %r15d
	leal	1023(%r15), %eax
	addl	$2046, %r15d                    # imm = 0x7FE
	testl	%eax, %eax
	cmovnsl	%eax, %r15d
	sarl	$10, %r15d
	movl	$.L.str.1, %edi
	movl	%r15d, %esi
	movl	$1024, %edx                     # imm = 0x400
	xorl	%eax, %eax
	callq	printf
	movl	$dev_N, %edi
	movl	$8, %edx
	movq	%r14, %rsi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movl	$dev_mu, %edi
	movl	$8, %edx
	leaq	96(%rsp), %rsi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movl	$dev_P, %edi
	movl	$8, %edx
	leaq	88(%rsp), %rsi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movl	$dev_w, %edi
	movl	$8, %edx
	leaq	80(%rsp), %rsi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movl	$dev_dx, %edi
	movl	$8, %edx
	movq	%r13, %rsi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movslq	(%r14), %rax
	leaq	-1(%rax), %rsi
	imulq	%rax, %rsi
	shlq	$3, %rsi
	leaq	40(%rsp), %rdi
	callq	hipMalloc
	movslq	(%r14), %rax
	leaq	-1(%rax), %rsi
	imulq	%rax, %rsi
	shlq	$3, %rsi
	leaq	32(%rsp), %rdi
	callq	hipMalloc
	movl	(%r14), %esi
	decl	%esi
	imull	%esi, %esi
	shlq	$3, %rsi
	leaq	16(%rsp), %rdi
	callq	hipMalloc
	addq	%rbx, %r15
	addq	$-1024, %r15                    # imm = 0xFC00
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_2
# %bb.1:
	movq	40(%rsp), %rdi
	movl	4(%rsp), %eax
	leal	-1(%rax), %esi
	imull	%eax, %esi
	callq	_Z19__device_stub__initPdi
.LBB5_2:
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_4
# %bb.3:
	movq	32(%rsp), %rdi
	movl	4(%rsp), %eax
	leal	-1(%rax), %esi
	imull	%eax, %esi
	callq	_Z19__device_stub__initPdi
.LBB5_4:
	xorl	%r14d, %r14d
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_6
# %bb.5:
	movq	16(%rsp), %rdi
	movl	4(%rsp), %esi
	decl	%esi
	imull	%esi, %esi
	callq	_Z19__device_stub__initPdi
.LBB5_6:
	leaq	8(%rsp), %rdi
	movl	$8192, %esi                     # imm = 0x2000
	callq	hipMalloc
	movabsq	$4607182418800017408, %rax      # imm = 0x3FF0000000000000
	movq	%rax, 56(%rsp)
	callq	clock
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	testl	%ebp, %ebp
	jle	.LBB5_26
# %bb.7:
	movsd	.LCPI5_0(%rip), %xmm0           # xmm0 = mem[0],zero
	ucomisd	24(%rsp), %xmm0                 # 8-byte Folded Reload
	jbe	.LBB5_26
# %bb.8:                                # %.lr.ph.preheader
	xorl	%r14d, %r14d
	leaq	-1023(%rbx), %r13
.LBB5_9:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_21 Depth 2
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_11
# %bb.10:                               #   in Loop: Header=BB5_9 Depth=1
	movq	8(%rsp), %rdi
	movl	$1024, %esi                     # imm = 0x400
	callq	_Z19__device_stub__initPdi
.LBB5_11:                               #   in Loop: Header=BB5_9 Depth=1
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_13
# %bb.12:                               #   in Loop: Header=BB5_9 Depth=1
	movq	40(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
	xorl	%ecx, %ecx
	callq	_Z22__device_stub__updateUPdS_S_i
.LBB5_13:                               #   in Loop: Header=BB5_9 Depth=1
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_15
# %bb.14:                               #   in Loop: Header=BB5_9 Depth=1
	movq	40(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
	movl	$1, %ecx
	callq	_Z22__device_stub__updateUPdS_S_i
.LBB5_15:                               #   in Loop: Header=BB5_9 Depth=1
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_17
# %bb.16:                               #   in Loop: Header=BB5_9 Depth=1
	movq	32(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
	xorl	%ecx, %ecx
	callq	_Z22__device_stub__updateVPdS_S_i
.LBB5_17:                               #   in Loop: Header=BB5_9 Depth=1
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_19
# %bb.18:                               #   in Loop: Header=BB5_9 Depth=1
	movq	32(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
	movl	$1, %ecx
	callq	_Z22__device_stub__updateVPdS_S_i
.LBB5_19:                               #   in Loop: Header=BB5_9 Depth=1
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_20
# %bb.27:                               #   in Loop: Header=BB5_9 Depth=1
	movq	40(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rcx
	callq	_Z22__device_stub__updatePPdS_S_S_
.LBB5_20:                               # %.preheader
                                        #   in Loop: Header=BB5_9 Depth=1
	movl	$1, %r12d
.LBB5_21:                               #   Parent Loop BB5_9 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_23
# %bb.22:                               #   in Loop: Header=BB5_21 Depth=2
	movq	8(%rsp), %rdi
	movl	%r12d, %esi
	callq	_Z25__device_stub__findmaxresPdi
.LBB5_23:                               #   in Loop: Header=BB5_21 Depth=2
	leal	(%r12,%r12), %eax
	cmpl	$512, %r12d                     # imm = 0x200
	movl	%eax, %r12d
	jb	.LBB5_21
# %bb.24:                               #   in Loop: Header=BB5_9 Depth=1
	movq	8(%rsp), %rsi
	movl	$8, %edx
	leaq	56(%rsp), %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	incl	%r14d
	cmpl	%ebp, %r14d
	jge	.LBB5_26
# %bb.25:                               #   in Loop: Header=BB5_9 Depth=1
	movsd	56(%rsp), %xmm0                 # xmm0 = mem[0],zero
	ucomisd	24(%rsp), %xmm0                 # 8-byte Folded Reload
	ja	.LBB5_9
.LBB5_26:                               # %._crit_edge
	callq	clock
	subq	48(%rsp), %rax                  # 8-byte Folded Reload
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rax, %xmm0
	divss	.LCPI5_1(%rip), %xmm0
	movss	%xmm0, 24(%rsp)                 # 4-byte Spill
	movl	$.L.str.2, %edi
	movl	$.L.str.3, %esi
	callq	fopen
	movq	%rax, %rbx
	xorps	%xmm0, %xmm0
	cvtsi2ss	%r14d, %xmm0
	movl	4(%rsp), %edx
	movss	24(%rsp), %xmm1                 # 4-byte Reload
                                        # xmm1 = mem[0],zero,zero,zero
	divss	%xmm0, %xmm1
	xorps	%xmm0, %xmm0
	cvtss2sd	%xmm1, %xmm0
	movsd	%xmm0, 48(%rsp)                 # 8-byte Spill
	movl	$.L.str.4, %esi
	movq	%rax, %rdi
	movb	$1, %al
	callq	fprintf
	movq	%rbx, %rdi
	callq	fclose
	xorps	%xmm0, %xmm0
	cvtss2sd	24(%rsp), %xmm0                 # 4-byte Folded Reload
	movl	$.L.str.5, %edi
	movsd	48(%rsp), %xmm1                 # 8-byte Reload
                                        # xmm1 = mem[0],zero
	movb	$2, %al
	callq	printf
	movsd	56(%rsp), %xmm0                 # xmm0 = mem[0],zero
	movl	$.L.str.6, %edi
	movl	%r14d, %esi
	movb	$1, %al
	callq	printf
	movslq	4(%rsp), %rax
	leaq	-1(%rax), %r15
	movq	%r15, %r12
	imulq	%rax, %r12
	shlq	$3, %r12
	movq	%r12, %rdi
	callq	malloc
	movq	%rax, %rbx
	movq	%r12, %rdi
	callq	malloc
	movq	%rax, %r14
	imull	%r15d, %r15d
	shlq	$3, %r15
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %r15
	movq	40(%rsp), %rsi
	movq	%rbx, %rdi
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	32(%rsp), %rsi
	movslq	4(%rsp), %rax
	leaq	-1(%rax), %rdx
	imulq	%rax, %rdx
	shlq	$3, %rdx
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rsi
	movl	4(%rsp), %edx
	decl	%edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.L.str.7, %edi
	movl	$.L.str.8, %esi
	callq	fopen
	movq	%rax, %r12
	movslq	4(%rsp), %rax
	leaq	-1(%rax), %rdx
	imulq	%rax, %rdx
	movl	$8, %esi
	movq	%rbx, %rdi
	movq	%r12, %rcx
	callq	fwrite
	movq	%r12, %rdi
	callq	fclose
	movl	$.L.str.9, %edi
	movl	$.L.str.8, %esi
	callq	fopen
	movq	%rax, %r12
	movslq	4(%rsp), %rax
	leaq	-1(%rax), %rdx
	imulq	%rax, %rdx
	movl	$8, %esi
	movq	%r14, %rdi
	movq	%r12, %rcx
	callq	fwrite
	movq	%r12, %rdi
	callq	fclose
	movl	$.L.str.10, %edi
	movl	$.L.str.8, %esi
	callq	fopen
	movq	%rax, %r12
	movl	4(%rsp), %edx
	decl	%edx
	imull	%edx, %edx
	movl	$8, %esi
	movq	%r15, %rdi
	movq	%rax, %rcx
	callq	fwrite
	movq	%r12, %rdi
	callq	fclose
	movq	%rbx, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$1592, %rsp                     # imm = 0x638
	.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_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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z4initPdi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z7updateUPdS_S_i, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z7updateVPdS_S_i, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z7updatePPdS_S_S_, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z10findmaxresPdi, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$dev_N, %esi
	movl	$.L__unnamed_6, %edx
	movl	$.L__unnamed_6, %ecx
	movl	$4, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$dev_mu, %esi
	movl	$.L__unnamed_7, %edx
	movl	$.L__unnamed_7, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$dev_P, %esi
	movl	$.L__unnamed_8, %edx
	movl	$.L__unnamed_8, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$dev_w, %esi
	movl	$.L__unnamed_9, %edx
	movl	$.L__unnamed_9, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$0, 8(%rsp)
	movl	$1, (%rsp)
	movl	$dev_dx, %esi
	movl	$.L__unnamed_10, %edx
	movl	$.L__unnamed_10, %ecx
	movl	$8, %r9d
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	callq	__hipRegisterVar
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.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	dev_N,@object                   # @dev_N
	.local	dev_N
	.comm	dev_N,4,4
	.type	dev_mu,@object                  # @dev_mu
	.local	dev_mu
	.comm	dev_mu,8,8
	.type	dev_P,@object                   # @dev_P
	.local	dev_P
	.comm	dev_P,8,8
	.type	dev_w,@object                   # @dev_w
	.local	dev_w
	.comm	dev_w,8,8
	.type	dev_dx,@object                  # @dev_dx
	.local	dev_dx
	.comm	dev_dx,8,8
	.type	_Z4initPdi,@object              # @_Z4initPdi
	.section	.rodata,"a",@progbits
	.globl	_Z4initPdi
	.p2align	3, 0x0
_Z4initPdi:
	.quad	_Z19__device_stub__initPdi
	.size	_Z4initPdi, 8

	.type	_Z7updateUPdS_S_i,@object       # @_Z7updateUPdS_S_i
	.globl	_Z7updateUPdS_S_i
	.p2align	3, 0x0
_Z7updateUPdS_S_i:
	.quad	_Z22__device_stub__updateUPdS_S_i
	.size	_Z7updateUPdS_S_i, 8

	.type	_Z7updateVPdS_S_i,@object       # @_Z7updateVPdS_S_i
	.globl	_Z7updateVPdS_S_i
	.p2align	3, 0x0
_Z7updateVPdS_S_i:
	.quad	_Z22__device_stub__updateVPdS_S_i
	.size	_Z7updateVPdS_S_i, 8

	.type	_Z7updatePPdS_S_S_,@object      # @_Z7updatePPdS_S_S_
	.globl	_Z7updatePPdS_S_S_
	.p2align	3, 0x0
_Z7updatePPdS_S_S_:
	.quad	_Z22__device_stub__updatePPdS_S_S_
	.size	_Z7updatePPdS_S_S_, 8

	.type	_Z10findmaxresPdi,@object       # @_Z10findmaxresPdi
	.globl	_Z10findmaxresPdi
	.p2align	3, 0x0
_Z10findmaxresPdi:
	.quad	_Z25__device_stub__findmaxresPdi
	.size	_Z10findmaxresPdi, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"N = %d\nmu = %lf\nP = %lf\nw = %lf\ntau = %e\nK = %d\n"
	.size	.L.str, 49

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"BlocksPerGrid = %d\nThreadsPerBlock = %d\n"
	.size	.L.str.1, 41

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"runtime.dat"
	.size	.L.str.2, 12

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

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"%d %f\n"
	.size	.L.str.4, 7

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Runtime = %f seconds\nOne iteration runtime = %f seconds\n"
	.size	.L.str.5, 57

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"number of iterations = %d\nresidual = %e\n"
	.size	.L.str.6, 41

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"stokesU.out"
	.size	.L.str.7, 12

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

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"stokesV.out"
	.size	.L.str.9, 12

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"stokesP.out"
	.size	.L.str.10, 12

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z4initPdi"
	.size	.L__unnamed_1, 11

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z7updateUPdS_S_i"
	.size	.L__unnamed_2, 18

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z7updateVPdS_S_i"
	.size	.L__unnamed_3, 18

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z7updatePPdS_S_S_"
	.size	.L__unnamed_4, 19

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"_Z10findmaxresPdi"
	.size	.L__unnamed_5, 18

	.type	.L__unnamed_6,@object           # @5
.L__unnamed_6:
	.asciz	"dev_N"
	.size	.L__unnamed_6, 6

	.type	.L__unnamed_7,@object           # @6
.L__unnamed_7:
	.asciz	"dev_mu"
	.size	.L__unnamed_7, 7

	.type	.L__unnamed_8,@object           # @7
.L__unnamed_8:
	.asciz	"dev_P"
	.size	.L__unnamed_8, 6

	.type	.L__unnamed_9,@object           # @8
.L__unnamed_9:
	.asciz	"dev_w"
	.size	.L__unnamed_9, 6

	.type	.L__unnamed_10,@object          # @9
.L__unnamed_10:
	.asciz	"dev_dx"
	.size	.L__unnamed_10, 7

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

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

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z19__device_stub__initPdi
	.addrsig_sym _Z22__device_stub__updateUPdS_S_i
	.addrsig_sym _Z22__device_stub__updateVPdS_S_i
	.addrsig_sym _Z22__device_stub__updatePPdS_S_S_
	.addrsig_sym _Z25__device_stub__findmaxresPdi
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym dev_N
	.addrsig_sym dev_mu
	.addrsig_sym dev_P
	.addrsig_sym dev_w
	.addrsig_sym dev_dx
	.addrsig_sym _Z4initPdi
	.addrsig_sym _Z7updateUPdS_S_i
	.addrsig_sym _Z7updateVPdS_S_i
	.addrsig_sym _Z7updatePPdS_S_S_
	.addrsig_sym _Z10findmaxresPdi
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z4initPdi              ; -- Begin function _Z4initPdi
	.globl	_Z4initPdi
	.p2align	8
	.type	_Z4initPdi,@function
_Z4initPdi:                             ; @_Z4initPdi
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x1c
	s_load_b32 s3, s[0:1], 0x8
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s3, v1
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_load_b64 s[0:1], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[1:2]
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, v2
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b64 v[0:1], v[2:3], off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z4initPdi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.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 4
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z4initPdi, .Lfunc_end0-_Z4initPdi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 132
; NumSgprs: 18
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 4
; 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
	.protected	_Z7updateUPdS_S_i       ; -- Begin function _Z7updateUPdS_S_i
	.globl	_Z7updateUPdS_S_i
	.p2align	8
	.type	_Z7updateUPdS_S_i,@function
_Z7updateUPdS_S_i:                      ; @_Z7updateUPdS_S_i
; %bb.0:
	s_load_b32 s4, s[0:1], 0x2c
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, dev_N@rel32@lo+4
	s_addc_u32 s3, s3, dev_N@rel32@hi+12
	s_load_b32 s10, s[2:3], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[5:6], null, s15, s2, v[0:1]
	s_add_i32 s8, s10, -1
	s_mul_i32 s2, s8, s10
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_i32_e32 vcc_lo, s2, v5
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB1_35
; %bb.1:
	s_ashr_i32 s2, s8, 31
	v_ashrrev_i32_e32 v6, 31, v5
	s_add_i32 s3, s8, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_xor_b32 s3, s3, s2
	v_add_nc_u32_e32 v3, v5, v6
	v_cvt_f32_u32_e32 v1, s3
	s_sub_i32 s4, 0, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v3, v3, v6
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v1, v1
	v_mul_lo_u32 v2, s4, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v1, v2
	v_add_nc_u32_e32 v1, v1, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v1, v3, v1
	v_mul_lo_u32 v2, v1, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v2, v3, v2
	v_add_nc_u32_e32 v3, 1, v1
	v_subrev_nc_u32_e32 v4, s3, v2
	v_cmp_le_u32_e32 vcc_lo, s3, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v2, v2, v4 :: v_dual_cndmask_b32 v1, v1, v3
	v_xor_b32_e32 v4, s2, v6
	s_load_b32 s2, s[0:1], 0x18
	v_cmp_le_u32_e32 vcc_lo, s3, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, 1, v1
	v_cndmask_b32_e32 v1, v1, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v1, v1, v4
	v_sub_nc_u32_e32 v12, v1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v1, v12, s8
	v_sub_nc_u32_e32 v10, v5, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v1, v10, v12
	v_lshrrev_b32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v2, v1, v2
	v_and_b32_e32 v2, -2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v1, v1, v2
	s_waitcnt lgkmcnt(0)
	v_cmp_eq_u32_e32 vcc_lo, s2, v1
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB1_35
; %bb.2:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[2:3], s[0:1], 0x10
	v_mad_u64_u32 v[3:4], null, v12, s8, v[10:11]
	v_add_nc_u32_e32 v13, s8, v5
                                        ; implicit-def: $vgpr8_vgpr9
                                        ; implicit-def: $vgpr1_vgpr2
	s_mov_b32 s0, exec_lo
	v_cmpx_ne_u32_e32 0, v12
	s_xor_b32 s11, exec_lo, s0
	s_cbranch_execz .LBB1_24
; %bb.3:
	v_subrev_nc_u32_e32 v11, s10, v5
	s_delay_alu instid0(VALU_DEP_4)
	v_subrev_nc_u32_e32 v7, s10, v3
	v_cmp_ne_u32_e32 vcc_lo, 0, v10
	s_mov_b32 s1, exec_lo
                                        ; implicit-def: $vgpr8_vgpr9
                                        ; implicit-def: $vgpr1_vgpr2
	v_cmpx_le_i32_e64 s8, v12
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB1_13
; %bb.4:
                                        ; implicit-def: $vgpr8_vgpr9
                                        ; implicit-def: $vgpr1_vgpr2
	s_and_saveexec_b32 s0, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s12, exec_lo, s0
	s_cbranch_execz .LBB1_10
; %bb.5:
	v_ashrrev_i32_e32 v12, 31, v11
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, dev_mu@rel32@lo+4
	s_addc_u32 s9, s9, dev_mu@rel32@hi+12
	s_mov_b32 s13, exec_lo
	s_load_b64 s[8:9], s[8:9], 0x0
                                        ; implicit-def: $vgpr8_vgpr9
	v_lshlrev_b64 v[1:2], 3, v[11:12]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v1, s0, s4, v1
	v_add_co_ci_u32_e64 v2, s0, s5, v2, s0
	s_add_i32 s0, s10, -2
	global_load_b64 v[3:4], v[1:2], off offset:8
                                        ; implicit-def: $vgpr1_vgpr2
	v_cmpx_le_i32_e64 s0, v10
	s_xor_b32 s13, exec_lo, s13
	s_cbranch_execz .LBB1_7
; %bb.6:
	v_ashrrev_i32_e32 v8, 31, v7
	v_lshlrev_b64 v[1:2], 3, v[5:6]
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, dev_dx@rel32@lo+4
	s_addc_u32 s15, s15, dev_dx@rel32@hi+12
	s_load_b64 s[14:15], s[14:15], 0x0
	v_lshlrev_b64 v[7:8], 3, v[7:8]
	v_add_co_u32 v1, s0, s4, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v2, s0, s5, v2, s0
	v_add_co_u32 v12, s0, s6, v7
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v13, s0, s7, v8, s0
	global_load_b128 v[8:11], v[1:2], off offset:-8
	global_load_b64 v[1:2], v[12:13], off offset:8
	s_waitcnt vmcnt(1)
	v_add_f64 v[3:4], v[3:4], v[8:9]
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[7:8], s[14:15], s[14:15]
	s_waitcnt vmcnt(0)
	v_add_f64 v[1:2], -v[1:2], 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[3:4], v[10:11], -4.0, v[3:4]
	v_mul_f64 v[1:2], v[7:8], v[1:2]
                                        ; implicit-def: $vgpr7
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[8:9], s[8:9], v[3:4], -v[1:2]
	v_dual_mov_b32 v1, v10 :: v_dual_mov_b32 v2, v11
                                        ; implicit-def: $vgpr3_vgpr4
.LBB1_7:                                ; %Flow
	s_and_not1_saveexec_b32 s13, s13
	s_cbranch_execz .LBB1_9
; %bb.8:
	v_lshlrev_b64 v[1:2], 3, v[5:6]
	v_ashrrev_i32_e32 v8, 31, v7
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, dev_dx@rel32@lo+4
	s_addc_u32 s15, s15, dev_dx@rel32@hi+12
	s_load_b64 s[14:15], s[14:15], 0x0
	v_lshlrev_b64 v[7:8], 3, v[7:8]
	v_add_co_u32 v1, s0, s4, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v2, s0, s5, v2, s0
	v_add_co_u32 v14, s0, s6, v7
	global_load_b64 v[12:13], v[1:2], off offset:8
	v_add_co_ci_u32_e64 v15, s0, s7, v8, s0
	global_load_b128 v[8:11], v[1:2], off offset:-8
	global_load_b64 v[1:2], v[14:15], off offset:8
	s_waitcnt vmcnt(2)
	v_add_f64 v[3:4], v[3:4], v[12:13]
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[12:13], s[14:15], s[14:15]
	s_waitcnt vmcnt(0)
	v_add_f64 v[1:2], -v[1:2], 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[3:4], v[3:4], v[8:9]
	v_mul_f64 v[1:2], v[12:13], v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[3:4], 0xc0080000, v[10:11], v[3:4]
	v_fma_f64 v[8:9], s[8:9], v[3:4], -v[1:2]
	v_dual_mov_b32 v1, v10 :: v_dual_mov_b32 v2, v11
.LBB1_9:                                ; %Flow133
	s_or_b32 exec_lo, exec_lo, s13
                                        ; implicit-def: $vgpr11
                                        ; implicit-def: $vgpr7
.LBB1_10:                               ; %Flow134
	s_and_not1_saveexec_b32 s8, s12
	s_cbranch_execz .LBB1_12
; %bb.11:
	v_ashrrev_i32_e32 v12, 31, v11
	v_ashrrev_i32_e32 v8, 31, v7
	s_waitcnt vmcnt(0)
	v_lshlrev_b64 v[3:4], 3, v[5:6]
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, dev_dx@rel32@lo+4
	s_addc_u32 s13, s13, dev_dx@rel32@hi+12
	s_load_b64 s[12:13], s[12:13], 0x0
	v_lshlrev_b64 v[1:2], 3, v[11:12]
	v_lshlrev_b64 v[7:8], 3, v[7:8]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v1, s0, s4, v1
	v_add_co_ci_u32_e64 v2, s0, s5, v2, s0
	v_add_co_u32 v3, s0, s4, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e64 v4, s0, s5, v4, s0
	v_add_co_u32 v7, s0, s6, v7
	v_add_co_ci_u32_e64 v8, s0, s7, v8, s0
	s_clause 0x1
	global_load_b64 v[9:10], v[1:2], off offset:8
	global_load_b128 v[1:4], v[3:4], off
	global_load_b64 v[7:8], v[7:8], off offset:8
	v_add_f64 v[11:12], s[12:13], s[12:13]
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, dev_mu@rel32@lo+4
	s_addc_u32 s13, s13, dev_mu@rel32@hi+12
	s_load_b64 s[12:13], s[12:13], 0x0
	s_waitcnt vmcnt(1)
	v_add_f64 v[3:4], v[9:10], v[3:4]
	s_waitcnt vmcnt(0)
	v_add_f64 v[7:8], -v[7:8], 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[3:4], v[1:2], -4.0, v[3:4]
	v_mul_f64 v[7:8], v[11:12], v[7:8]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[8:9], s[12:13], v[3:4], -v[7:8]
.LBB1_12:                               ; %Flow135
	s_or_b32 exec_lo, exec_lo, s8
                                        ; implicit-def: $vgpr13
                                        ; implicit-def: $vgpr11
                                        ; implicit-def: $vgpr3_vgpr4
                                        ; implicit-def: $vgpr7
                                        ; implicit-def: $vgpr10
.LBB1_13:                               ; %Flow141
	s_and_not1_saveexec_b32 s8, s1
	s_cbranch_execz .LBB1_23
; %bb.14:
	v_ashrrev_i32_e32 v12, 31, v11
                                        ; implicit-def: $vgpr8_vgpr9
                                        ; implicit-def: $vgpr1_vgpr2
	s_and_saveexec_b32 s0, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s9, exec_lo, s0
	s_cbranch_execz .LBB1_20
; %bb.15:
	v_ashrrev_i32_e32 v14, 31, v13
	v_lshlrev_b64 v[8:9], 3, v[11:12]
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, dev_mu@rel32@lo+4
	s_addc_u32 s1, s1, dev_mu@rel32@hi+12
	s_add_i32 s12, s10, -2
	s_load_b64 s[0:1], s[0:1], 0x0
	v_lshlrev_b64 v[1:2], 3, v[13:14]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_add_co_u32 v8, vcc_lo, s4, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s5, v9, vcc_lo
	v_cmp_le_i32_e32 vcc_lo, s12, v10
	s_clause 0x1
	global_load_b64 v[1:2], v[1:2], off
	global_load_b64 v[8:9], v[8:9], off offset:8
	s_waitcnt vmcnt(0)
	v_add_f64 v[11:12], v[1:2], v[8:9]
                                        ; implicit-def: $vgpr8_vgpr9
                                        ; implicit-def: $vgpr1_vgpr2
	s_and_saveexec_b32 s12, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s12, exec_lo, s12
	s_cbranch_execz .LBB1_17
; %bb.16:
	v_ashrrev_i32_e32 v4, 31, v3
	v_lshlrev_b64 v[1:2], 3, v[5:6]
	v_ashrrev_i32_e32 v8, 31, v7
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, dev_dx@rel32@lo+4
	s_addc_u32 s15, s15, dev_dx@rel32@hi+12
	s_load_b64 s[14:15], s[14:15], 0x0
	v_lshlrev_b64 v[3:4], 3, v[3:4]
	v_lshlrev_b64 v[7:8], 3, v[7:8]
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v9, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v4, vcc_lo
	v_add_co_u32 v7, vcc_lo, s6, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	global_load_b128 v[1:4], v[1:2], off offset:-8
	s_clause 0x1
	global_load_b64 v[9:10], v[9:10], off
	global_load_b64 v[7:8], v[7:8], off offset:8
	s_waitcnt vmcnt(2)
	v_add_f64 v[1:2], v[11:12], v[1:2]
                                        ; implicit-def: $vgpr11_vgpr12
	s_waitcnt vmcnt(0)
	v_add_f64 v[7:8], v[9:10], -v[7:8]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[1:2], 0xc0140000, v[3:4], v[1:2]
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[7:8], s[14:15], v[7:8]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[8:9], s[0:1], v[1:2], -v[7:8]
	v_dual_mov_b32 v1, v3 :: v_dual_mov_b32 v2, v4
                                        ; implicit-def: $vgpr3_vgpr4
                                        ; implicit-def: $vgpr7
.LBB1_17:                               ; %Flow137
	s_and_not1_saveexec_b32 s12, s12
	s_cbranch_execz .LBB1_19
; %bb.18:
	v_lshlrev_b64 v[1:2], 3, v[5:6]
	v_ashrrev_i32_e32 v4, 31, v3
	v_ashrrev_i32_e32 v8, 31, v7
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, dev_dx@rel32@lo+4
	s_addc_u32 s15, s15, dev_dx@rel32@hi+12
	s_load_b64 s[14:15], s[14:15], 0x0
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_lshlrev_b64 v[3:4], 3, v[3:4]
	v_lshlrev_b64 v[7:8], 3, v[7:8]
	global_load_b64 v[9:10], v[1:2], off offset:8
	v_add_co_u32 v13, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v14, vcc_lo, s7, v4, vcc_lo
	v_add_co_u32 v7, vcc_lo, s6, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	global_load_b128 v[1:4], v[1:2], off offset:-8
	s_clause 0x1
	global_load_b64 v[13:14], v[13:14], off
	global_load_b64 v[7:8], v[7:8], off offset:8
	s_waitcnt vmcnt(3)
	v_add_f64 v[9:10], v[11:12], v[9:10]
	s_waitcnt vmcnt(0)
	v_add_f64 v[7:8], v[13:14], -v[7:8]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[1:2], v[9:10], v[1:2]
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[7:8], s[14:15], v[7:8]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[1:2], v[3:4], -4.0, v[1:2]
	v_fma_f64 v[8:9], s[0:1], v[1:2], -v[7:8]
	v_dual_mov_b32 v1, v3 :: v_dual_mov_b32 v2, v4
.LBB1_19:                               ; %Flow138
	s_or_b32 exec_lo, exec_lo, s12
                                        ; implicit-def: $vgpr13
                                        ; implicit-def: $vgpr11
                                        ; implicit-def: $vgpr3_vgpr4
                                        ; implicit-def: $vgpr7
.LBB1_20:                               ; %Flow139
	s_and_not1_saveexec_b32 s0, s9
	s_cbranch_execz .LBB1_22
; %bb.21:
	v_ashrrev_i32_e32 v14, 31, v13
	s_delay_alu instid0(VALU_DEP_2)
	v_lshlrev_b64 v[8:9], 3, v[11:12]
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v4, 31, v3
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, dev_dx@rel32@lo+4
	s_addc_u32 s13, s13, dev_dx@rel32@hi+12
	s_load_b64 s[12:13], s[12:13], 0x0
	v_lshlrev_b64 v[1:2], 3, v[13:14]
	v_lshlrev_b64 v[3:4], 3, v[3:4]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_add_co_u32 v8, vcc_lo, s4, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s5, v9, vcc_lo
	s_clause 0x1
	global_load_b64 v[10:11], v[1:2], off
	global_load_b64 v[12:13], v[8:9], off offset:8
	v_lshlrev_b64 v[1:2], 3, v[5:6]
	v_ashrrev_i32_e32 v8, 31, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[7:8], 3, v[7:8]
	v_add_co_u32 v1, vcc_lo, s4, v1
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_add_co_u32 v14, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v15, vcc_lo, s7, v4, vcc_lo
	v_add_co_u32 v7, vcc_lo, s6, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	global_load_b128 v[1:4], v[1:2], off
	s_clause 0x1
	global_load_b64 v[14:15], v[14:15], off
	global_load_b64 v[7:8], v[7:8], off offset:8
	s_waitcnt vmcnt(3)
	v_add_f64 v[9:10], v[10:11], v[12:13]
	s_waitcnt vmcnt(0)
	v_add_f64 v[7:8], v[14:15], -v[7:8]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[3:4], v[9:10], v[3:4]
	v_mul_f64 v[7:8], s[12:13], v[7:8]
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, dev_mu@rel32@lo+4
	s_addc_u32 s13, s13, dev_mu@rel32@hi+12
	s_load_b64 s[12:13], s[12:13], 0x0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[3:4], 0xc0140000, v[1:2], v[3:4]
	s_waitcnt lgkmcnt(0)
	v_fma_f64 v[8:9], s[12:13], v[3:4], -v[7:8]
.LBB1_22:                               ; %Flow140
	s_or_b32 exec_lo, exec_lo, s0
.LBB1_23:                               ; %Flow142
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s8
                                        ; implicit-def: $vgpr13
                                        ; implicit-def: $vgpr3_vgpr4
                                        ; implicit-def: $vgpr10
.LBB1_24:                               ; %Flow147
	s_and_not1_saveexec_b32 s8, s11
	s_cbranch_execz .LBB1_34
; %bb.25:
	v_ashrrev_i32_e32 v14, 31, v13
                                        ; implicit-def: $vgpr8_vgpr9
                                        ; implicit-def: $vgpr1_vgpr2
	s_mov_b32 s0, exec_lo
	v_cmpx_ne_u32_e32 0, v10
	s_xor_b32 s9, exec_lo, s0
	s_cbranch_execz .LBB1_31
; %bb.26:
	v_lshlrev_b64 v[1:2], 3, v[13:14]
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, dev_mu@rel32@lo+4
	s_addc_u32 s1, s1, dev_mu@rel32@hi+12
	s_add_i32 s10, s10, -2
	s_load_b64 s[0:1], s[0:1], 0x0
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v4, 31, v3
                                        ; implicit-def: $vgpr8_vgpr9
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_cmp_le_i32_e32 vcc_lo, s10, v10
	global_load_b64 v[11:12], v[1:2], off
                                        ; implicit-def: $vgpr1_vgpr2
	s_and_saveexec_b32 s10, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s10, exec_lo, s10
	s_cbranch_execz .LBB1_28
; %bb.27:
	v_lshlrev_b64 v[1:2], 3, v[5:6]
	v_lshlrev_b64 v[3:4], 3, v[3:4]
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, dev_dx@rel32@lo+4
	s_addc_u32 s13, s13, dev_dx@rel32@hi+12
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, dev_P@rel32@lo+4
	s_addc_u32 s15, s15, dev_P@rel32@hi+12
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_add_co_u32 v7, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v4, vcc_lo
	global_load_b128 v[1:4], v[1:2], off offset:-8
	global_load_b64 v[7:8], v[7:8], off
	s_load_b64 s[12:13], s[12:13], 0x0
	s_load_b64 s[14:15], s[14:15], 0x0
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[9:10], s[12:13], s[12:13]
	s_waitcnt vmcnt(1)
	v_add_f64 v[1:2], v[11:12], v[1:2]
	s_waitcnt vmcnt(0)
	v_add_f64 v[7:8], v[7:8], -s[14:15]
                                        ; implicit-def: $vgpr11_vgpr12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[1:2], v[3:4], -4.0, v[1:2]
	v_mul_f64 v[7:8], v[9:10], v[7:8]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[8:9], s[0:1], v[1:2], -v[7:8]
	v_dual_mov_b32 v1, v3 :: v_dual_mov_b32 v2, v4
                                        ; implicit-def: $vgpr3_vgpr4
.LBB1_28:                               ; %Flow143
	s_and_not1_saveexec_b32 s10, s10
	s_cbranch_execz .LBB1_30
; %bb.29:
	v_lshlrev_b64 v[1:2], 3, v[5:6]
	v_lshlrev_b64 v[3:4], 3, v[3:4]
	s_getpc_b64 s[12:13]
	s_add_u32 s12, s12, dev_dx@rel32@lo+4
	s_addc_u32 s13, s13, dev_dx@rel32@hi+12
	s_getpc_b64 s[14:15]
	s_add_u32 s14, s14, dev_P@rel32@lo+4
	s_addc_u32 s15, s15, dev_P@rel32@hi+12
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_add_co_u32 v9, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v4, vcc_lo
	s_clause 0x1
	global_load_b64 v[7:8], v[1:2], off offset:8
	global_load_b128 v[1:4], v[1:2], off offset:-8
	global_load_b64 v[9:10], v[9:10], off
	s_load_b64 s[12:13], s[12:13], 0x0
	s_load_b64 s[14:15], s[14:15], 0x0
	s_waitcnt vmcnt(2)
	v_add_f64 v[7:8], v[11:12], v[7:8]
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[11:12], s[12:13], s[12:13]
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_add_f64 v[1:2], v[7:8], v[1:2]
	s_waitcnt vmcnt(0)
	v_add_f64 v[7:8], v[9:10], -s[14:15]
	v_fma_f64 v[1:2], 0xc0080000, v[3:4], v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[7:8], v[11:12], v[7:8]
	v_fma_f64 v[8:9], s[0:1], v[1:2], -v[7:8]
	v_dual_mov_b32 v1, v3 :: v_dual_mov_b32 v2, v4
.LBB1_30:                               ; %Flow144
	s_or_b32 exec_lo, exec_lo, s10
                                        ; implicit-def: $vgpr13
                                        ; implicit-def: $vgpr3_vgpr4
.LBB1_31:                               ; %Flow145
	s_and_not1_saveexec_b32 s0, s9
	s_cbranch_execz .LBB1_33
; %bb.32:
	v_lshlrev_b64 v[1:2], 3, v[13:14]
	s_waitcnt vmcnt(0)
	v_mov_b32_e32 v4, 0
	v_lshlrev_b64 v[7:8], 3, v[5:6]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[3:4], 3, v[3:4]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_add_co_u32 v7, vcc_lo, s4, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s5, v8, vcc_lo
	v_add_co_u32 v9, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v4, vcc_lo
	s_clause 0x1
	global_load_b64 v[11:12], v[1:2], off
	global_load_b128 v[1:4], v[7:8], off
	global_load_b64 v[7:8], v[9:10], off
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, dev_dx@rel32@lo+4
	s_addc_u32 s7, s7, dev_dx@rel32@hi+12
	s_getpc_b64 s[10:11]
	s_add_u32 s10, s10, dev_P@rel32@lo+4
	s_addc_u32 s11, s11, dev_P@rel32@hi+12
	s_load_b64 s[6:7], s[6:7], 0x0
	s_load_b64 s[10:11], s[10:11], 0x0
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[9:10], s[6:7], s[6:7]
	s_getpc_b64 s[6:7]
	s_add_u32 s6, s6, dev_mu@rel32@lo+4
	s_addc_u32 s7, s7, dev_mu@rel32@hi+12
	s_load_b64 s[6:7], s[6:7], 0x0
	s_waitcnt vmcnt(1)
	v_add_f64 v[3:4], v[11:12], v[3:4]
	s_waitcnt vmcnt(0)
	v_add_f64 v[7:8], v[7:8], -s[10:11]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[3:4], v[1:2], -4.0, v[3:4]
	v_mul_f64 v[7:8], v[9:10], v[7:8]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[8:9], s[6:7], v[3:4], -v[7:8]
.LBB1_33:                               ; %Flow146
	s_or_b32 exec_lo, exec_lo, s0
.LBB1_34:
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s8
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, dev_w@rel32@lo+4
	s_addc_u32 s1, s1, dev_w@rel32@hi+12
	s_waitcnt vmcnt(0)
	v_lshlrev_b64 v[3:4], 3, v[5:6]
	s_load_b64 s[0:1], s[0:1], 0x0
	v_lshlrev_b32_e32 v5, 3, v0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	v_fma_f64 v[1:2], v[8:9], s[0:1], v[1:2]
	global_store_b64 v[3:4], v[1:2], off
	global_load_b64 v[0:1], v5, s[2:3]
	v_max_f64 v[2:3], |v[8:9]|, |v[8:9]|
	s_waitcnt vmcnt(0)
	v_max_f64 v[0:1], v[0:1], v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_max_f64 v[0:1], v[2:3], v[0:1]
	global_store_b64 v5, v[0:1], s[2:3]
.LBB1_35:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z7updateUPdS_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 16
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z7updateUPdS_S_i, .Lfunc_end1-_Z7updateUPdS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 3044
; NumSgprs: 18
; NumVgprs: 16
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 16
; 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
	.protected	_Z7updateVPdS_S_i       ; -- Begin function _Z7updateVPdS_S_i
	.globl	_Z7updateVPdS_S_i
	.p2align	8
	.type	_Z7updateVPdS_S_i,@function
_Z7updateVPdS_S_i:                      ; @_Z7updateVPdS_S_i
; %bb.0:
	s_load_b32 s4, s[0:1], 0x2c
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, dev_N@rel32@lo+4
	s_addc_u32 s3, s3, dev_N@rel32@hi+12
	s_load_b32 s8, s[2:3], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[5:6], null, s15, s2, v[0:1]
	s_add_i32 s3, s8, -1
	s_mul_i32 s2, s3, s8
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_i32_e32 vcc_lo, s2, v5
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB2_12
; %bb.1:
	s_ashr_i32 s9, s8, 31
	v_ashrrev_i32_e32 v6, 31, v5
	s_add_i32 s2, s8, s9
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_xor_b32 s2, s2, s9
	v_add_nc_u32_e32 v3, v5, v6
	v_cvt_f32_u32_e32 v1, s2
	s_sub_i32 s4, 0, s2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v3, v3, v6
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v1, v1
	v_mul_lo_u32 v2, s4, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v1, v2
	v_add_nc_u32_e32 v1, v1, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v1, v3, v1
	v_mul_lo_u32 v2, v1, s2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v2, v3, v2
	v_add_nc_u32_e32 v3, 1, v1
	v_subrev_nc_u32_e32 v4, s2, v2
	v_cmp_le_u32_e32 vcc_lo, s2, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v2, v2, v4 :: v_dual_cndmask_b32 v1, v1, v3
	v_xor_b32_e32 v4, s9, v6
	v_cmp_le_u32_e32 vcc_lo, s2, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, 1, v1
	v_cndmask_b32_e32 v1, v1, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v1, v1, v4
	v_sub_nc_u32_e32 v14, v1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v1, v14, s8
	v_sub_nc_u32_e32 v1, v5, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_lt_i32_e32 vcc_lo, 0, v1
	v_cmp_gt_i32_e64 s2, s3, v1
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 exec_lo, exec_lo, s2
	s_cbranch_execz .LBB2_12
; %bb.2:
	v_add_nc_u32_e32 v2, v14, v1
	s_load_b32 s2, s[0:1], 0x18
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v3, 31, v2
	v_add_nc_u32_e32 v3, v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v3, -2, v3
	v_sub_nc_u32_e32 v2, v2, v3
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_eq_u32_e32 vcc_lo, s2, v2
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB2_12
; %bb.3:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_mad_u64_u32 v[7:8], null, v14, s3, v[1:2]
                                        ; implicit-def: $vgpr3_vgpr4
                                        ; implicit-def: $vgpr1_vgpr2
	s_mov_b32 s2, exec_lo
	v_cmpx_ne_u32_e32 0, v14
	s_xor_b32 s10, exec_lo, s2
	s_cbranch_execz .LBB2_9
; %bb.4:
	v_lshlrev_b64 v[1:2], 3, v[5:6]
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, dev_mu@rel32@lo+4
	s_addc_u32 s3, s3, dev_mu@rel32@hi+12
	v_subrev_nc_u32_e32 v13, s8, v5
	s_load_b64 s[2:3], s[2:3], 0x0
	s_add_i32 s11, s8, -2
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v9, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v10, vcc_lo, s5, v2, vcc_lo
	v_cmp_le_i32_e32 vcc_lo, s11, v14
	v_ashrrev_i32_e32 v14, 31, v13
	s_clause 0x1
	global_load_b64 v[1:2], v[9:10], off offset:8
	global_load_b64 v[3:4], v[9:10], off offset:-8
	s_waitcnt vmcnt(0)
	v_add_f64 v[11:12], v[1:2], v[3:4]
                                        ; implicit-def: $vgpr3_vgpr4
                                        ; implicit-def: $vgpr1_vgpr2
	s_and_saveexec_b32 s11, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s11, exec_lo, s11
	s_cbranch_execz .LBB2_6
; %bb.5:
	v_lshlrev_b64 v[1:2], 3, v[13:14]
	v_ashrrev_i32_e32 v8, 31, v7
                                        ; implicit-def: $vgpr13
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	s_clause 0x1
	global_load_b64 v[1:2], v[1:2], off
	global_load_b64 v[3:4], v[9:10], off
                                        ; implicit-def: $vgpr9
	s_waitcnt vmcnt(1)
	v_add_f64 v[1:2], v[11:12], v[1:2]
                                        ; implicit-def: $vgpr11_vgpr12
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[1:2], 0xc0080000, v[3:4], v[1:2]
	v_mul_f64 v[1:2], s[2:3], v[1:2]
.LBB2_6:                                ; %Flow
	s_and_not1_saveexec_b32 s11, s11
	s_cbranch_execz .LBB2_8
; %bb.7:
	s_lshl_b64 s[12:13], s[8:9], 3
	v_lshlrev_b64 v[3:4], 3, v[13:14]
	v_add_co_u32 v1, vcc_lo, v9, s12
	v_add_co_ci_u32_e32 v2, vcc_lo, s13, v10, vcc_lo
	v_ashrrev_i32_e32 v8, 31, v7
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v3, vcc_lo, s4, v3
	global_load_b64 v[1:2], v[1:2], off
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	s_clause 0x1
	global_load_b64 v[13:14], v[3:4], off
	global_load_b64 v[3:4], v[9:10], off
	s_waitcnt vmcnt(2)
	v_add_f64 v[1:2], v[11:12], v[1:2]
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_f64 v[1:2], v[1:2], v[13:14]
	s_waitcnt vmcnt(0)
	v_fma_f64 v[1:2], v[3:4], -4.0, v[1:2]
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f64 v[1:2], s[2:3], v[1:2]
.LBB2_8:                                ; %Flow87
	s_or_b32 exec_lo, exec_lo, s11
.LBB2_9:                                ; %Flow88
	s_and_not1_saveexec_b32 s2, s10
	s_cbranch_execz .LBB2_11
; %bb.10:
	v_lshlrev_b64 v[1:2], 3, v[5:6]
	s_lshl_b64 s[8:9], s[8:9], 3
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v8, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v9, vcc_lo, s5, v2, vcc_lo
	s_clause 0x1
	global_load_b64 v[10:11], v[8:9], off offset:8
	global_load_b128 v[1:4], v[8:9], off offset:-8
	v_add_co_u32 v8, vcc_lo, v8, s8
	v_add_co_ci_u32_e32 v9, vcc_lo, s9, v9, vcc_lo
	s_getpc_b64 s[8:9]
	s_add_u32 s8, s8, dev_mu@rel32@lo+4
	s_addc_u32 s9, s9, dev_mu@rel32@hi+12
	s_load_b64 s[8:9], s[8:9], 0x0
	global_load_b64 v[8:9], v[8:9], off
	s_waitcnt vmcnt(1)
	v_add_f64 v[1:2], v[10:11], v[1:2]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[1:2], v[1:2], v[8:9]
	v_mov_b32_e32 v8, 0
	v_fma_f64 v[1:2], 0xc0080000, v[3:4], v[1:2]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f64 v[1:2], s[8:9], v[1:2]
.LBB2_11:
	s_or_b32 exec_lo, exec_lo, s2
	v_lshlrev_b64 v[9:10], 3, v[7:8]
	v_ashrrev_i32_e32 v8, 31, v7
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, dev_dx@rel32@lo+4
	s_addc_u32 s3, s3, dev_dx@rel32@hi+12
	v_lshlrev_b64 v[5:6], 3, v[5:6]
	s_load_b64 s[2:3], s[2:3], 0x0
	v_lshlrev_b64 v[7:8], 3, v[7:8]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v9, vcc_lo, s6, v9
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v10, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v7, vcc_lo, s6, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	v_add_co_u32 v5, vcc_lo, s4, v5
	s_clause 0x1
	global_load_b64 v[9:10], v[9:10], off
	global_load_b64 v[7:8], v[7:8], off offset:-8
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v6, vcc_lo
	s_waitcnt vmcnt(0)
	v_add_f64 v[7:8], v[9:10], -v[7:8]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[1:2], -s[2:3], v[7:8], v[1:2]
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, dev_w@rel32@lo+4
	s_addc_u32 s3, s3, dev_w@rel32@hi+12
	v_lshlrev_b32_e32 v7, 3, v0
	s_load_b64 s[2:3], s[2:3], 0x0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_fma_f64 v[3:4], v[1:2], s[2:3], v[3:4]
	global_store_b64 v[5:6], v[3:4], off
	global_load_b64 v[3:4], v7, s[0:1]
	s_waitcnt vmcnt(0)
	v_max_f64 v[3:4], v[3:4], v[3:4]
	v_max_f64 v[0:1], |v[1:2]|, v[3:4]
	global_store_b64 v7, v[0:1], s[0:1]
.LBB2_12:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z7updateVPdS_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 15
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end2:
	.size	_Z7updateVPdS_S_i, .Lfunc_end2-_Z7updateVPdS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1164
; NumSgprs: 18
; NumVgprs: 15
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 15
; Occupancy: 16
; WaveLimiterHint : 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
	.protected	_Z7updatePPdS_S_S_      ; -- Begin function _Z7updatePPdS_S_S_
	.globl	_Z7updatePPdS_S_S_
	.p2align	8
	.type	_Z7updatePPdS_S_S_,@function
_Z7updatePPdS_S_S_:                     ; @_Z7updatePPdS_S_S_
; %bb.0:
	s_load_b32 s4, s[0:1], 0x2c
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, dev_N@rel32@lo+4
	s_addc_u32 s3, s3, dev_N@rel32@hi+12
	s_load_b32 s8, s[2:3], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_add_i32 s9, s8, -1
	s_mul_i32 s2, s9, s9
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_cmp_gt_i32_e32 vcc_lo, s2, v1
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB3_2
; %bb.1:
	s_ashr_i32 s2, s9, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s3, s9, s2
	s_xor_b32 s3, s3, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cvt_f32_u32_e32 v2, s3
	s_sub_i32 s4, 0, s3
	v_rcp_iflag_f32_e32 v2, v2
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v2, 0x4f7ffffe, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cvt_u32_f32_e32 v3, v2
	v_ashrrev_i32_e32 v2, 31, v1
	v_mul_lo_u32 v4, s4, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v1, v2
	v_xor_b32_e32 v5, v5, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v4, v3, v4
	v_add_nc_u32_e32 v3, v3, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v3, v5, v3
	v_mul_lo_u32 v4, v3, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v4, v5, v4
	v_subrev_nc_u32_e32 v6, s3, v4
	v_cmp_le_u32_e32 vcc_lo, s3, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v4, v4, v6 :: v_dual_add_nc_u32 v5, 1, v3
	v_cndmask_b32_e32 v3, v3, v5, vcc_lo
	v_xor_b32_e32 v6, s2, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_le_u32_e32 vcc_lo, s3, v4
	v_add_nc_u32_e32 v5, 1, v3
	s_load_b256 s[0:7], s[0:1], 0x0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v3, v3, v5, vcc_lo
	v_xor_b32_e32 v3, v3, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v8, v3, v6
	v_mul_lo_u32 v3, v8, s9
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v3, v1, v3
	v_lshlrev_b64 v[1:2], 3, v[1:2]
	v_mad_u64_u32 v[4:5], null, v8, s9, v[3:4]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[6:7], null, v8, s8, v[3:4]
	v_add_nc_u32_e32 v8, s9, v4
	v_ashrrev_i32_e32 v5, 31, v4
	v_ashrrev_i32_e32 v9, 31, v8
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v7, 31, v6
	v_lshlrev_b64 v[3:4], 3, v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[5:6], 3, v[6:7]
	v_lshlrev_b64 v[7:8], 3, v[8:9]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v3, vcc_lo, s0, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s1, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v7, vcc_lo, s0, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s1, v8, vcc_lo
	v_add_co_u32 v5, vcc_lo, s2, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s3, v6, vcc_lo
	s_clause 0x1
	global_load_b64 v[7:8], v[7:8], off
	global_load_b64 v[9:10], v[3:4], off
	global_load_b128 v[3:6], v[5:6], off
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, dev_w@rel32@lo+4
	s_addc_u32 s1, s1, dev_w@rel32@hi+12
	s_load_b64 s[0:1], s[0:1], 0x0
	global_load_b64 v[11:12], v[1:2], off
	s_waitcnt vmcnt(2)
	v_add_f64 v[7:8], v[7:8], -v[9:10]
	s_waitcnt vmcnt(1)
	v_add_f64 v[3:4], v[5:6], -v[3:4]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_add_f64 v[3:4], -v[7:8], -v[3:4]
	v_lshlrev_b32_e32 v7, 3, v0
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_fma_f64 v[5:6], s[0:1], v[3:4], v[11:12]
	global_store_b64 v[1:2], v[5:6], off
	global_load_b64 v[0:1], v7, s[6:7]
	s_waitcnt vmcnt(0)
	v_max_f64 v[0:1], v[0:1], v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_max_f64 v[0:1], |v[3:4]|, v[0:1]
	global_store_b64 v7, v[0:1], s[6:7]
.LBB3_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z7updatePPdS_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 13
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end3:
	.size	_Z7updatePPdS_S_S_, .Lfunc_end3-_Z7updatePPdS_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 600
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; 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
	.protected	_Z10findmaxresPdi       ; -- Begin function _Z10findmaxresPdi
	.globl	_Z10findmaxresPdi
	.p2align	8
	.type	_Z10findmaxresPdi,@function
_Z10findmaxresPdi:                      ; @_Z10findmaxresPdi
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x8
	s_load_b32 s3, s[0:1], 0x1c
	s_waitcnt lgkmcnt(0)
	s_lshl_b32 s4, s2, 1
	s_bfe_i32 s5, s2, 0x1001e
	s_and_b32 s3, s3, 0xffff
	s_add_i32 s4, s4, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_xor_b32 s4, s4, s5
	v_cvt_f32_u32_e32 v1, s4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v3, v1
	v_mad_u64_u32 v[1:2], null, s15, s3, v[0:1]
	s_sub_i32 s3, 0, s4
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	v_mul_lo_u32 v0, s3, v3
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v2, 31, v1
	v_mul_hi_u32 v0, v3, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v4, v1, v2
	v_xor_b32_e32 v4, v4, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v3, v0
	v_mul_hi_u32 v0, v4, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, v0, s4
	v_sub_nc_u32_e32 v0, v4, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v3, s4, v0
	v_cmp_le_u32_e32 vcc_lo, s4, v0
	v_cndmask_b32_e32 v0, v0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v3, s4, v0
	v_cmp_le_u32_e32 vcc_lo, s4, v0
	v_cndmask_b32_e32 v0, v0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v0, v0, v2
	v_sub_nc_u32_e32 v0, v0, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB4_2
; %bb.1:
	s_load_b64 s[0:1], s[0:1], 0x0
	v_add_nc_u32_e32 v3, s2, v1
	v_lshlrev_b64 v[0:1], 3, v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v4, 31, v3
	v_lshlrev_b64 v[3:4], 3, v[3:4]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s0, v3
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v4, vcc_lo
	v_add_co_u32 v0, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_clause 0x1
	global_load_b64 v[2:3], v[2:3], off
	global_load_b64 v[4:5], v[0:1], off
	s_waitcnt vmcnt(1)
	v_max_f64 v[2:3], v[2:3], v[2:3]
	s_waitcnt vmcnt(0)
	v_max_f64 v[4:5], v[4:5], v[4:5]
	s_delay_alu instid0(VALU_DEP_1)
	v_max_f64 v[2:3], v[4:5], v[2:3]
	global_store_b64 v[0:1], v[2:3], off
.LBB4_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z10findmaxresPdi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.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 6
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end4:
	.size	_Z10findmaxresPdi, .Lfunc_end4-_Z10findmaxresPdi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 372
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; 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
	.protected	dev_N                   ; @dev_N
	.type	dev_N,@object
	.section	.bss,"aw",@nobits
	.globl	dev_N
	.p2align	2, 0x0
dev_N:
	.long	0                               ; 0x0
	.size	dev_N, 4

	.protected	dev_mu                  ; @dev_mu
	.type	dev_mu,@object
	.globl	dev_mu
	.p2align	3, 0x0
dev_mu:
	.quad	0x0000000000000000              ; double 0
	.size	dev_mu, 8

	.protected	dev_P                   ; @dev_P
	.type	dev_P,@object
	.globl	dev_P
	.p2align	3, 0x0
dev_P:
	.quad	0x0000000000000000              ; double 0
	.size	dev_P, 8

	.protected	dev_w                   ; @dev_w
	.type	dev_w,@object
	.globl	dev_w
	.p2align	3, 0x0
dev_w:
	.quad	0x0000000000000000              ; double 0
	.size	dev_w, 8

	.protected	dev_dx                  ; @dev_dx
	.type	dev_dx,@object
	.globl	dev_dx
	.p2align	3, 0x0
dev_dx:
	.quad	0x0000000000000000              ; double 0
	.size	dev_dx, 8

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.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 dev_N
	.addrsig_sym dev_mu
	.addrsig_sym dev_P
	.addrsig_sym dev_w
	.addrsig_sym dev_dx
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z4initPdi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z4initPdi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z7updateUPdS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z7updateUPdS_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     16
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z7updateVPdS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z7updateVPdS_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     15
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .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
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z7updatePPdS_S_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z7updatePPdS_S_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z10findmaxresPdi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z10findmaxresPdi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
