//nvcc -ptx EM4.cu -ccbin "F:Visual Studio\VC\Tools\MSVC\14.12.25827\bin\Hostx64\x64"
__device__ void EM1( double * x, 
                     double * y, 
                     double * z, 
                     double * vx,
                     double * vy,
                     double * vz,
                     double * r,
                     double * phi,
                     double * vr,
                     double * Er,
                     double * Ez,
                     double * Hphi,
                     double * charge,
                     double * m,
                     double * E,
                     int *parDelete,
                     const int parNum,
                     const double Rp,
                     const double Lp,
                     const double PHIp,
                     const int gridR,
                     const int gridZ,
                     const double mu,
                     const double c,
                     const double dr,
                     const double dz,
                     const double dt ) {
    int globalBlockIndex = blockIdx.x + blockIdx.y * gridDim.x;
    int localThreadIdx = threadIdx.x + blockDim.x * threadIdx.y;
    int threadsPerBlock = blockDim.x*blockDim.y;
    int n = localThreadIdx + globalBlockIndex*threadsPerBlock;
    if ( n >= parNum || parDelete[n]==1){
        return;
    }
    double r0 = r[n];
    double z0 = z[n];
    int ar,az,a1,a2,a3,a4;
    double Er0,Ez0,Hphi0;

    // Er
    ar = floor(r0/dr-0.5);
    az = floor(z0/dz);
    a1 = ar + az * gridR;
    a2 = (ar+1) + az * gridR;
    a3 = ar + (az+1) * gridR;
    a4 = (ar+1) + (az+1) * gridR;
    if( ar<0 ){
        Er0 = ( Er[a2]*(az*dz+dz-z0) + Er[a4]*(z0-az*dz) )/dz;
    }
    else if( ar >= gridR-1 ){
        Er0 = ( Er[a1]*(az*dz+dz-z0) + Er[a3]*(z0-az*dz) )/dz;
    }
    else{
        Er0 = ( Er[a1]*((ar+1.5)*dr-r0)*((az+1)*dz-z0)
                 + Er[a2]*(r0-(ar+0.5)*dr)*((az+1)*dz-z0)
                 + Er[a3]*((ar+1.5)*dr-r0)*(z0-az*dz)
                 + Er[a4]*(r0-(ar+0.5)*dr)*(z0-az*dz) )/(dr*dz);
        
    }
    // Ez
    ar = floor(r0/dr);
    az = floor(z0/dz-0.5);
    a1 = ar + az * (gridR+1);
    a2 = (ar+1) + az * (gridR+1);
    a3 = ar + (az+1) * (gridR+1);
    a4 = (ar+1) + (az+1) * (gridR+1);
    if( az<0 ){
        Ez0 = ( Ez[a3]*(ar*dr+dr-r0) + Ez[a4]*(r0-ar*dr) )/dr;
    }
    else if( az >= gridZ-1 ){
        Ez0 = ( Ez[a1]*(ar*dr+dr-r0) + Ez[a2]*(r0-ar*dr) )/dr;
    }
    else{
        Ez0 = ( Ez[a1]*((ar+1)*dr-r0)*((az+1.5)*dz-z0)
                 + Ez[a2]*(r0-ar*dr)*((az+1.5)*dz-z0)
                 + Ez[a3]*((ar+1)*dr-r0)*(z0-(az+0.5)*dz)
                 + Ez[a4]*(r0-ar*dr)*(z0-(az+0.5)*dz) )/(dr*dz);
        
    }

    // Hphi
    ar = floor(r0/dr-0.5);
    az = floor(z0/dz-0.5);
    a1 = ar + az * gridR;
    a2 = (ar+1) + az * gridR;
    a3 = ar + (az+1) * gridR;
    a4 = (ar+1) + (az+1) * gridR;
    if( ar<0 ){
        if( az<0 ){
            Hphi0 = Hphi[a4];
        }
        else if( az>=gridZ-1 ){
            Hphi0 = Hphi[a2];
        }
        else{
            Hphi0 = ( Hphi[a2]*((az+1.5)*dz-z0) + Hphi[a4]*(z0-(az+0.5)*dz) )/dz;
        }
    }
    else if( ar>=gridR-1 ){
        if( az<0 ){
            Hphi0 = Hphi[a3];
        }
        else if( az>=gridZ-1 ){
            Hphi0 = Hphi[a1];
        }
        else{
            Hphi0 = ( Hphi[a1]*((az+1.5)*dz-z0) + Hphi[a3]*(z0-(az+0.5)*dz) )/dz;
        }
    }
    else if( az<0 ){
        Hphi0 = ( Hphi[a3]*((ar+1.5)*dr-r0) + Hphi[a4]*(r0-(ar+0.5)*dr) )/dr;
    }
    else if( az>=gridZ-1 ){
        Hphi0 = ( Hphi[a1]*((ar+1.5)*dr-r0) + Hphi[a2]*(r0-(ar+0.5)*dr) )/dr;
    }
    else{
        Hphi0 = ( Hphi[a1]*((ar+1.5)*dr-r0)*((az+1.5)*dz-z0)
               + Hphi[a2]*(r0-(ar+0.5)*dr)*((az+1.5)*dz-z0)
               + Hphi[a3]*((ar+1.5)*dr-r0)*(z0-(az+0.5)*dz)
               + Hphi[a4]*(r0-(ar+0.5)*dr)*(z0-(az+0.5)*dz) )/(dr*dz);
    }

    //F
    double Fx,Fy,Fz,Fr;
    Fr = charge[n] * (Er0 + (-vz[n]*Hphi0*mu));
    Fz = charge[n] * (Ez0 + (vr[n]*Hphi0*mu));
    Fx = Fr*cos(phi[n]);
    Fy = Fr*sin(phi[n]);


    //v
    double gamma,ux,uy,uz;
    gamma = 1/sqrt( 1-( vx[n]*vx[n] + vy[n]*vy[n] + vz[n]*vz[n] )/(c*c) );
    ux = gamma * vx[n] + Fx/m[n]*dt;
    uy = gamma * vy[n] + Fy/m[n]*dt;
    uz = gamma * vz[n] + Fz/m[n]*dt;
    gamma = sqrt( 1+ (ux*ux + uy*uy + uz*uz)/(c*c) );

    E[n] = (gamma-1)*m[n]*c*c;

    vx[n] = ux/gamma;
    vy[n] = uy/gamma;
    vz[n] = uz/gamma;

    x[n] = x[n] + 0.5*dt*vx[n];
    y[n] = y[n] + 0.5*dt*vy[n];
    z[n] = z[n] + 0.5*dt*vz[n];

    r[n] = sqrt( x[n]*x[n] + y[n]*y[n] );
    phi[n] = atan( y[n]/x[n] );
    if (r[n] > Rp){
        parDelete[n] = 1;
    }
    double vx1;
    while(phi[n]<0){
        phi[n] = phi[n] + PHIp;
        vx1 = vx[n] * cos(PHIp) - vy[n] * sin(PHIp);
        vy[n] = vx[n] * sin(PHIp) + vy[n] * cos(PHIp);
        vx[n] = vx1;
    }
    while(phi[n]>PHIp){
        phi[n] = phi[n] - PHIp;
        vx1 = vx[n] * cos(PHIp) + vy[n] * sin(PHIp);
        vy[n] = -vx[n] * sin(PHIp) + vy[n] * cos(PHIp);
        vx[n] = vx1;
    }
    x[n] = r[n] * cos(phi[n]);
    y[n] = r[n] * sin(phi[n]);
    if (z[n]>Lp){
        z[n] = 2*Lp - z[n];
        vz[n] = -vz[n];
    }
    if (z[n]<0){
        z[n] = -z[n];
        vz[n] = -vz[n];
    }
}


__global__ void processMandelbrotElement( 
                     double * x, 
                     double * y, 
                     double * z, 
                     double * vx,
                     double * vy,
                     double * vz,
                     double * r,
                     double * phi,
                     double * vr,
                     double * Er,
                     double * Ez,
                     double * Hphi,
                     double * charge,
                     double * m,
                     double * E,
                     int *parDelete,
                     const int parNum,
                     const double Rp,
                     const double Lp,
                     const double PHIp,
                     const int gridR,
                     const int gridZ,
                     const double mu,
                     const double c,
                     const double dr,
                     const double dz,
                     const double dt ) {
        EM1(x, y, z, vx, vy, vz, r, phi, vr, Er, Ez, Hphi, charge, m, E, parDelete, parNum, Rp, Lp, PHIp, gridR, gridZ, mu, c, dr, dz, dt);
}

	.file	"tmpxft_002ee94e_00000000-6_EM4.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2011:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2011:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z3EM1PdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.type	_Z3EM1PdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, @function
_Z3EM1PdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd:
.LFB2008:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2008:
	.size	_Z3EM1PdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, .-_Z3EM1PdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.globl	_Z85__device_stub__Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_PiidddiidddddPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.type	_Z85__device_stub__Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_PiidddiidddddPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, @function
_Z85__device_stub__Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_PiidddiidddddPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd:
.LFB2033:
	.cfi_startproc
	endbr64
	subq	$488, %rsp
	.cfi_def_cfa_offset 496
	movq	496(%rsp), %rax
	movq	%rdi, 184(%rsp)
	leaq	208(%rsp), %rdi
	movq	%rsi, 176(%rsp)
	leaq	220(%rsp), %rsi
	movq	%rax, 136(%rsp)
	movq	504(%rsp), %rax
	movq	%rdx, 168(%rsp)
	leaq	192(%rsp), %rdx
	movq	%rax, 128(%rsp)
	movq	512(%rsp), %rax
	movq	%rcx, 160(%rsp)
	leaq	200(%rsp), %rcx
	movq	%rax, 120(%rsp)
	movq	520(%rsp), %rax
	movq	%r8, 152(%rsp)
	movq	%rax, 112(%rsp)
	movq	528(%rsp), %rax
	movq	%r9, 144(%rsp)
	movq	%rax, 104(%rsp)
	movq	536(%rsp), %rax
	movsd	%xmm0, 56(%rsp)
	movq	%rax, 96(%rsp)
	movq	544(%rsp), %rax
	movsd	%xmm1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	552(%rsp), %rax
	movsd	%xmm2, 40(%rsp)
	movq	%rax, 80(%rsp)
	movsd	%xmm3, 32(%rsp)
	movsd	%xmm4, 24(%rsp)
	movsd	%xmm5, 16(%rsp)
	movsd	%xmm6, 8(%rsp)
	movsd	%xmm7, (%rsp)
	movq	560(%rsp), %rax
	movq	%rax, 72(%rsp)
	movq	568(%rsp), %rax
	movq	%rax, 64(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 472(%rsp)
	xorl	%eax, %eax
	leaq	184(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	176(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	168(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	160(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	152(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	144(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	136(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	128(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	120(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 336(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 344(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 352(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 360(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 368(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 376(%rsp)
	leaq	576(%rsp), %rax
	movq	%rax, 384(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 392(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 400(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 408(%rsp)
	leaq	584(%rsp), %rax
	movq	%rax, 416(%rsp)
	leaq	592(%rsp), %rax
	movq	%rax, 424(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 432(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 440(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 448(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 456(%rsp)
	movq	%rsp, %rax
	movq	%rax, 464(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 208(%rsp)
	movl	$1, 216(%rsp)
	movq	%rax, 220(%rsp)
	movl	$1, 228(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L4
	pushq	200(%rsp)
	.cfi_def_cfa_offset 504
	leaq	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd(%rip), %rdi
	pushq	200(%rsp)
	.cfi_def_cfa_offset 512
	movq	236(%rsp), %rcx
	movl	244(%rsp), %r8d
	movq	224(%rsp), %rsi
	movl	232(%rsp), %edx
	leaq	272(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 504
	popq	%rdx
	.cfi_def_cfa_offset 496
.L4:
	movq	472(%rsp), %rax
	subq	%fs:40, %rax
	je	.L6
	call	__stack_chk_fail@PLT
.L6:
	addq	$488, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2033:
	.size	_Z85__device_stub__Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_PiidddiidddddPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, .-_Z85__device_stub__Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_PiidddiidddddPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.globl	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.type	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, @function
_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd:
.LFB2034:
	.cfi_startproc
	endbr64
	jmp	_Z85__device_stub__Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_PiidddiidddddPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.cfi_endproc
.LFE2034:
	.size	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, .-_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2036:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2036:
	.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
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
//nvcc -ptx EM4.cu -ccbin "F:Visual Studio\VC\Tools\MSVC\14.12.25827\bin\Hostx64\x64"
__device__ void EM1( double * x, 
                     double * y, 
                     double * z, 
                     double * vx,
                     double * vy,
                     double * vz,
                     double * r,
                     double * phi,
                     double * vr,
                     double * Er,
                     double * Ez,
                     double * Hphi,
                     double * charge,
                     double * m,
                     double * E,
                     int *parDelete,
                     const int parNum,
                     const double Rp,
                     const double Lp,
                     const double PHIp,
                     const int gridR,
                     const int gridZ,
                     const double mu,
                     const double c,
                     const double dr,
                     const double dz,
                     const double dt ) {
    int globalBlockIndex = blockIdx.x + blockIdx.y * gridDim.x;
    int localThreadIdx = threadIdx.x + blockDim.x * threadIdx.y;
    int threadsPerBlock = blockDim.x*blockDim.y;
    int n = localThreadIdx + globalBlockIndex*threadsPerBlock;
    if ( n >= parNum || parDelete[n]==1){
        return;
    }
    double r0 = r[n];
    double z0 = z[n];
    int ar,az,a1,a2,a3,a4;
    double Er0,Ez0,Hphi0;

    // Er
    ar = floor(r0/dr-0.5);
    az = floor(z0/dz);
    a1 = ar + az * gridR;
    a2 = (ar+1) + az * gridR;
    a3 = ar + (az+1) * gridR;
    a4 = (ar+1) + (az+1) * gridR;
    if( ar<0 ){
        Er0 = ( Er[a2]*(az*dz+dz-z0) + Er[a4]*(z0-az*dz) )/dz;
    }
    else if( ar >= gridR-1 ){
        Er0 = ( Er[a1]*(az*dz+dz-z0) + Er[a3]*(z0-az*dz) )/dz;
    }
    else{
        Er0 = ( Er[a1]*((ar+1.5)*dr-r0)*((az+1)*dz-z0)
                 + Er[a2]*(r0-(ar+0.5)*dr)*((az+1)*dz-z0)
                 + Er[a3]*((ar+1.5)*dr-r0)*(z0-az*dz)
                 + Er[a4]*(r0-(ar+0.5)*dr)*(z0-az*dz) )/(dr*dz);
        
    }
    // Ez
    ar = floor(r0/dr);
    az = floor(z0/dz-0.5);
    a1 = ar + az * (gridR+1);
    a2 = (ar+1) + az * (gridR+1);
    a3 = ar + (az+1) * (gridR+1);
    a4 = (ar+1) + (az+1) * (gridR+1);
    if( az<0 ){
        Ez0 = ( Ez[a3]*(ar*dr+dr-r0) + Ez[a4]*(r0-ar*dr) )/dr;
    }
    else if( az >= gridZ-1 ){
        Ez0 = ( Ez[a1]*(ar*dr+dr-r0) + Ez[a2]*(r0-ar*dr) )/dr;
    }
    else{
        Ez0 = ( Ez[a1]*((ar+1)*dr-r0)*((az+1.5)*dz-z0)
                 + Ez[a2]*(r0-ar*dr)*((az+1.5)*dz-z0)
                 + Ez[a3]*((ar+1)*dr-r0)*(z0-(az+0.5)*dz)
                 + Ez[a4]*(r0-ar*dr)*(z0-(az+0.5)*dz) )/(dr*dz);
        
    }

    // Hphi
    ar = floor(r0/dr-0.5);
    az = floor(z0/dz-0.5);
    a1 = ar + az * gridR;
    a2 = (ar+1) + az * gridR;
    a3 = ar + (az+1) * gridR;
    a4 = (ar+1) + (az+1) * gridR;
    if( ar<0 ){
        if( az<0 ){
            Hphi0 = Hphi[a4];
        }
        else if( az>=gridZ-1 ){
            Hphi0 = Hphi[a2];
        }
        else{
            Hphi0 = ( Hphi[a2]*((az+1.5)*dz-z0) + Hphi[a4]*(z0-(az+0.5)*dz) )/dz;
        }
    }
    else if( ar>=gridR-1 ){
        if( az<0 ){
            Hphi0 = Hphi[a3];
        }
        else if( az>=gridZ-1 ){
            Hphi0 = Hphi[a1];
        }
        else{
            Hphi0 = ( Hphi[a1]*((az+1.5)*dz-z0) + Hphi[a3]*(z0-(az+0.5)*dz) )/dz;
        }
    }
    else if( az<0 ){
        Hphi0 = ( Hphi[a3]*((ar+1.5)*dr-r0) + Hphi[a4]*(r0-(ar+0.5)*dr) )/dr;
    }
    else if( az>=gridZ-1 ){
        Hphi0 = ( Hphi[a1]*((ar+1.5)*dr-r0) + Hphi[a2]*(r0-(ar+0.5)*dr) )/dr;
    }
    else{
        Hphi0 = ( Hphi[a1]*((ar+1.5)*dr-r0)*((az+1.5)*dz-z0)
               + Hphi[a2]*(r0-(ar+0.5)*dr)*((az+1.5)*dz-z0)
               + Hphi[a3]*((ar+1.5)*dr-r0)*(z0-(az+0.5)*dz)
               + Hphi[a4]*(r0-(ar+0.5)*dr)*(z0-(az+0.5)*dz) )/(dr*dz);
    }

    //F
    double Fx,Fy,Fz,Fr;
    Fr = charge[n] * (Er0 + (-vz[n]*Hphi0*mu));
    Fz = charge[n] * (Ez0 + (vr[n]*Hphi0*mu));
    Fx = Fr*cos(phi[n]);
    Fy = Fr*sin(phi[n]);


    //v
    double gamma,ux,uy,uz;
    gamma = 1/sqrt( 1-( vx[n]*vx[n] + vy[n]*vy[n] + vz[n]*vz[n] )/(c*c) );
    ux = gamma * vx[n] + Fx/m[n]*dt;
    uy = gamma * vy[n] + Fy/m[n]*dt;
    uz = gamma * vz[n] + Fz/m[n]*dt;
    gamma = sqrt( 1+ (ux*ux + uy*uy + uz*uz)/(c*c) );

    E[n] = (gamma-1)*m[n]*c*c;

    vx[n] = ux/gamma;
    vy[n] = uy/gamma;
    vz[n] = uz/gamma;

    x[n] = x[n] + 0.5*dt*vx[n];
    y[n] = y[n] + 0.5*dt*vy[n];
    z[n] = z[n] + 0.5*dt*vz[n];

    r[n] = sqrt( x[n]*x[n] + y[n]*y[n] );
    phi[n] = atan( y[n]/x[n] );
    if (r[n] > Rp){
        parDelete[n] = 1;
    }
    double vx1;
    while(phi[n]<0){
        phi[n] = phi[n] + PHIp;
        vx1 = vx[n] * cos(PHIp) - vy[n] * sin(PHIp);
        vy[n] = vx[n] * sin(PHIp) + vy[n] * cos(PHIp);
        vx[n] = vx1;
    }
    while(phi[n]>PHIp){
        phi[n] = phi[n] - PHIp;
        vx1 = vx[n] * cos(PHIp) + vy[n] * sin(PHIp);
        vy[n] = -vx[n] * sin(PHIp) + vy[n] * cos(PHIp);
        vx[n] = vx1;
    }
    x[n] = r[n] * cos(phi[n]);
    y[n] = r[n] * sin(phi[n]);
    if (z[n]>Lp){
        z[n] = 2*Lp - z[n];
        vz[n] = -vz[n];
    }
    if (z[n]<0){
        z[n] = -z[n];
        vz[n] = -vz[n];
    }
}


__global__ void processMandelbrotElement( 
                     double * x, 
                     double * y, 
                     double * z, 
                     double * vx,
                     double * vy,
                     double * vz,
                     double * r,
                     double * phi,
                     double * vr,
                     double * Er,
                     double * Ez,
                     double * Hphi,
                     double * charge,
                     double * m,
                     double * E,
                     int *parDelete,
                     const int parNum,
                     const double Rp,
                     const double Lp,
                     const double PHIp,
                     const int gridR,
                     const int gridZ,
                     const double mu,
                     const double c,
                     const double dr,
                     const double dz,
                     const double dt ) {
        EM1(x, y, z, vx, vy, vz, r, phi, vr, Er, Ez, Hphi, charge, m, E, parDelete, parNum, Rp, Lp, PHIp, gridR, gridZ, mu, c, dr, dz, dt);
}

	.text
	.file	"EM4.hip"
	.globl	_Z39__device_stub__processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd # -- Begin function _Z39__device_stub__processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.type	_Z39__device_stub__processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd,@function
_Z39__device_stub__processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd: # @_Z39__device_stub__processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.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	$376, %rsp                      # imm = 0x178
	.cfi_def_cfa_offset 432
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	leaq	120(%rsp), %r11
	movq	%rdi, (%r11)
	leaq	112(%rsp), %r14
	movq	%rsi, (%r14)
	leaq	104(%rsp), %r15
	movq	%rdx, (%r15)
	leaq	96(%rsp), %r12
	movq	%rcx, (%r12)
	leaq	88(%rsp), %r13
	movq	%r8, (%r13)
	leaq	80(%rsp), %rbp
	movq	%r9, (%rbp)
	leaq	72(%rsp), %rcx
	movsd	%xmm0, (%rcx)
	leaq	64(%rsp), %rsi
	movsd	%xmm1, (%rsi)
	leaq	56(%rsp), %r8
	movsd	%xmm2, (%r8)
	leaq	48(%rsp), %rax
	movsd	%xmm3, (%rax)
	leaq	40(%rsp), %rdx
	movsd	%xmm4, (%rdx)
	leaq	32(%rsp), %rdi
	movsd	%xmm5, (%rdi)
	leaq	24(%rsp), %r9
	movsd	%xmm6, (%r9)
	leaq	16(%rsp), %r10
	movsd	%xmm7, (%r10)
	leaq	160(%rsp), %rbx
	movq	%r11, (%rbx)
	movq	%r14, 8(%rbx)
	movq	%r15, 16(%rbx)
	movq	%r12, 24(%rbx)
	movq	%r13, 32(%rbx)
	movq	%rbp, 40(%rbx)
	leaq	432(%rsp), %r11
	movq	%r11, 48(%rbx)
	leaq	440(%rsp), %r11
	movq	%r11, 56(%rbx)
	leaq	448(%rsp), %r11
	movq	%r11, 64(%rbx)
	leaq	456(%rsp), %r11
	movq	%r11, 72(%rbx)
	leaq	464(%rsp), %r11
	movq	%r11, 80(%rbx)
	leaq	472(%rsp), %r11
	movq	%r11, 88(%rbx)
	leaq	480(%rsp), %r11
	movq	%r11, 96(%rbx)
	leaq	488(%rsp), %r11
	movq	%r11, 104(%rbx)
	leaq	496(%rsp), %r11
	movq	%r11, 112(%rbx)
	leaq	504(%rsp), %r11
	movq	%r11, 120(%rbx)
	leaq	512(%rsp), %r11
	movq	%r11, 128(%rbx)
	movq	%rcx, 136(%rbx)
	movq	%rsi, 144(%rbx)
	movq	%r8, 152(%rbx)
	leaq	520(%rsp), %rcx
	movq	%rcx, 160(%rbx)
	leaq	528(%rsp), %rcx
	movq	%rcx, 168(%rbx)
	movq	%rax, 176(%rbx)
	movq	%rdx, 184(%rbx)
	movq	%rdi, 192(%rbx)
	movq	%r9, 200(%rbx)
	movq	%r10, 208(%rbx)
	leaq	144(%rsp), %r14
	leaq	128(%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	$_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$392, %rsp                      # imm = 0x188
	.cfi_adjust_cfa_offset -392
	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_end0:
	.size	_Z39__device_stub__processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, .Lfunc_end0-_Z39__device_stub__processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB1_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB1_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end1:
	.size	__hip_module_ctor, .Lfunc_end1-__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	.LBB2_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
.LBB2_2:
	retq
.Lfunc_end2:
	.size	__hip_module_dtor, .Lfunc_end2-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd,@object # @_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.section	.rodata,"a",@progbits
	.globl	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.p2align	3, 0x0
_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd:
	.quad	_Z39__device_stub__processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.size	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, 8

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd"
	.size	.L__unnamed_1, 72

	.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 _Z39__device_stub__processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd ; -- Begin function _Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.globl	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
	.p2align	8
	.type	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd,@function
_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd: ; @_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
; %bb.0:
	s_load_b32 s4, s[0:1], 0xd0
	s_add_u32 s2, s0, 0xd0
	s_addc_u32 s3, s1, 0
	s_load_b32 s5, s[0:1], 0xdc
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s6, s4, s15
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_add_i32 s6, s6, s14
	s_cmp_lt_u32 s14, s4
	s_cselect_b32 s4, 12, 18
	v_mov_b32_e32 v1, s4
	global_load_u16 v4, v1, s[2:3]
	v_bfe_u32 v1, v0, 10, 10
	s_load_b32 s2, s[0:1], 0x80
	s_lshr_b32 s3, s5, 16
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s6, s3, v[1:2]
	v_and_b32_e32 v3, 0x3ff, v0
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[0:1], null, v2, v4, v[3:4]
	s_waitcnt lgkmcnt(0)
	v_cmp_gt_i32_e32 vcc_lo, s2, v0
	s_and_saveexec_b32 s2, vcc_lo
	s_cbranch_execz .LBB0_76
; %bb.1:
	s_load_b512 s[4:19], s[0:1], 0x40
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v6, vcc_lo, s18, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v7, vcc_lo, s19, v3, vcc_lo
	global_load_b32 v2, v[6:7], off
	s_waitcnt vmcnt(0)
	v_cmp_ne_u32_e32 vcc_lo, 1, v2
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_76
; %bb.2:
	s_load_b512 s[36:51], s[0:1], 0x0
	v_lshlrev_b64 v[10:11], 3, v[0:1]
	s_clause 0x1
	s_load_b256 s[20:27], s[0:1], 0xa8
	s_load_b64 s[18:19], s[0:1], 0xa0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_u32 v2, vcc_lo, s48, v10
	v_add_co_ci_u32_e32 v3, vcc_lo, s49, v11, vcc_lo
	v_add_co_u32 v0, vcc_lo, s40, v10
	v_add_co_ci_u32_e32 v1, vcc_lo, s41, v11, vcc_lo
	global_load_b64 v[8:9], v[2:3], off
	global_load_b64 v[4:5], v[0:1], off
	s_waitcnt vmcnt(1)
	v_div_scale_f64 v[12:13], null, s[24:25], s[24:25], v[8:9]
	s_waitcnt vmcnt(0)
	v_div_scale_f64 v[14:15], null, s[26:27], s[26:27], v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[16:17], v[12:13]
	v_rcp_f64_e32 v[18:19], v[14:15]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[20:21], -v[12:13], v[16:17], 1.0
	v_fma_f64 v[22:23], -v[14:15], v[18:19], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[16:17], v[16:17], v[20:21], v[16:17]
	v_fma_f64 v[18:19], v[18:19], v[22:23], v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[20:21], -v[12:13], v[16:17], 1.0
	v_fma_f64 v[22:23], -v[14:15], v[18:19], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[16:17], v[16:17], v[20:21], v[16:17]
	v_div_scale_f64 v[20:21], vcc_lo, v[8:9], s[24:25], v[8:9]
	v_fma_f64 v[18:19], v[18:19], v[22:23], v[18:19]
	v_div_scale_f64 v[22:23], s2, v[4:5], s[26:27], v[4:5]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[24:25], v[20:21], v[16:17]
	v_mul_f64 v[26:27], v[22:23], v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[12:13], -v[12:13], v[24:25], v[20:21]
	v_fma_f64 v[14:15], -v[14:15], v[26:27], v[22:23]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_fmas_f64 v[12:13], v[12:13], v[16:17], v[24:25]
	s_mov_b32 vcc_lo, s2
	v_div_fmas_f64 v[16:17], v[14:15], v[18:19], v[26:27]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_fixup_f64 v[14:15], v[12:13], s[24:25], v[8:9]
	v_div_fixup_f64 v[17:18], v[16:17], s[26:27], v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[12:13], v[14:15], -0.5
	v_floor_f64_e32 v[19:20], v[17:18]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_floor_f64_e32 v[12:13], v[12:13]
	v_cvt_i32_f64_e32 v25, v[19:20]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_i32_f64_e32 v28, v[12:13]
                                        ; implicit-def: $vgpr12_vgpr13
	v_mul_lo_u32 v19, v25, s18
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v26, s18, v19
	v_cmp_lt_i32_e64 s3, -1, v28
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_saveexec_b32 s2, s3
	s_xor_b32 s2, exec_lo, s2
	s_cbranch_execz .LBB0_8
; %bb.3:
	v_add_nc_u32_e32 v12, v19, v28
	v_add_nc_u32_e32 v21, v26, v28
	s_add_i32 s28, s18, -1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v13, 31, v12
	v_ashrrev_i32_e32 v22, 31, v21
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[12:13], 3, v[12:13]
	v_add_co_u32 v23, vcc_lo, s6, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	v_add_co_ci_u32_e32 v24, vcc_lo, s7, v13, vcc_lo
	v_cmp_le_i32_e32 vcc_lo, s28, v28
                                        ; implicit-def: $vgpr12_vgpr13
	global_load_b64 v[19:20], v[23:24], off
	s_and_saveexec_b32 s28, vcc_lo
	s_xor_b32 s28, exec_lo, s28
	s_cbranch_execz .LBB0_5
; %bb.4:
	v_lshlrev_b64 v[12:13], 3, v[21:22]
	v_cvt_f64_i32_e32 v[21:22], v25
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v12, vcc_lo, s6, v12
	v_add_co_ci_u32_e32 v13, vcc_lo, s7, v13, vcc_lo
	global_load_b64 v[12:13], v[12:13], off
	v_fma_f64 v[23:24], v[21:22], s[26:27], s[26:27]
	v_fma_f64 v[21:22], -v[21:22], s[26:27], v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[23:24], v[23:24], -v[4:5]
	s_waitcnt vmcnt(0)
	v_mul_f64 v[12:13], v[21:22], v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[12:13], v[23:24], v[19:20], v[12:13]
	v_div_scale_f64 v[19:20], null, s[26:27], s[26:27], v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[21:22], v[19:20]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[23:24], -v[19:20], v[21:22], 1.0
	v_fma_f64 v[21:22], v[21:22], v[23:24], v[21:22]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[23:24], -v[19:20], v[21:22], 1.0
	v_fma_f64 v[21:22], v[21:22], v[23:24], v[21:22]
	v_div_scale_f64 v[23:24], vcc_lo, v[12:13], s[26:27], v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[25:26], v[23:24], v[21:22]
	v_fma_f64 v[19:20], -v[19:20], v[25:26], v[23:24]
                                        ; implicit-def: $vgpr23_vgpr24
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[19:20], v[19:20], v[21:22], v[25:26]
                                        ; implicit-def: $vgpr25
                                        ; implicit-def: $vgpr21
                                        ; implicit-def: $vgpr26
	v_div_fixup_f64 v[12:13], v[19:20], s[26:27], v[12:13]
                                        ; implicit-def: $vgpr19_vgpr20
.LBB0_5:                                ; %Flow153
	s_and_not1_saveexec_b32 s28, s28
	s_cbranch_execz .LBB0_7
; %bb.6:
	global_load_b64 v[12:13], v[23:24], off offset:8
	v_lshlrev_b64 v[21:22], 3, v[21:22]
	v_ashrrev_i32_e32 v16, 31, v26
	v_cvt_f64_i32_e32 v[33:34], v25
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v21, vcc_lo, s6, v21
	v_add_co_ci_u32_e32 v22, vcc_lo, s7, v22, vcc_lo
	global_load_b64 v[21:22], v[21:22], off
	v_ashrrev_i32_e32 v24, 31, v28
	v_add_co_u32 v23, vcc_lo, v26, v28
	v_cvt_f64_i32_e32 v[26:27], v28
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v24, vcc_lo, v16, v24, vcc_lo
	v_add_nc_u32_e32 v16, 1, v25
	v_lshlrev_b64 v[23:24], 3, v[23:24]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_f64_i32_e32 v[31:32], v16
	v_add_co_u32 v23, vcc_lo, s6, v23
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v24, vcc_lo, s7, v24, vcc_lo
	global_load_b64 v[23:24], v[23:24], off offset:8
	v_fma_f64 v[33:34], -v[33:34], s[26:27], v[4:5]
	v_add_f64 v[29:30], v[26:27], 0.5
	v_add_f64 v[26:27], 0x3ff80000, v[26:27]
	v_fma_f64 v[31:32], v[31:32], s[26:27], -v[4:5]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[29:30], -v[29:30], s[24:25], v[8:9]
	v_fma_f64 v[26:27], v[26:27], s[24:25], -v[8:9]
	s_waitcnt vmcnt(3)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_mul_f64 v[19:20], v[26:27], v[19:20]
	s_waitcnt vmcnt(2)
	v_mul_f64 v[12:13], v[29:30], v[12:13]
	s_waitcnt vmcnt(1)
	v_mul_f64 v[21:22], v[26:27], v[21:22]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[12:13], v[31:32], v[12:13]
	v_fma_f64 v[12:13], v[31:32], v[19:20], v[12:13]
	s_waitcnt vmcnt(0)
	v_mul_f64 v[19:20], v[29:30], v[23:24]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[12:13], v[33:34], v[21:22], v[12:13]
	v_mul_f64 v[21:22], s[24:25], s[26:27]
	v_fma_f64 v[12:13], v[33:34], v[19:20], v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_scale_f64 v[19:20], null, v[21:22], v[21:22], v[12:13]
	v_rcp_f64_e32 v[23:24], v[19:20]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[25:26], -v[19:20], v[23:24], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[23:24], v[23:24], v[25:26], v[23:24]
	v_fma_f64 v[25:26], -v[19:20], v[23:24], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[23:24], v[23:24], v[25:26], v[23:24]
	v_div_scale_f64 v[25:26], vcc_lo, v[12:13], v[21:22], v[12:13]
	v_mul_f64 v[29:30], v[25:26], v[23:24]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[19:20], -v[19:20], v[29:30], v[25:26]
	v_div_fmas_f64 v[19:20], v[19:20], v[23:24], v[29:30]
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fixup_f64 v[12:13], v[19:20], v[21:22], v[12:13]
.LBB0_7:                                ; %Flow154
	s_or_b32 exec_lo, exec_lo, s28
                                        ; implicit-def: $vgpr19
                                        ; implicit-def: $vgpr25
                                        ; implicit-def: $vgpr26
.LBB0_8:                                ; %Flow155
	s_or_saveexec_b32 s2, s2
	v_ashrrev_i32_e32 v16, 31, v28
	s_xor_b32 exec_lo, exec_lo, s2
	s_cbranch_execz .LBB0_10
; %bb.9:
	v_ashrrev_i32_e32 v13, 31, v26
	v_add_co_u32 v12, vcc_lo, v26, v28
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v20, 31, v19
	v_cvt_f64_i32_e32 v[21:22], v25
	v_add_co_ci_u32_e32 v13, vcc_lo, v13, v16, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[12:13], 3, v[12:13]
	v_add_co_u32 v12, vcc_lo, s6, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v13, vcc_lo, s7, v13, vcc_lo
	v_add_co_u32 v19, vcc_lo, v28, v19
	v_add_co_ci_u32_e32 v20, vcc_lo, v16, v20, vcc_lo
	global_load_b64 v[12:13], v[12:13], off offset:8
	v_lshlrev_b64 v[19:20], 3, v[19:20]
	v_add_co_u32 v19, vcc_lo, s6, v19
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v20, vcc_lo, s7, v20, vcc_lo
	global_load_b64 v[19:20], v[19:20], off offset:8
	v_fma_f64 v[23:24], v[21:22], s[26:27], s[26:27]
	v_fma_f64 v[21:22], -v[21:22], s[26:27], v[4:5]
	v_add_f64 v[23:24], v[23:24], -v[4:5]
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_f64 v[12:13], v[21:22], v[12:13]
	s_waitcnt vmcnt(0)
	v_fma_f64 v[12:13], v[23:24], v[19:20], v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_scale_f64 v[19:20], null, s[26:27], s[26:27], v[12:13]
	v_rcp_f64_e32 v[21:22], v[19:20]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[23:24], -v[19:20], v[21:22], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[21:22], v[21:22], v[23:24], v[21:22]
	v_fma_f64 v[23:24], -v[19:20], v[21:22], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[21:22], v[21:22], v[23:24], v[21:22]
	v_div_scale_f64 v[23:24], vcc_lo, v[12:13], s[26:27], v[12:13]
	v_mul_f64 v[25:26], v[23:24], v[21:22]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[19:20], -v[19:20], v[25:26], v[23:24]
	v_div_fmas_f64 v[19:20], v[19:20], v[21:22], v[25:26]
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fixup_f64 v[12:13], v[19:20], s[26:27], v[12:13]
.LBB0_10:
	s_or_b32 exec_lo, exec_lo, s2
	v_add_f64 v[17:18], v[17:18], -0.5
	v_floor_f64_e32 v[14:15], v[14:15]
	s_add_i32 s2, s18, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_floor_f64_e32 v[17:18], v[17:18]
	v_cvt_i32_f64_e32 v23, v[14:15]
                                        ; implicit-def: $vgpr14_vgpr15
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_i32_f64_e32 v29, v[17:18]
	v_mul_lo_u32 v18, s2, v29
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v24, s2, v18
	v_cmp_lt_i32_e64 s2, -1, v29
	v_add_nc_u32_e32 v17, v24, v23
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_saveexec_b32 s6, s2
	s_xor_b32 s6, exec_lo, s6
	s_cbranch_execz .LBB0_16
; %bb.11:
	v_add_nc_u32_e32 v14, v18, v23
	s_add_i32 s7, s19, -1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v15, 31, v14
	v_lshlrev_b64 v[14:15], 3, v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v21, vcc_lo, s8, v14
	v_add_co_ci_u32_e32 v22, vcc_lo, s9, v15, vcc_lo
	v_cmp_le_i32_e32 vcc_lo, s7, v29
                                        ; implicit-def: $vgpr14_vgpr15
	global_load_b64 v[19:20], v[21:22], off
	s_and_saveexec_b32 s7, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s7, exec_lo, s7
	s_cbranch_execz .LBB0_13
; %bb.12:
	global_load_b64 v[14:15], v[21:22], off offset:8
	v_cvt_f64_i32_e32 v[17:18], v23
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[21:22], v[17:18], s[24:25], s[24:25]
	v_fma_f64 v[17:18], -v[17:18], s[24:25], v[8:9]
	v_add_f64 v[21:22], v[21:22], -v[8:9]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[14:15], v[17:18], v[14:15]
	v_fma_f64 v[14:15], v[21:22], v[19:20], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_scale_f64 v[17:18], null, s[24:25], s[24:25], v[14:15]
	v_rcp_f64_e32 v[19:20], v[17:18]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[21:22], -v[17:18], v[19:20], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[19:20], v[19:20], v[21:22], v[19:20]
	v_fma_f64 v[21:22], -v[17:18], v[19:20], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[19:20], v[19:20], v[21:22], v[19:20]
	v_div_scale_f64 v[21:22], vcc_lo, v[14:15], s[24:25], v[14:15]
	v_mul_f64 v[23:24], v[21:22], v[19:20]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[17:18], -v[17:18], v[23:24], v[21:22]
                                        ; implicit-def: $vgpr21_vgpr22
	v_div_fmas_f64 v[17:18], v[17:18], v[19:20], v[23:24]
                                        ; implicit-def: $vgpr23
                                        ; implicit-def: $vgpr19_vgpr20
                                        ; implicit-def: $vgpr24
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fixup_f64 v[14:15], v[17:18], s[24:25], v[14:15]
                                        ; implicit-def: $vgpr17
.LBB0_13:                               ; %Flow150
	s_and_not1_saveexec_b32 s7, s7
	s_cbranch_execz .LBB0_15
; %bb.14:
	global_load_b64 v[14:15], v[21:22], off offset:8
	v_ashrrev_i32_e32 v18, 31, v17
	v_ashrrev_i32_e32 v25, 31, v23
	v_cvt_f64_i32_e32 v[26:27], v23
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[17:18], 3, v[17:18]
	v_add_co_u32 v17, vcc_lo, s8, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v18, vcc_lo, s9, v18, vcc_lo
	global_load_b64 v[17:18], v[17:18], off
	v_ashrrev_i32_e32 v22, 31, v24
	v_add_co_u32 v21, vcc_lo, v24, v23
	v_add_nc_u32_e32 v23, 1, v23
	v_add_co_ci_u32_e32 v22, vcc_lo, v22, v25, vcc_lo
	v_cvt_f64_i32_e32 v[24:25], v29
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_f64_i32_e32 v[30:31], v23
	v_lshlrev_b64 v[21:22], 3, v[21:22]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v21, vcc_lo, s8, v21
	v_add_co_ci_u32_e32 v22, vcc_lo, s9, v22, vcc_lo
	v_fma_f64 v[26:27], -v[26:27], s[24:25], v[8:9]
	global_load_b64 v[21:22], v[21:22], off offset:8
	v_add_f64 v[32:33], 0x3ff80000, v[24:25]
	v_fma_f64 v[30:31], v[30:31], s[24:25], -v[8:9]
	v_add_f64 v[23:24], v[24:25], 0.5
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[32:33], v[32:33], s[26:27], -v[4:5]
	s_waitcnt vmcnt(3)
	v_mul_f64 v[19:20], v[30:31], v[19:20]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_fma_f64 v[23:24], -v[23:24], s[26:27], v[4:5]
	s_waitcnt vmcnt(2)
	v_mul_f64 v[14:15], v[26:27], v[14:15]
	s_waitcnt vmcnt(1)
	v_mul_f64 v[17:18], v[30:31], v[17:18]
	v_mul_f64 v[14:15], v[32:33], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f64 v[14:15], v[32:33], v[19:20], v[14:15]
	s_waitcnt vmcnt(0)
	v_mul_f64 v[19:20], v[26:27], v[21:22]
	v_fma_f64 v[14:15], v[23:24], v[17:18], v[14:15]
	v_mul_f64 v[17:18], s[24:25], s[26:27]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[14:15], v[23:24], v[19:20], v[14:15]
	v_div_scale_f64 v[19:20], null, v[17:18], v[17:18], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[21:22], v[19:20]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[23:24], -v[19:20], v[21:22], 1.0
	v_fma_f64 v[21:22], v[21:22], v[23:24], v[21:22]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[23:24], -v[19:20], v[21:22], 1.0
	v_fma_f64 v[21:22], v[21:22], v[23:24], v[21:22]
	v_div_scale_f64 v[23:24], vcc_lo, v[14:15], v[17:18], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[25:26], v[23:24], v[21:22]
	v_fma_f64 v[19:20], -v[19:20], v[25:26], v[23:24]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[19:20], v[19:20], v[21:22], v[25:26]
	v_div_fixup_f64 v[14:15], v[19:20], v[17:18], v[14:15]
.LBB0_15:                               ; %Flow151
	s_or_b32 exec_lo, exec_lo, s7
                                        ; implicit-def: $vgpr17
                                        ; implicit-def: $vgpr23
                                        ; implicit-def: $vgpr24
.LBB0_16:                               ; %Flow152
	s_and_not1_saveexec_b32 s6, s6
	s_cbranch_execz .LBB0_18
; %bb.17:
	v_ashrrev_i32_e32 v15, 31, v24
	v_ashrrev_i32_e32 v18, 31, v23
	v_add_co_u32 v14, vcc_lo, v24, v23
	s_waitcnt vmcnt(0)
	v_cvt_f64_i32_e32 v[19:20], v23
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v15, vcc_lo, v15, v18, vcc_lo
	v_ashrrev_i32_e32 v18, 31, v17
	v_lshlrev_b64 v[14:15], 3, v[14:15]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[17:18], 3, v[17:18]
	v_add_co_u32 v14, vcc_lo, s8, v14
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v15, vcc_lo, s9, v15, vcc_lo
	v_add_co_u32 v17, vcc_lo, s8, v17
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v18, vcc_lo, s9, v18, vcc_lo
	s_clause 0x1
	global_load_b64 v[14:15], v[14:15], off offset:8
	global_load_b64 v[17:18], v[17:18], off
	v_fma_f64 v[21:22], v[19:20], s[24:25], s[24:25]
	v_fma_f64 v[19:20], -v[19:20], s[24:25], v[8:9]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[21:22], v[21:22], -v[8:9]
	s_waitcnt vmcnt(1)
	v_mul_f64 v[14:15], v[19:20], v[14:15]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[14:15], v[21:22], v[17:18], v[14:15]
	v_div_scale_f64 v[17:18], null, s[24:25], s[24:25], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[19:20], v[17:18]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[21:22], -v[17:18], v[19:20], 1.0
	v_fma_f64 v[19:20], v[19:20], v[21:22], v[19:20]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[21:22], -v[17:18], v[19:20], 1.0
	v_fma_f64 v[19:20], v[19:20], v[21:22], v[19:20]
	v_div_scale_f64 v[21:22], vcc_lo, v[14:15], s[24:25], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[23:24], v[21:22], v[19:20]
	v_fma_f64 v[17:18], -v[17:18], v[23:24], v[21:22]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[17:18], v[17:18], v[19:20], v[23:24]
	v_div_fixup_f64 v[14:15], v[17:18], s[24:25], v[14:15]
.LBB0_18:
	s_or_b32 exec_lo, exec_lo, s6
	v_mul_lo_u32 v31, v29, s18
                                        ; implicit-def: $vgpr20_vgpr21
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_nc_u32_e32 v30, s18, v31
	s_and_saveexec_b32 s6, s3
	s_xor_b32 s3, exec_lo, s6
	s_cbranch_execz .LBB0_38
; %bb.19:
	s_add_i32 s6, s18, -1
	v_add_nc_u32_e32 v18, v31, v28
	v_add_nc_u32_e32 v17, v30, v28
	v_cmp_le_i32_e32 vcc_lo, s6, v28
                                        ; implicit-def: $vgpr20_vgpr21
	s_and_saveexec_b32 s6, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s6, exec_lo, s6
	s_cbranch_execz .LBB0_27
; %bb.20:
                                        ; implicit-def: $vgpr20_vgpr21
	s_and_saveexec_b32 s7, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s7, exec_lo, s7
	s_cbranch_execz .LBB0_24
; %bb.21:
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v19, 31, v18
	s_add_i32 s8, s19, -1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[8:9], 3, v[18:19]
	v_add_co_u32 v8, vcc_lo, s10, v8
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v9, vcc_lo, s11, v9, vcc_lo
	v_cmp_gt_i32_e32 vcc_lo, s8, v29
	global_load_b64 v[20:21], v[8:9], off
	s_and_saveexec_b32 s8, vcc_lo
	s_cbranch_execz .LBB0_23
; %bb.22:
	v_ashrrev_i32_e32 v18, 31, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[8:9], 3, v[17:18]
	v_cvt_f64_i32_e32 v[16:17], v29
	v_add_co_u32 v8, vcc_lo, s10, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v9, vcc_lo, s11, v9, vcc_lo
	global_load_b64 v[8:9], v[8:9], off
	v_add_f64 v[18:19], v[16:17], 0.5
	v_add_f64 v[16:17], 0x3ff80000, v[16:17]
	v_fma_f64 v[18:19], -v[18:19], s[26:27], v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[4:5], v[16:17], s[26:27], -v[4:5]
	s_waitcnt vmcnt(0)
	v_mul_f64 v[8:9], v[18:19], v[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[4:5], v[4:5], v[20:21], v[8:9]
	v_div_scale_f64 v[8:9], null, s[26:27], s[26:27], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[16:17], v[8:9]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[18:19], -v[8:9], v[16:17], 1.0
	v_fma_f64 v[16:17], v[16:17], v[18:19], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[18:19], -v[8:9], v[16:17], 1.0
	v_fma_f64 v[16:17], v[16:17], v[18:19], v[16:17]
	v_div_scale_f64 v[18:19], vcc_lo, v[4:5], s[26:27], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[20:21], v[18:19], v[16:17]
	v_fma_f64 v[8:9], -v[8:9], v[20:21], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[8:9], v[8:9], v[16:17], v[20:21]
	v_div_fixup_f64 v[20:21], v[8:9], s[26:27], v[4:5]
.LBB0_23:                               ; %Flow136
	s_or_b32 exec_lo, exec_lo, s8
                                        ; implicit-def: $vgpr17
.LBB0_24:                               ; %Flow137
	s_and_not1_saveexec_b32 s7, s7
	s_cbranch_execz .LBB0_26
; %bb.25:
	v_ashrrev_i32_e32 v18, 31, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[4:5], 3, v[17:18]
	v_add_co_u32 v4, vcc_lo, s10, v4
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v5, vcc_lo
	global_load_b64 v[20:21], v[4:5], off
.LBB0_26:                               ; %Flow138
	s_or_b32 exec_lo, exec_lo, s7
                                        ; implicit-def: $vgpr8_vgpr9
                                        ; implicit-def: $vgpr4_vgpr5
                                        ; implicit-def: $vgpr28
                                        ; implicit-def: $vgpr29
                                        ; implicit-def: $vgpr31
                                        ; implicit-def: $vgpr30
                                        ; implicit-def: $vgpr18
                                        ; implicit-def: $vgpr17
                                        ; implicit-def: $vgpr15_vgpr16
.LBB0_27:                               ; %Flow143
	s_and_not1_saveexec_b32 s6, s6
	s_cbranch_execz .LBB0_37
; %bb.28:
	v_cvt_f64_i32_e32 v[22:23], v28
                                        ; implicit-def: $vgpr20_vgpr21
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_f64 v[24:25], 0x3ff80000, v[22:23]
	s_and_saveexec_b32 s7, s2
	s_xor_b32 s7, exec_lo, s7
	s_cbranch_execz .LBB0_34
; %bb.29:
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v19, 31, v18
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[26:27], v[24:25], s[24:25], -v[8:9]
	s_add_i32 s8, s19, -1
                                        ; implicit-def: $vgpr20_vgpr21
	v_lshlrev_b64 v[18:19], 3, v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v18, vcc_lo, s10, v18
	v_add_co_ci_u32_e32 v19, vcc_lo, s11, v19, vcc_lo
	v_cmp_le_i32_e32 vcc_lo, s8, v29
	global_load_b64 v[18:19], v[18:19], off
	s_waitcnt vmcnt(0)
	v_mul_f64 v[24:25], v[26:27], v[18:19]
	s_and_saveexec_b32 s8, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s8, exec_lo, s8
	s_cbranch_execz .LBB0_31
; %bb.30:
	v_ashrrev_i32_e32 v5, 31, v31
	v_add_co_u32 v4, vcc_lo, v31, v28
                                        ; implicit-def: $vgpr26_vgpr27
                                        ; implicit-def: $vgpr28
                                        ; implicit-def: $vgpr29
                                        ; implicit-def: $vgpr31
                                        ; implicit-def: $vgpr30
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, v5, v16, vcc_lo
	v_add_f64 v[16:17], v[22:23], 0.5
                                        ; implicit-def: $vgpr22_vgpr23
	v_lshlrev_b64 v[4:5], 3, v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, s10, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v5, vcc_lo
	global_load_b64 v[4:5], v[4:5], off offset:8
	v_fma_f64 v[8:9], -v[16:17], s[24:25], v[8:9]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[4:5], v[8:9], v[4:5], v[24:25]
                                        ; implicit-def: $vgpr24_vgpr25
	v_div_scale_f64 v[8:9], null, s[24:25], s[24:25], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[16:17], v[8:9]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[18:19], -v[8:9], v[16:17], 1.0
	v_fma_f64 v[16:17], v[16:17], v[18:19], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[18:19], -v[8:9], v[16:17], 1.0
	v_fma_f64 v[16:17], v[16:17], v[18:19], v[16:17]
	v_div_scale_f64 v[18:19], vcc_lo, v[4:5], s[24:25], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[20:21], v[18:19], v[16:17]
	v_fma_f64 v[8:9], -v[8:9], v[20:21], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[8:9], v[8:9], v[16:17], v[20:21]
                                        ; implicit-def: $vgpr17
                                        ; implicit-def: $vgpr15_vgpr16
	v_div_fixup_f64 v[20:21], v[8:9], s[24:25], v[4:5]
                                        ; implicit-def: $vgpr8_vgpr9
                                        ; implicit-def: $vgpr4_vgpr5
.LBB0_31:                               ; %Flow139
	s_and_not1_saveexec_b32 s8, s8
	s_cbranch_execz .LBB0_33
; %bb.32:
	v_ashrrev_i32_e32 v19, 31, v31
	v_add_co_u32 v18, vcc_lo, v31, v28
	v_ashrrev_i32_e32 v21, 31, v30
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v19, vcc_lo, v19, v16, vcc_lo
	v_lshlrev_b64 v[18:19], 3, v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v18, vcc_lo, s10, v18
	v_add_co_ci_u32_e32 v19, vcc_lo, s11, v19, vcc_lo
	global_load_b64 v[19:20], v[18:19], off offset:8
	v_ashrrev_i32_e32 v18, 31, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[17:18], 3, v[17:18]
	v_add_co_u32 v17, vcc_lo, s10, v17
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v18, vcc_lo, s11, v18, vcc_lo
	v_add_co_u32 v30, vcc_lo, v30, v28
	v_add_co_ci_u32_e32 v31, vcc_lo, v21, v16, vcc_lo
	global_load_b64 v[17:18], v[17:18], off
	v_cvt_f64_i32_e32 v[28:29], v29
	v_add_f64 v[21:22], v[22:23], 0.5
	v_lshlrev_b64 v[30:31], 3, v[30:31]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v30, vcc_lo, s10, v30
	v_add_co_ci_u32_e32 v31, vcc_lo, s11, v31, vcc_lo
	global_load_b64 v[30:31], v[30:31], off offset:8
	v_add_f64 v[32:33], 0x3ff80000, v[28:29]
	v_fma_f64 v[8:9], -v[21:22], s[24:25], v[8:9]
	v_add_f64 v[28:29], v[28:29], 0.5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[21:22], v[32:33], s[26:27], -v[4:5]
	v_fma_f64 v[4:5], -v[28:29], s[26:27], v[4:5]
	s_waitcnt vmcnt(2)
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_3)
	v_mul_f64 v[19:20], v[8:9], v[19:20]
	s_waitcnt vmcnt(1)
	v_mul_f64 v[16:17], v[26:27], v[17:18]
	s_waitcnt vmcnt(0)
	v_mul_f64 v[8:9], v[8:9], v[30:31]
	v_mul_f64 v[19:20], v[21:22], v[19:20]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[18:19], v[21:22], v[24:25], v[19:20]
	v_fma_f64 v[16:17], v[4:5], v[16:17], v[18:19]
	v_mul_f64 v[18:19], s[24:25], s[26:27]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[4:5], v[4:5], v[8:9], v[16:17]
	v_div_scale_f64 v[8:9], null, v[18:19], v[18:19], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[16:17], v[8:9]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[20:21], -v[8:9], v[16:17], 1.0
	v_fma_f64 v[16:17], v[16:17], v[20:21], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], -v[8:9], v[16:17], 1.0
	v_fma_f64 v[16:17], v[16:17], v[20:21], v[16:17]
	v_div_scale_f64 v[20:21], vcc_lo, v[4:5], v[18:19], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[22:23], v[20:21], v[16:17]
	v_fma_f64 v[8:9], -v[8:9], v[22:23], v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[8:9], v[8:9], v[16:17], v[22:23]
	v_div_fixup_f64 v[20:21], v[8:9], v[18:19], v[4:5]
.LBB0_33:                               ; %Flow140
	s_or_b32 exec_lo, exec_lo, s8
                                        ; implicit-def: $vgpr17
                                        ; implicit-def: $vgpr24_vgpr25
                                        ; implicit-def: $vgpr8_vgpr9
                                        ; implicit-def: $vgpr30
                                        ; implicit-def: $vgpr28
                                        ; implicit-def: $vgpr22_vgpr23
                                        ; implicit-def: $vgpr15_vgpr16
.LBB0_34:                               ; %Flow141
	s_and_not1_saveexec_b32 s7, s7
	s_cbranch_execz .LBB0_36
; %bb.35:
	v_ashrrev_i32_e32 v5, 31, v30
	v_add_co_u32 v4, vcc_lo, v30, v28
	v_ashrrev_i32_e32 v18, 31, v17
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, v5, v16, vcc_lo
	v_lshlrev_b64 v[16:17], 3, v[17:18]
	s_waitcnt vmcnt(0)
	v_add_f64 v[18:19], v[22:23], 0.5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[4:5], 3, v[4:5]
	v_add_co_u32 v4, vcc_lo, s10, v4
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v5, vcc_lo
	v_add_co_u32 v16, vcc_lo, s10, v16
	v_add_co_ci_u32_e32 v17, vcc_lo, s11, v17, vcc_lo
	s_clause 0x1
	global_load_b64 v[4:5], v[4:5], off offset:8
	global_load_b64 v[16:17], v[16:17], off
	v_fma_f64 v[18:19], -v[18:19], s[24:25], v[8:9]
	v_fma_f64 v[8:9], v[24:25], s[24:25], -v[8:9]
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mul_f64 v[4:5], v[18:19], v[4:5]
	s_waitcnt vmcnt(0)
	v_fma_f64 v[4:5], v[8:9], v[16:17], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_scale_f64 v[8:9], null, s[24:25], s[24:25], v[4:5]
	v_rcp_f64_e32 v[16:17], v[8:9]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[18:19], -v[8:9], v[16:17], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[16:17], v[16:17], v[18:19], v[16:17]
	v_fma_f64 v[18:19], -v[8:9], v[16:17], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[16:17], v[16:17], v[18:19], v[16:17]
	v_div_scale_f64 v[18:19], vcc_lo, v[4:5], s[24:25], v[4:5]
	v_mul_f64 v[20:21], v[18:19], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[8:9], -v[8:9], v[20:21], v[18:19]
	v_div_fmas_f64 v[8:9], v[8:9], v[16:17], v[20:21]
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fixup_f64 v[20:21], v[8:9], s[24:25], v[4:5]
.LBB0_36:                               ; %Flow142
	s_or_b32 exec_lo, exec_lo, s7
.LBB0_37:                               ; %Flow144
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s6
                                        ; implicit-def: $vgpr30
                                        ; implicit-def: $vgpr28
                                        ; implicit-def: $vgpr29
                                        ; implicit-def: $vgpr31
                                        ; implicit-def: $vgpr4_vgpr5
                                        ; implicit-def: $vgpr15_vgpr16
.LBB0_38:                               ; %Flow149
	s_and_not1_saveexec_b32 s3, s3
	s_cbranch_execz .LBB0_46
; %bb.39:
                                        ; implicit-def: $vgpr20_vgpr21
	s_and_saveexec_b32 s6, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s2, exec_lo, s6
	s_cbranch_execz .LBB0_43
; %bb.40:
	v_ashrrev_i32_e32 v9, 31, v31
	v_add_co_u32 v8, vcc_lo, v31, v28
	s_add_i32 s6, s19, -1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v9, vcc_lo, v9, v16, vcc_lo
	v_lshlrev_b64 v[8:9], 3, v[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v8, vcc_lo, s10, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s11, v9, vcc_lo
	v_cmp_gt_i32_e32 vcc_lo, s6, v29
	global_load_b64 v[20:21], v[8:9], off offset:8
	s_and_saveexec_b32 s6, vcc_lo
	s_cbranch_execz .LBB0_42
; %bb.41:
	v_ashrrev_i32_e32 v9, 31, v30
	v_add_co_u32 v8, vcc_lo, v30, v28
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v9, vcc_lo, v9, v16, vcc_lo
	v_cvt_f64_i32_e32 v[16:17], v29
	v_lshlrev_b64 v[8:9], 3, v[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v8, vcc_lo, s10, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s11, v9, vcc_lo
	global_load_b64 v[8:9], v[8:9], off offset:8
	s_waitcnt vmcnt(2)
	v_add_f64 v[18:19], v[16:17], 0.5
	v_add_f64 v[16:17], 0x3ff80000, v[16:17]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[18:19], -v[18:19], s[26:27], v[4:5]
	v_fma_f64 v[4:5], v[16:17], s[26:27], -v[4:5]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[8:9], v[18:19], v[8:9]
	v_fma_f64 v[4:5], v[4:5], v[20:21], v[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_scale_f64 v[8:9], null, s[26:27], s[26:27], v[4:5]
	v_rcp_f64_e32 v[16:17], v[8:9]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[18:19], -v[8:9], v[16:17], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[16:17], v[16:17], v[18:19], v[16:17]
	v_fma_f64 v[18:19], -v[8:9], v[16:17], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[16:17], v[16:17], v[18:19], v[16:17]
	v_div_scale_f64 v[18:19], vcc_lo, v[4:5], s[26:27], v[4:5]
	v_mul_f64 v[20:21], v[18:19], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[8:9], -v[8:9], v[20:21], v[18:19]
	v_div_fmas_f64 v[8:9], v[8:9], v[16:17], v[20:21]
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fixup_f64 v[20:21], v[8:9], s[26:27], v[4:5]
.LBB0_42:                               ; %Flow145
	s_or_b32 exec_lo, exec_lo, s6
                                        ; implicit-def: $vgpr30
                                        ; implicit-def: $vgpr28
                                        ; implicit-def: $vgpr15_vgpr16
.LBB0_43:                               ; %Flow146
	s_and_not1_saveexec_b32 s2, s2
	s_cbranch_execz .LBB0_45
; %bb.44:
	v_ashrrev_i32_e32 v5, 31, v30
	v_add_co_u32 v4, vcc_lo, v30, v28
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v5, vcc_lo, v5, v16, vcc_lo
	v_lshlrev_b64 v[4:5], 3, v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, s10, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v5, vcc_lo
	global_load_b64 v[20:21], v[4:5], off offset:8
.LBB0_45:                               ; %Flow147
	s_or_b32 exec_lo, exec_lo, s2
.LBB0_46:
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s3
	v_add_co_u32 v8, vcc_lo, s50, v10
	v_add_co_ci_u32_e32 v9, vcc_lo, s51, v11, vcc_lo
	v_add_co_u32 v16, vcc_lo, s12, v10
	v_add_co_ci_u32_e32 v17, vcc_lo, s13, v11, vcc_lo
	global_load_b64 v[28:29], v[8:9], off
	v_add_co_u32 v4, vcc_lo, s46, v10
	v_add_co_ci_u32_e32 v5, vcc_lo, s47, v11, vcc_lo
	v_add_co_u32 v18, vcc_lo, s4, v10
	s_waitcnt vmcnt(1)
	v_add_co_ci_u32_e32 v19, vcc_lo, s5, v11, vcc_lo
	global_load_b64 v[24:25], v[16:17], off
	global_load_b64 v[22:23], v[4:5], off
	global_load_b64 v[26:27], v[18:19], off
                                        ; implicit-def: $vgpr42
                                        ; implicit-def: $vgpr30_vgpr31
                                        ; implicit-def: $vgpr32_vgpr33
	s_waitcnt vmcnt(3)
	v_cmp_ngt_f64_e64 s3, 0x41d00000, |v[28:29]|
	v_trig_preop_f64 v[38:39], |v[28:29]|, 0
	v_trig_preop_f64 v[18:19], |v[28:29]|, 1
	v_ldexp_f64 v[40:41], |v[28:29]|, 0xffffff80
	v_trig_preop_f64 v[16:17], |v[28:29]|, 2
	v_and_b32_e32 v44, 0x7fffffff, v29
	s_and_saveexec_b32 s2, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s4, exec_lo, s2
	s_cbranch_execz .LBB0_48
; %bb.47:
	v_cmp_le_f64_e64 vcc_lo, 0x7b000000, |v[28:29]|
	v_mov_b32_e32 v55, 0
	s_mov_b32 s6, 0x54442d18
	s_mov_b32 s7, 0x3ff921fb
	s_mov_b32 s8, 0x33145c07
	s_mov_b32 s9, 0x3c91a626
	v_cndmask_b32_e32 v31, v44, v41, vcc_lo
	v_cndmask_b32_e32 v30, v28, v40, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f64 v[32:33], v[38:39], v[30:31]
	v_mul_f64 v[34:35], v[18:19], v[30:31]
	v_fma_f64 v[36:37], v[38:39], v[30:31], -v[32:33]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[53:54], v[18:19], v[30:31], -v[34:35]
	v_add_f64 v[42:43], v[34:35], v[36:37]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[45:46], v[42:43], -v[34:35]
	v_add_f64 v[49:50], v[32:33], v[42:43]
	v_add_f64 v[47:48], v[42:43], -v[45:46]
	v_add_f64 v[36:37], v[36:37], -v[45:46]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_ldexp_f64 v[45:46], v[49:50], -2
	v_add_f64 v[32:33], v[49:50], -v[32:33]
	v_add_f64 v[34:35], v[34:35], -v[47:48]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_neq_f64_e64 vcc_lo, 0x7ff00000, |v[45:46]|
	v_add_f64 v[32:33], v[42:43], -v[32:33]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_f64 v[34:35], v[36:37], v[34:35]
	v_fract_f64_e32 v[36:37], v[45:46]
	v_dual_cndmask_b32 v37, 0, v37 :: v_dual_cndmask_b32 v36, 0, v36
	v_mul_f64 v[51:52], v[16:17], v[30:31]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ldexp_f64 v[36:37], v[36:37], 2
	v_add_f64 v[47:48], v[51:52], v[53:54]
	v_fma_f64 v[30:31], v[16:17], v[30:31], -v[51:52]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[42:43], v[47:48], v[34:35]
	v_add_f64 v[45:46], v[32:33], v[42:43]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[49:50], v[45:46], v[36:37]
	v_add_f64 v[32:33], v[45:46], -v[32:33]
	v_cmp_gt_f64_e32 vcc_lo, 0, v[49:50]
	v_add_f64 v[49:50], v[47:48], -v[51:52]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[32:33], v[42:43], -v[32:33]
	v_cndmask_b32_e64 v56, 0, 0x40100000, vcc_lo
	v_add_f64 v[60:61], v[47:48], -v[49:50]
	v_add_f64 v[49:50], v[53:54], -v[49:50]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_f64 v[36:37], v[36:37], v[55:56]
	v_add_f64 v[56:57], v[42:43], -v[47:48]
	v_add_f64 v[53:54], v[51:52], -v[60:61]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[58:59], v[45:46], v[36:37]
	v_add_f64 v[62:63], v[42:43], -v[56:57]
	v_add_f64 v[34:35], v[34:35], -v[56:57]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f64 v[49:50], v[49:50], v[53:54]
	v_cvt_i32_f64_e32 v58, v[58:59]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[47:48], v[47:48], -v[62:63]
	v_cvt_f64_i32_e32 v[56:57], v58
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[34:35], v[34:35], v[47:48]
	v_add_f64 v[36:37], v[36:37], -v[56:57]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[34:35], v[49:50], v[34:35]
	v_add_f64 v[47:48], v[45:46], v[36:37]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[30:31], v[30:31], v[34:35]
	v_add_f64 v[34:35], v[47:48], -v[36:37]
	v_cmp_le_f64_e32 vcc_lo, 0.5, v[47:48]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[30:31], v[32:33], v[30:31]
	v_add_f64 v[32:33], v[45:46], -v[34:35]
	v_cndmask_b32_e64 v56, 0, 0x3ff00000, vcc_lo
	v_add_co_ci_u32_e64 v42, s2, 0, v58, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[30:31], v[30:31], v[32:33]
	v_add_f64 v[32:33], v[47:48], -v[55:56]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[34:35], v[32:33], v[30:31]
	v_mul_f64 v[36:37], v[34:35], s[6:7]
	v_add_f64 v[32:33], v[34:35], -v[32:33]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[45:46], v[34:35], s[6:7], -v[36:37]
	v_add_f64 v[30:31], v[30:31], -v[32:33]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[32:33], v[34:35], s[8:9], v[45:46]
	v_fma_f64 v[32:33], v[30:31], s[6:7], v[32:33]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[30:31], v[36:37], v[32:33]
	v_add_f64 v[34:35], v[30:31], -v[36:37]
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[32:33], v[32:33], -v[34:35]
.LBB0_48:                               ; %Flow135
	s_or_saveexec_b32 s2, s4
	s_load_b64 s[6:7], s[0:1], 0xc8
	s_xor_b32 exec_lo, exec_lo, s2
	s_cbranch_execz .LBB0_50
; %bb.49:
	s_mov_b32 s4, 0x6dc9c883
	s_mov_b32 s5, 0x3fe45f30
	s_mov_b32 s9, 0xbc91a626
	v_mul_f64 v[30:31], |v[28:29]|, s[4:5]
	s_mov_b32 s4, 0x54442d18
	s_mov_b32 s5, 0xbff921fb
	s_mov_b32 s8, 0x33145c00
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rndne_f64_e32 v[34:35], v[30:31]
	v_fma_f64 v[30:31], v[34:35], s[4:5], |v[28:29]|
	v_mul_f64 v[32:33], v[34:35], s[8:9]
	s_mov_b32 s4, 0x252049c0
	s_mov_b32 s5, 0xb97b839a
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[42:43], v[34:35], s[8:9], v[30:31]
	v_add_f64 v[36:37], v[30:31], v[32:33]
	s_mov_b32 s9, 0x3c91a626
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[30:31], v[30:31], -v[36:37]
	v_add_f64 v[36:37], v[36:37], -v[42:43]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[30:31], v[30:31], v[32:33]
	v_fma_f64 v[32:33], v[34:35], s[8:9], v[32:33]
	v_add_f64 v[30:31], v[36:37], v[30:31]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[30:31], v[30:31], -v[32:33]
	v_fma_f64 v[32:33], v[34:35], s[4:5], v[30:31]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[30:31], v[42:43], v[32:33]
	v_add_f64 v[36:37], v[30:31], -v[42:43]
	v_cvt_i32_f64_e32 v42, v[34:35]
	s_delay_alu instid0(VALU_DEP_2)
	v_add_f64 v[32:33], v[32:33], -v[36:37]
.LBB0_50:                               ; %_ZL3cosd.exit.i
	s_or_b32 exec_lo, exec_lo, s2
                                        ; implicit-def: $vgpr43
                                        ; implicit-def: $vgpr34_vgpr35
                                        ; implicit-def: $vgpr36_vgpr37
	s_and_saveexec_b32 s2, s3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s3, exec_lo, s2
	s_cbranch_execz .LBB0_52
; %bb.51:
	v_cmp_le_f64_e64 vcc_lo, 0x7b000000, |v[28:29]|
	v_mov_b32_e32 v53, 0
	s_mov_b32 s4, 0x54442d18
	s_mov_b32 s5, 0x3ff921fb
	s_mov_b32 s8, 0x33145c07
	s_mov_b32 s9, 0x3c91a626
	v_cndmask_b32_e32 v35, v44, v41, vcc_lo
	v_cndmask_b32_e32 v34, v28, v40, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f64 v[36:37], v[38:39], v[34:35]
	v_mul_f64 v[40:41], v[18:19], v[34:35]
	v_fma_f64 v[38:39], v[38:39], v[34:35], -v[36:37]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[18:19], v[18:19], v[34:35], -v[40:41]
	v_add_f64 v[43:44], v[40:41], v[38:39]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[45:46], v[43:44], -v[40:41]
	v_add_f64 v[49:50], v[36:37], v[43:44]
	v_add_f64 v[47:48], v[43:44], -v[45:46]
	v_add_f64 v[38:39], v[38:39], -v[45:46]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_ldexp_f64 v[45:46], v[49:50], -2
	v_add_f64 v[36:37], v[49:50], -v[36:37]
	v_add_f64 v[40:41], v[40:41], -v[47:48]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_neq_f64_e64 vcc_lo, 0x7ff00000, |v[45:46]|
	v_add_f64 v[36:37], v[43:44], -v[36:37]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_f64 v[38:39], v[38:39], v[40:41]
	v_fract_f64_e32 v[40:41], v[45:46]
	v_dual_cndmask_b32 v41, 0, v41 :: v_dual_cndmask_b32 v40, 0, v40
	v_mul_f64 v[51:52], v[16:17], v[34:35]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ldexp_f64 v[40:41], v[40:41], 2
	v_add_f64 v[47:48], v[51:52], v[18:19]
	v_fma_f64 v[16:17], v[16:17], v[34:35], -v[51:52]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[43:44], v[47:48], v[38:39]
	v_add_f64 v[45:46], v[36:37], v[43:44]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[49:50], v[45:46], v[40:41]
	v_add_f64 v[34:35], v[45:46], -v[36:37]
	v_cmp_gt_f64_e32 vcc_lo, 0, v[49:50]
	v_add_f64 v[49:50], v[47:48], -v[51:52]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[34:35], v[43:44], -v[34:35]
	v_cndmask_b32_e64 v54, 0, 0x40100000, vcc_lo
	v_add_f64 v[58:59], v[47:48], -v[49:50]
	v_add_f64 v[18:19], v[18:19], -v[49:50]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_f64 v[40:41], v[40:41], v[53:54]
	v_add_f64 v[54:55], v[43:44], -v[47:48]
	v_add_f64 v[49:50], v[51:52], -v[58:59]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[56:57], v[45:46], v[40:41]
	v_add_f64 v[60:61], v[43:44], -v[54:55]
	v_add_f64 v[38:39], v[38:39], -v[54:55]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f64 v[18:19], v[18:19], v[49:50]
	v_cvt_i32_f64_e32 v56, v[56:57]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[47:48], v[47:48], -v[60:61]
	v_cvt_f64_i32_e32 v[54:55], v56
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[38:39], v[38:39], v[47:48]
	v_add_f64 v[40:41], v[40:41], -v[54:55]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[18:19], v[18:19], v[38:39]
	v_add_f64 v[36:37], v[45:46], v[40:41]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[16:17], v[16:17], v[18:19]
	v_add_f64 v[18:19], v[36:37], -v[40:41]
	v_cmp_le_f64_e32 vcc_lo, 0.5, v[36:37]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[16:17], v[34:35], v[16:17]
	v_add_f64 v[18:19], v[45:46], -v[18:19]
	v_cndmask_b32_e64 v54, 0, 0x3ff00000, vcc_lo
	v_add_co_ci_u32_e64 v43, s2, 0, v56, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[16:17], v[16:17], v[18:19]
	v_add_f64 v[18:19], v[36:37], -v[53:54]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[34:35], v[18:19], v[16:17]
	v_mul_f64 v[36:37], v[34:35], s[4:5]
	v_add_f64 v[18:19], v[34:35], -v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[38:39], v[34:35], s[4:5], -v[36:37]
	v_add_f64 v[16:17], v[16:17], -v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[18:19], v[34:35], s[8:9], v[38:39]
	v_fma_f64 v[16:17], v[16:17], s[4:5], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[34:35], v[36:37], v[16:17]
	v_add_f64 v[18:19], v[34:35], -v[36:37]
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[36:37], v[16:17], -v[18:19]
.LBB0_52:                               ; %Flow134
	s_and_not1_saveexec_b32 s2, s3
	s_cbranch_execz .LBB0_54
; %bb.53:
	s_mov_b32 s4, 0x6dc9c883
	s_mov_b32 s5, 0x3fe45f30
	s_mov_b32 s9, 0xbc91a626
	v_mul_f64 v[16:17], |v[28:29]|, s[4:5]
	s_mov_b32 s4, 0x54442d18
	s_mov_b32 s5, 0xbff921fb
	s_mov_b32 s8, 0x33145c00
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rndne_f64_e32 v[16:17], v[16:17]
	v_fma_f64 v[18:19], v[16:17], s[4:5], |v[28:29]|
	v_mul_f64 v[34:35], v[16:17], s[8:9]
	s_mov_b32 s4, 0x252049c0
	s_mov_b32 s5, 0xb97b839a
	v_cvt_i32_f64_e32 v43, v[16:17]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[38:39], v[16:17], s[8:9], v[18:19]
	v_add_f64 v[36:37], v[18:19], v[34:35]
	s_mov_b32 s9, 0x3c91a626
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[18:19], v[18:19], -v[36:37]
	v_add_f64 v[36:37], v[36:37], -v[38:39]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[18:19], v[18:19], v[34:35]
	v_fma_f64 v[34:35], v[16:17], s[8:9], v[34:35]
	v_add_f64 v[18:19], v[36:37], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[18:19], v[18:19], -v[34:35]
	v_fma_f64 v[18:19], v[16:17], s[4:5], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[34:35], v[38:39], v[18:19]
	v_add_f64 v[36:37], v[34:35], -v[38:39]
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[36:37], v[18:19], -v[36:37]
.LBB0_54:                               ; %_ZL3sind.exit.i
	s_or_b32 exec_lo, exec_lo, s2
	v_add_co_u32 v16, vcc_lo, s44, v10
	v_add_co_ci_u32_e32 v17, vcc_lo, s45, v11, vcc_lo
	v_add_co_u32 v18, vcc_lo, s42, v10
	v_add_co_ci_u32_e32 v19, vcc_lo, s43, v11, vcc_lo
	global_load_b64 v[38:39], v[16:17], off
	v_mul_f64 v[46:47], s[22:23], s[22:23]
	s_mov_b32 s2, 0xb42fdfa7
	global_load_b64 v[40:41], v[18:19], off
	s_mov_b32 s4, 0xf9a43bb8
	s_mov_b32 s3, 0xbe5ae600
	s_mov_b32 s5, 0x3de5e0b2
	s_mov_b32 s10, 0x796cde01
	s_mov_b32 s11, 0x3ec71de3
	s_mov_b32 s8, 0x46cc5e42
	s_mov_b32 s9, 0xbda907db
	v_mul_f64 v[74:75], v[32:33], 0.5
	v_mul_f64 v[82:83], v[36:37], 0.5
	s_waitcnt vmcnt(1)
	v_mul_f64 v[44:45], v[38:39], v[38:39]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[44:45], v[40:41], v[40:41], v[44:45]
	v_fma_f64 v[44:45], v[22:23], v[22:23], v[44:45]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f64 v[48:49], null, v[46:47], v[46:47], v[44:45]
	v_div_scale_f64 v[54:55], vcc_lo, v[44:45], v[46:47], v[44:45]
	v_rcp_f64_e32 v[50:51], v[48:49]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[52:53], -v[48:49], v[50:51], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[50:51], v[50:51], v[52:53], v[50:51]
	v_fma_f64 v[52:53], -v[48:49], v[50:51], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[50:51], v[50:51], v[52:53], v[50:51]
	v_mul_f64 v[52:53], v[54:55], v[50:51]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[48:49], -v[48:49], v[52:53], v[54:55]
	v_div_fmas_f64 v[48:49], v[48:49], v[50:51], v[52:53]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_fixup_f64 v[44:45], v[48:49], v[46:47], v[44:45]
	v_mul_f64 v[48:49], v[30:31], v[30:31]
	v_add_f64 v[44:45], -v[44:45], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_mul_f64 v[60:61], v[48:49], 0.5
	v_mul_f64 v[76:77], v[30:31], -v[48:49]
	v_cmp_gt_f64_e32 vcc_lo, 0x10000000, v[44:45]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[70:71], -v[60:61], 1.0
	v_cndmask_b32_e64 v50, 0, 1, vcc_lo
	v_add_f64 v[80:81], -v[70:71], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b32_e32 v52, 8, v50
	v_mul_f64 v[50:51], v[34:35], v[34:35]
	v_ldexp_f64 v[44:45], v[44:45], v52
	v_fma_f64 v[52:53], v[48:49], s[4:5], s[2:3]
	v_add_f64 v[60:61], v[80:81], -v[60:61]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_4)
	v_fma_f64 v[54:55], v[50:51], s[4:5], s[2:3]
	s_mov_b32 s4, 0x9037ab78
	s_mov_b32 s5, 0x3e21eeb6
	v_mul_f64 v[64:65], v[50:51], 0.5
	v_rsq_f64_e32 v[56:57], v[44:45]
	v_fma_f64 v[52:53], v[48:49], v[52:53], s[10:11]
	v_fma_f64 v[58:59], v[48:49], s[8:9], s[4:5]
	v_fma_f64 v[66:67], v[50:51], s[8:9], s[4:5]
	s_mov_b32 s4, 0x19e83e5c
	s_mov_b32 s5, 0xbf2a01a0
	v_add_co_u32 v62, s2, s14, v10
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v63, s2, s15, v11, s2
	s_mov_b32 s2, 0xa17f65f6
	s_mov_b32 s3, 0xbe927e4f
	s_mov_b32 s8, 0x19f4ec90
	global_load_b64 v[62:63], v[62:63], off
	s_mov_b32 s9, 0x3efa01a0
	v_mul_f64 v[84:85], v[34:35], -v[50:51]
	v_fma_f64 v[60:61], v[30:31], -v[32:33], v[60:61]
	v_fma_f64 v[54:55], v[50:51], v[54:55], s[10:11]
	v_add_f64 v[72:73], -v[64:65], 1.0
	v_fma_f64 v[52:53], v[48:49], v[52:53], s[4:5]
	v_mul_f64 v[68:69], v[44:45], v[56:57]
	v_mul_f64 v[56:57], v[56:57], 0.5
	v_fma_f64 v[58:59], v[48:49], v[58:59], s[2:3]
	v_fma_f64 v[66:67], v[50:51], v[66:67], s[2:3]
	s_mov_b32 s2, 0x11110bb3
	s_mov_b32 s3, 0x3f811111
	v_fma_f64 v[54:55], v[50:51], v[54:55], s[4:5]
	v_add_f64 v[86:87], -v[72:73], 1.0
	v_fma_f64 v[52:53], v[48:49], v[52:53], s[2:3]
	v_fma_f64 v[78:79], -v[56:57], v[68:69], 0.5
	v_fma_f64 v[58:59], v[48:49], v[58:59], s[8:9]
	v_fma_f64 v[66:67], v[50:51], v[66:67], s[8:9]
	v_fma_f64 v[54:55], v[50:51], v[54:55], s[2:3]
	s_mov_b32 s2, 0x16c16967
	s_mov_b32 s3, 0xbf56c16c
	v_add_f64 v[64:65], v[86:87], -v[64:65]
	v_fma_f64 v[52:53], v[76:77], v[52:53], v[74:75]
	v_mul_f64 v[74:75], v[48:49], v[48:49]
	v_fma_f64 v[68:69], v[68:69], v[78:79], v[68:69]
	v_fma_f64 v[58:59], v[48:49], v[58:59], s[2:3]
	v_fma_f64 v[66:67], v[50:51], v[66:67], s[2:3]
	v_fma_f64 v[56:57], v[56:57], v[78:79], v[56:57]
	s_mov_b32 s2, 0x55555555
	s_mov_b32 s3, 0x3fa55555
	v_fma_f64 v[54:55], v[84:85], v[54:55], v[82:83]
	v_fma_f64 v[64:65], v[34:35], -v[36:37], v[64:65]
	v_fma_f64 v[32:33], v[48:49], v[52:53], -v[32:33]
	v_fma_f64 v[78:79], -v[68:69], v[68:69], v[44:45]
	v_fma_f64 v[58:59], v[48:49], v[58:59], s[2:3]
	v_mul_f64 v[48:49], v[50:51], v[50:51]
	v_fma_f64 v[52:53], v[50:51], v[66:67], s[2:3]
	s_mov_b32 s3, 0xbfc55555
	v_fma_f64 v[36:37], v[50:51], v[54:55], -v[36:37]
	v_fma_f64 v[32:33], v[76:77], s[2:3], v[32:33]
	v_fma_f64 v[50:51], v[78:79], v[56:57], v[68:69]
	v_fma_f64 v[54:55], v[74:75], v[58:59], v[60:61]
	v_mul_f64 v[58:59], v[20:21], v[22:23]
	v_fma_f64 v[48:49], v[48:49], v[52:53], v[64:65]
	v_mul_f64 v[20:21], v[20:21], v[26:27]
	v_fma_f64 v[36:37], v[84:85], s[2:3], v[36:37]
	v_cmp_class_f64_e64 s3, v[28:29], 0x1f8
	v_lshlrev_b32_e32 v28, 30, v43
	v_add_f64 v[30:31], v[30:31], -v[32:33]
	v_fma_f64 v[52:53], -v[50:51], v[50:51], v[44:45]
	v_add_f64 v[54:55], v[70:71], v[54:55]
	v_fma_f64 v[12:13], -v[58:59], s[20:21], v[12:13]
	v_add_f64 v[32:33], v[72:73], v[48:49]
	v_and_b32_e32 v48, 1, v42
	v_lshlrev_b32_e32 v42, 30, v42
	v_xor_b32_e32 v28, v28, v29
	v_fma_f64 v[14:15], v[20:21], s[20:21], v[14:15]
	s_delay_alu instid0(VALU_DEP_4)
	v_cmp_eq_u32_e64 s2, 0, v48
	v_add_f64 v[34:35], v[34:35], -v[36:37]
	v_xor_b32_e32 v29, 0x80000000, v31
	v_and_b32_e32 v31, 0x80000000, v42
	v_fma_f64 v[36:37], v[52:53], v[56:57], v[50:51]
	v_and_b32_e32 v42, 0x80000000, v28
	v_mul_f64 v[12:13], v[24:25], v[12:13]
	v_cndmask_b32_e64 v28, v29, v55, s2
	v_cndmask_b32_e64 v29, 0, 0xffffff80, vcc_lo
	v_cmp_class_f64_e64 vcc_lo, v[44:45], 0x260
	v_cndmask_b32_e64 v30, v30, v54, s2
	v_mul_f64 v[14:15], v[24:25], v[14:15]
	v_xor_b32_e32 v31, v28, v31
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v30, 0, v30, s3
	v_cndmask_b32_e64 v31, 0x7ff80000, v31, s3
	v_ldexp_f64 v[28:29], v[36:37], v29
	v_and_b32_e32 v43, 1, v43
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_mul_f64 v[30:31], v[12:13], v[30:31]
	s_waitcnt vmcnt(0)
	v_div_scale_f64 v[50:51], null, v[62:63], v[62:63], v[14:15]
	v_cndmask_b32_e32 v28, v28, v44, vcc_lo
	v_cmp_eq_u32_e64 s4, 0, v43
	v_cndmask_b32_e32 v29, v29, v45, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v33, v33, v35, s4
	v_cndmask_b32_e64 v32, v32, v34, s4
	v_div_scale_f64 v[26:27], null, v[28:29], v[28:29], 1.0
	v_div_scale_f64 v[56:57], vcc_lo, 1.0, v[28:29], 1.0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_xor_b32_e32 v33, v33, v42
	v_cndmask_b32_e64 v32, 0, v32, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v33, 0x7ff80000, v33, s3
	v_mul_f64 v[12:13], v[12:13], v[32:33]
	v_div_scale_f64 v[32:33], null, v[62:63], v[62:63], v[30:31]
	v_rcp_f64_e32 v[20:21], v[26:27]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_scale_f64 v[34:35], null, v[62:63], v[62:63], v[12:13]
	v_rcp_f64_e32 v[36:37], v[32:33]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[24:25], -v[26:27], v[20:21], 1.0
	v_rcp_f64_e32 v[42:43], v[34:35]
	v_fma_f64 v[44:45], -v[32:33], v[36:37], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_3)
	v_fma_f64 v[20:21], v[20:21], v[24:25], v[20:21]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[48:49], -v[34:35], v[42:43], 1.0
	v_fma_f64 v[24:25], v[36:37], v[44:45], v[36:37]
	v_fma_f64 v[44:45], -v[26:27], v[20:21], 1.0
	v_fma_f64 v[36:37], v[42:43], v[48:49], v[42:43]
	v_rcp_f64_e32 v[42:43], v[50:51]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[48:49], -v[32:33], v[24:25], 1.0
	v_fma_f64 v[20:21], v[20:21], v[44:45], v[20:21]
	v_div_scale_f64 v[44:45], s2, v[30:31], v[62:63], v[30:31]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[52:53], -v[34:35], v[36:37], 1.0
	v_fma_f64 v[24:25], v[24:25], v[48:49], v[24:25]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[54:55], -v[50:51], v[42:43], 1.0
	v_div_scale_f64 v[48:49], s3, v[12:13], v[62:63], v[12:13]
	v_fma_f64 v[36:37], v[36:37], v[52:53], v[36:37]
	v_mul_f64 v[52:53], v[56:57], v[20:21]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_fma_f64 v[42:43], v[42:43], v[54:55], v[42:43]
	v_mul_f64 v[54:55], v[44:45], v[24:25]
	v_mul_f64 v[58:59], v[48:49], v[36:37]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[26:27], -v[26:27], v[52:53], v[56:57]
	v_fma_f64 v[60:61], -v[50:51], v[42:43], 1.0
	s_delay_alu instid0(VALU_DEP_4)
	v_fma_f64 v[32:33], -v[32:33], v[54:55], v[44:45]
	v_div_scale_f64 v[44:45], s4, v[14:15], v[62:63], v[14:15]
	v_fma_f64 v[34:35], -v[34:35], v[58:59], v[48:49]
	v_div_fmas_f64 v[20:21], v[26:27], v[20:21], v[52:53]
	s_mov_b32 vcc_lo, s2
	v_fma_f64 v[42:43], v[42:43], v[60:61], v[42:43]
	v_div_fmas_f64 v[24:25], v[32:33], v[24:25], v[54:55]
	s_mov_b32 vcc_lo, s3
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_div_fmas_f64 v[26:27], v[34:35], v[36:37], v[58:59]
	v_div_fixup_f64 v[20:21], v[20:21], v[28:29], 1.0
	s_mov_b32 vcc_lo, s4
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_mul_f64 v[32:33], v[44:45], v[42:43]
	v_div_fixup_f64 v[24:25], v[24:25], v[62:63], v[30:31]
	s_mov_b32 s4, 0xb5e68a13
	s_mov_b32 s5, 0x3eeba404
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_div_fixup_f64 v[12:13], v[26:27], v[62:63], v[12:13]
	v_mul_f64 v[28:29], v[40:41], v[20:21]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[26:27], -v[50:51], v[32:33], v[44:45]
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[12:13], v[12:13], s[6:7]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_fmas_f64 v[26:27], v[26:27], v[42:43], v[32:33]
	v_fma_f64 v[12:13], v[38:39], v[20:21], v[12:13]
	v_mul_f64 v[20:21], v[22:23], v[20:21]
	v_fma_f64 v[22:23], v[24:25], s[6:7], v[28:29]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_div_fixup_f64 v[14:15], v[26:27], v[62:63], v[14:15]
	v_mul_f64 v[24:25], v[12:13], v[12:13]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[14:15], v[14:15], s[6:7], v[20:21]
	v_fma_f64 v[20:21], v[22:23], v[22:23], v[24:25]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[20:21], v[14:15], v[14:15], v[20:21]
	v_div_scale_f64 v[24:25], null, v[46:47], v[46:47], v[20:21]
	v_div_scale_f64 v[30:31], vcc_lo, v[20:21], v[46:47], v[20:21]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[26:27], v[24:25]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[28:29], -v[24:25], v[26:27], 1.0
	v_fma_f64 v[26:27], v[26:27], v[28:29], v[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[28:29], -v[24:25], v[26:27], 1.0
	v_fma_f64 v[26:27], v[26:27], v[28:29], v[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[28:29], v[30:31], v[26:27]
	v_fma_f64 v[24:25], -v[24:25], v[28:29], v[30:31]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[24:25], v[24:25], v[26:27], v[28:29]
	v_div_fixup_f64 v[20:21], v[24:25], v[46:47], v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[20:21], v[20:21], 1.0
	v_cmp_gt_f64_e32 vcc_lo, 0x10000000, v[20:21]
	v_cndmask_b32_e64 v24, 0, 1, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v24, 8, v24
	v_ldexp_f64 v[20:21], v[20:21], v24
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_rsq_f64_e32 v[24:25], v[20:21]
	s_waitcnt_depctr 0xfff
	v_mul_f64 v[26:27], v[20:21], v[24:25]
	v_mul_f64 v[24:25], v[24:25], 0.5
	v_fma_f64 v[28:29], -v[24:25], v[26:27], 0.5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[26:27], v[26:27], v[28:29], v[26:27]
	v_fma_f64 v[24:25], v[24:25], v[28:29], v[24:25]
	v_fma_f64 v[28:29], -v[26:27], v[26:27], v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[26:27], v[28:29], v[24:25], v[26:27]
	v_fma_f64 v[28:29], -v[26:27], v[26:27], v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f64 v[24:25], v[28:29], v[24:25], v[26:27]
	v_cndmask_b32_e64 v26, 0, 0xffffff80, vcc_lo
	v_cmp_class_f64_e64 vcc_lo, v[20:21], 0x260
	v_ldexp_f64 v[24:25], v[24:25], v26
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v21, v25, v21 :: v_dual_cndmask_b32 v20, v24, v20
	v_div_scale_f64 v[24:25], null, v[20:21], v[20:21], v[22:23]
	v_div_scale_f64 v[26:27], null, v[20:21], v[20:21], v[12:13]
	v_div_scale_f64 v[28:29], null, v[20:21], v[20:21], v[14:15]
	v_div_scale_f64 v[42:43], vcc_lo, v[22:23], v[20:21], v[22:23]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_rcp_f64_e32 v[30:31], v[24:25]
	v_rcp_f64_e32 v[32:33], v[26:27]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(TRANS32_DEP_3)
	v_rcp_f64_e32 v[34:35], v[28:29]
	v_fma_f64 v[36:37], -v[24:25], v[30:31], 1.0
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[38:39], -v[26:27], v[32:33], 1.0
	v_fma_f64 v[40:41], -v[28:29], v[34:35], 1.0
	v_fma_f64 v[30:31], v[30:31], v[36:37], v[30:31]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[32:33], v[32:33], v[38:39], v[32:33]
	v_fma_f64 v[34:35], v[34:35], v[40:41], v[34:35]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[36:37], -v[24:25], v[30:31], 1.0
	v_fma_f64 v[38:39], -v[26:27], v[32:33], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[40:41], -v[28:29], v[34:35], 1.0
	v_fma_f64 v[30:31], v[30:31], v[36:37], v[30:31]
	v_div_scale_f64 v[36:37], s2, v[12:13], v[20:21], v[12:13]
	s_delay_alu instid0(VALU_DEP_4)
	v_fma_f64 v[32:33], v[32:33], v[38:39], v[32:33]
	v_div_scale_f64 v[38:39], s3, v[14:15], v[20:21], v[14:15]
	v_fma_f64 v[34:35], v[34:35], v[40:41], v[34:35]
	v_add_f64 v[40:41], v[20:21], -1.0
	v_mul_f64 v[44:45], v[42:43], v[30:31]
	v_mul_f64 v[46:47], v[36:37], v[32:33]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_mul_f64 v[48:49], v[38:39], v[34:35]
	v_mul_f64 v[40:41], v[62:63], v[40:41]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[24:25], -v[24:25], v[44:45], v[42:43]
	v_fma_f64 v[26:27], -v[26:27], v[46:47], v[36:37]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[28:29], -v[28:29], v[48:49], v[38:39]
	v_mul_f64 v[36:37], v[40:41], s[22:23]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_div_fmas_f64 v[24:25], v[24:25], v[30:31], v[44:45]
	s_mov_b32 vcc_lo, s2
	v_div_fmas_f64 v[26:27], v[26:27], v[32:33], v[46:47]
	s_mov_b32 vcc_lo, s3
	s_mov_b32 s2, 0xbd3237f4
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_div_fmas_f64 v[28:29], v[28:29], v[34:35], v[48:49]
	v_mul_f64 v[30:31], v[36:37], s[22:23]
	s_mov_b32 s3, 0xbf23e260
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_div_fixup_f64 v[22:23], v[24:25], v[20:21], v[22:23]
	v_div_fixup_f64 v[24:25], v[26:27], v[20:21], v[12:13]
	s_delay_alu instid0(VALU_DEP_4)
	v_div_fixup_f64 v[14:15], v[28:29], v[20:21], v[14:15]
	v_add_co_u32 v20, vcc_lo, s16, v10
	v_add_co_ci_u32_e32 v21, vcc_lo, s17, v11, vcc_lo
	v_add_co_u32 v12, vcc_lo, s36, v10
	v_add_co_ci_u32_e32 v13, vcc_lo, s37, v11, vcc_lo
	v_add_co_u32 v10, vcc_lo, s38, v10
	v_add_co_ci_u32_e32 v11, vcc_lo, s39, v11, vcc_lo
	global_store_b64 v[20:21], v[30:31], off
	global_store_b64 v[18:19], v[22:23], off
	global_store_b64 v[16:17], v[24:25], off
	global_store_b64 v[4:5], v[14:15], off
	global_load_b64 v[14:15], v[12:13], off
	global_load_b64 v[20:21], v[18:19], off
	v_mul_f64 v[22:23], s[6:7], 0.5
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[14:15], v[22:23], v[20:21], v[14:15]
	global_store_b64 v[12:13], v[14:15], off
	global_load_b64 v[14:15], v[10:11], off
	global_load_b64 v[20:21], v[16:17], off
	s_waitcnt vmcnt(0)
	v_fma_f64 v[14:15], v[22:23], v[20:21], v[14:15]
	global_store_b64 v[10:11], v[14:15], off
	global_load_b64 v[14:15], v[0:1], off
	global_load_b64 v[20:21], v[4:5], off
	s_waitcnt vmcnt(0)
	v_fma_f64 v[14:15], v[22:23], v[20:21], v[14:15]
	global_store_b64 v[0:1], v[14:15], off
	global_load_b64 v[14:15], v[10:11], off
	global_load_b64 v[20:21], v[12:13], off
	s_waitcnt vmcnt(1)
	v_mul_f64 v[14:15], v[14:15], v[14:15]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[14:15], v[20:21], v[20:21], v[14:15]
	v_cmp_gt_f64_e32 vcc_lo, 0x10000000, v[14:15]
	v_cndmask_b32_e64 v20, 0, 1, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v20, 8, v20
	v_ldexp_f64 v[14:15], v[14:15], v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_rsq_f64_e32 v[20:21], v[14:15]
	s_waitcnt_depctr 0xfff
	v_mul_f64 v[22:23], v[14:15], v[20:21]
	v_mul_f64 v[20:21], v[20:21], 0.5
	v_fma_f64 v[24:25], -v[20:21], v[22:23], 0.5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f64 v[22:23], v[22:23], v[24:25], v[22:23]
	v_fma_f64 v[20:21], v[20:21], v[24:25], v[20:21]
	v_fma_f64 v[24:25], -v[22:23], v[22:23], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[22:23], v[24:25], v[20:21], v[22:23]
	v_fma_f64 v[24:25], -v[22:23], v[22:23], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f64 v[20:21], v[24:25], v[20:21], v[22:23]
	v_cndmask_b32_e64 v22, 0, 0xffffff80, vcc_lo
	v_cmp_class_f64_e64 vcc_lo, v[14:15], 0x260
	v_ldexp_f64 v[20:21], v[20:21], v22
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_cndmask_b32 v15, v21, v15 :: v_dual_cndmask_b32 v14, v20, v14
	global_store_b64 v[2:3], v[14:15], off
	global_load_b64 v[14:15], v[10:11], off
	global_load_b64 v[20:21], v[12:13], off
	s_waitcnt vmcnt(0)
	v_div_scale_f64 v[22:23], null, v[20:21], v[20:21], v[14:15]
	v_div_scale_f64 v[28:29], vcc_lo, v[14:15], v[20:21], v[14:15]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[24:25], v[22:23]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[26:27], -v[22:23], v[24:25], 1.0
	v_fma_f64 v[24:25], v[24:25], v[26:27], v[24:25]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[26:27], -v[22:23], v[24:25], 1.0
	v_fma_f64 v[24:25], v[24:25], v[26:27], v[24:25]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[26:27], v[28:29], v[24:25]
	v_fma_f64 v[22:23], -v[22:23], v[26:27], v[28:29]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[22:23], v[22:23], v[24:25], v[26:27]
	v_div_fixup_f64 v[14:15], v[22:23], v[20:21], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v20, v14 :: v_dual_and_b32 v21, 0x7fffffff, v15
	v_div_scale_f64 v[22:23], null, v[20:21], v[20:21], 1.0
	v_div_scale_f64 v[28:29], vcc_lo, 1.0, v[20:21], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[24:25], v[22:23]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[26:27], -v[22:23], v[24:25], 1.0
	v_fma_f64 v[24:25], v[24:25], v[26:27], v[24:25]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[26:27], -v[22:23], v[24:25], 1.0
	v_fma_f64 v[24:25], v[24:25], v[26:27], v[24:25]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[26:27], v[28:29], v[24:25]
	v_fma_f64 v[22:23], -v[22:23], v[26:27], v[28:29]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_fmas_f64 v[22:23], v[22:23], v[24:25], v[26:27]
	v_cmp_gt_f64_e64 vcc_lo, |v[14:15]|, 1.0
	v_div_fixup_f64 v[22:23], v[22:23], |v[14:15]|, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v21, v21, v23 :: v_dual_cndmask_b32 v20, v20, v22
	v_mul_f64 v[22:23], v[20:21], v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_fma_f64 v[24:25], v[22:23], s[4:5], s[2:3]
	s_mov_b32 s2, 0x69efb384
	s_mov_b32 s3, 0x3f4b2bb0
	s_mov_b32 s4, 0x336a0500
	s_mov_b32 s5, 0x3fedd9ad
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0xaf56de9b
	s_mov_b32 s3, 0xbf67952d
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0xa595c56f
	s_mov_b32 s3, 0x3f7d6d43
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0xa57d9582
	s_mov_b32 s3, 0xbf8c6ea4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x5f08b19f
	s_mov_b32 s3, 0x3f967e29
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0xfc27006a
	s_mov_b32 s3, 0xbf9e9ae6
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x5711927a
	s_mov_b32 s3, 0x3fa2c15b
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0xe82d3ff0
	s_mov_b32 s3, 0xbfa59976
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x6ef28734
	s_mov_b32 s3, 0x3fa82d5d
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x6a214619
	s_mov_b32 s3, 0xbfaae5ce
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x8427b883
	s_mov_b32 s3, 0x3fae1bb4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x8b207f05
	s_mov_b32 s3, 0xbfb110e4
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x57b87036
	s_mov_b32 s3, 0x3fb3b136
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x19378e4f
	s_mov_b32 s3, 0xbfb745d1
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x17e1913c
	s_mov_b32 s3, 0x3fbc71c7
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x92376b7d
	s_mov_b32 s3, 0xbfc24924
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x999952cc
	s_mov_b32 s3, 0x3fc99999
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0x55555523
	s_mov_b32 s3, 0xbfd55555
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[24:25], v[22:23], v[24:25], s[2:3]
	s_mov_b32 s2, 0xeeb562d6
	s_mov_b32 s3, 0x3ffaf154
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[22:23], v[22:23], v[24:25]
	v_fma_f64 v[20:21], v[20:21], v[22:23], v[20:21]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[22:23], s[4:5], s[2:3], -v[20:21]
	v_dual_cndmask_b32 v21, v21, v23 :: v_dual_cndmask_b32 v14, v20, v22
	s_delay_alu instid0(VALU_DEP_1)
	v_bfi_b32 v15, 0x7fffffff, v21, v15
	global_store_b64 v[8:9], v[14:15], off
	global_load_b64 v[20:21], v[2:3], off
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x88
	s_load_b64 s[2:3], s[0:1], 0x98
	s_mov_b32 s0, exec_lo
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_cmpx_lt_f64_e32 s[4:5], v[20:21]
	s_cbranch_execz .LBB0_56
; %bb.55:
	v_mov_b32_e32 v20, 1
	global_store_b32 v[6:7], v20, off
.LBB0_56:
	s_or_b32 exec_lo, exec_lo, s0
	v_trig_preop_f64 v[22:23], |s[2:3]|, 0
	v_trig_preop_f64 v[20:21], |s[2:3]|, 1
	v_ldexp_f64 v[24:25], |s[2:3]|, 0xffffff80
	v_trig_preop_f64 v[6:7], |s[2:3]|, 2
	v_cmp_class_f64_e64 vcc_lo, s[2:3], 0x1f8
	s_mov_b32 s1, exec_lo
	v_cmpx_gt_f64_e32 0, v[14:15]
	s_cbranch_execz .LBB0_60
; %bb.57:                               ; %.lr.ph.i
	v_cmp_le_f64_e64 s0, 0x7b000000, |s[2:3]|
	s_and_b32 s4, s3, 0x7fffffff
	s_mov_b32 s9, 0x3ff921fb
	s_mov_b32 s8, 0x54442d18
	s_mov_b32 s10, 0x6dc9c883
	s_mov_b32 s11, 0x3fe45f30
	s_mov_b32 s13, 0xbc91a626
	s_mov_b32 s12, 0x33145c00
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v27, s4, v25, s0
	v_cndmask_b32_e64 v26, s2, v24, s0
	s_mov_b32 s4, 0
	v_mul_f64 v[28:29], v[22:23], v[26:27]
	v_mul_f64 v[30:31], v[20:21], v[26:27]
	v_mul_f64 v[42:43], v[6:7], v[26:27]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[32:33], v[22:23], v[26:27], -v[28:29]
	v_fma_f64 v[44:45], v[20:21], v[26:27], -v[30:31]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[26:27], v[6:7], v[26:27], -v[42:43]
	v_add_f64 v[34:35], v[30:31], v[32:33]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[36:37], v[34:35], -v[30:31]
	v_add_f64 v[40:41], v[28:29], v[34:35]
	v_add_f64 v[38:39], v[34:35], -v[36:37]
	v_add_f64 v[32:33], v[32:33], -v[36:37]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_ldexp_f64 v[36:37], v[40:41], -2
	v_add_f64 v[28:29], v[40:41], -v[28:29]
	v_add_f64 v[30:31], v[30:31], -v[38:39]
	v_add_f64 v[38:39], v[42:43], v[44:45]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_neq_f64_e64 s0, 0x7ff00000, |v[36:37]|
	v_add_f64 v[28:29], v[34:35], -v[28:29]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[30:31], v[32:33], v[30:31]
	v_fract_f64_e32 v[32:33], v[36:37]
	v_add_f64 v[34:35], v[38:39], v[30:31]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v33, 0, v33, s0
	v_cndmask_b32_e64 v32, 0, v32, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_ldexp_f64 v[32:33], v[32:33], 2
	v_add_f64 v[36:37], v[28:29], v[34:35]
	v_add_f64 v[46:47], v[34:35], -v[38:39]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[40:41], v[36:37], v[32:33]
	v_add_f64 v[52:53], v[34:35], -v[46:47]
	v_add_f64 v[30:31], v[30:31], -v[46:47]
	v_add_f64 v[28:29], v[36:37], -v[28:29]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cmp_gt_f64_e64 s0, 0, v[40:41]
	v_add_f64 v[40:41], v[38:39], -v[42:43]
	v_add_f64 v[28:29], v[34:35], -v[28:29]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	s_and_b32 s0, s0, exec_lo
	s_cselect_b32 s5, 0x40100000, 0
	v_add_f64 v[50:51], v[38:39], -v[40:41]
	v_add_f64 v[32:33], v[32:33], s[4:5]
	v_add_f64 v[40:41], v[44:45], -v[40:41]
	v_add_f64 v[38:39], v[38:39], -v[52:53]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f64 v[44:45], v[42:43], -v[50:51]
	v_add_f64 v[48:49], v[36:37], v[32:33]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[30:31], v[30:31], v[38:39]
	v_add_f64 v[40:41], v[40:41], v[44:45]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_i32_f64_e32 v48, v[48:49]
	v_add_f64 v[30:31], v[40:41], v[30:31]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cvt_f64_i32_e32 v[46:47], v48
	v_readfirstlane_b32 s5, v48
	v_add_f64 v[26:27], v[26:27], v[30:31]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[32:33], v[32:33], -v[46:47]
	v_add_f64 v[26:27], v[28:29], v[26:27]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[38:39], v[36:37], v[32:33]
	v_add_f64 v[30:31], v[38:39], -v[32:33]
	v_cmp_le_f64_e64 s0, 0.5, v[38:39]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[28:29], v[36:37], -v[30:31]
	s_cmp_lg_u32 s0, 0
	v_mul_f64 v[36:37], |s[2:3]|, s[10:11]
	s_addc_u32 s14, s5, 0
	s_and_b32 s0, s0, exec_lo
	s_cselect_b32 s5, 0x3ff00000, 0
	s_mov_b32 s10, 0x33145c07
	v_add_f64 v[30:31], v[38:39], -s[4:5]
	s_mov_b32 s11, 0x3c91a626
	v_cmp_gt_f64_e64 s0, 0x41d00000, |s[2:3]|
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[26:27], v[26:27], v[28:29]
	v_add_f64 v[28:29], v[30:31], v[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f64 v[32:33], v[28:29], s[8:9]
	v_add_f64 v[30:31], v[28:29], -v[30:31]
	v_fma_f64 v[34:35], v[28:29], s[8:9], -v[32:33]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[26:27], v[26:27], -v[30:31]
	v_rndne_f64_e32 v[30:31], v[36:37]
	v_fma_f64 v[28:29], v[28:29], s[10:11], v[34:35]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[26:27], v[26:27], s[8:9], v[28:29]
	s_mov_b32 s9, 0xbff921fb
	v_mul_f64 v[28:29], v[30:31], s[12:13]
	v_fma_f64 v[34:35], v[30:31], s[8:9], |s[2:3]|
	s_mov_b32 s8, 0x252049c0
	s_mov_b32 s9, 0xb97b839a
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_add_f64 v[36:37], v[34:35], v[28:29]
	v_fma_f64 v[38:39], v[30:31], s[12:13], v[34:35]
	s_mov_b32 s13, s11
	s_mov_b32 s10, 0xf9a43bb8
	s_mov_b32 s11, 0x3de5e0b2
	v_add_f64 v[34:35], v[34:35], -v[36:37]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[36:37], v[36:37], -v[38:39]
	v_add_f64 v[34:35], v[34:35], v[28:29]
	v_fma_f64 v[28:29], v[30:31], s[12:13], v[28:29]
	s_mov_b32 s12, 0x46cc5e42
	s_mov_b32 s13, 0xbda907db
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[34:35], v[36:37], v[34:35]
	v_add_f64 v[28:29], v[34:35], -v[28:29]
	v_add_f64 v[34:35], v[32:33], v[26:27]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[28:29], v[30:31], s[8:9], v[28:29]
	v_add_f64 v[32:33], v[34:35], -v[32:33]
	s_mov_b32 s8, 0xb42fdfa7
	s_mov_b32 s9, 0xbe5ae600
	v_cvt_i32_f64_e32 v30, v[30:31]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[36:37], v[38:39], v[28:29]
	v_add_f64 v[26:27], v[26:27], -v[32:33]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_readfirstlane_b32 s5, v30
	v_cndmask_b32_e64 v33, v35, v37, s0
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v32, v34, v36, s0
	v_add_f64 v[38:39], v[36:37], -v[38:39]
	v_mul_f64 v[34:35], v[32:33], v[32:33]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[28:29], v[28:29], -v[38:39]
	v_fma_f64 v[36:37], v[34:35], s[10:11], s[8:9]
	s_mov_b32 s10, 0x796cde01
	s_mov_b32 s11, 0x3ec71de3
	s_mov_b32 s8, 0x9037ab78
	s_mov_b32 s9, 0x3e21eeb6
	v_mul_f64 v[40:41], v[34:35], 0.5
	v_fma_f64 v[38:39], v[34:35], s[12:13], s[8:9]
	s_mov_b32 s8, 0x19e83e5c
	s_mov_b32 s9, 0xbf2a01a0
	v_mul_f64 v[44:45], v[32:33], -v[34:35]
	v_cndmask_b32_e64 v27, v27, v29, s0
	v_cndmask_b32_e64 v26, v26, v28, s0
	s_and_b32 s0, s0, exec_lo
	s_cselect_b32 s0, s5, s14
	v_fma_f64 v[36:37], v[34:35], v[36:37], s[10:11]
	s_mov_b32 s10, 0xa17f65f6
	s_mov_b32 s11, 0xbe927e4f
	v_mul_f64 v[28:29], v[26:27], 0.5
	s_lshl_b32 s5, s0, 30
	v_add_f64 v[42:43], -v[40:41], 1.0
	v_fma_f64 v[38:39], v[34:35], v[38:39], s[10:11]
	s_mov_b32 s10, 0x11110bb3
	s_mov_b32 s11, 0x3f811111
	s_and_b32 s0, s0, 1
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f64 v[36:37], v[34:35], v[36:37], s[8:9]
	s_mov_b32 s8, 0x19f4ec90
	s_mov_b32 s9, 0x3efa01a0
	v_add_f64 v[46:47], -v[42:43], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f64 v[38:39], v[34:35], v[38:39], s[8:9]
	s_mov_b32 s8, 0x16c16967
	s_mov_b32 s9, 0xbf56c16c
	v_fma_f64 v[36:37], v[34:35], v[36:37], s[10:11]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[28:29], v[44:45], v[36:37], v[28:29]
	v_fma_f64 v[36:37], v[34:35], v[38:39], s[8:9]
	v_add_f64 v[38:39], v[46:47], -v[40:41]
	s_mov_b32 s9, 0x3fa55555
	s_mov_b32 s8, 0x55555555
	v_mul_f64 v[40:41], v[34:35], v[34:35]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[28:29], v[34:35], v[28:29], -v[26:27]
	v_fma_f64 v[34:35], v[34:35], v[36:37], s[8:9]
	s_delay_alu instid0(VALU_DEP_4)
	v_fma_f64 v[26:27], v[32:33], -v[26:27], v[38:39]
	s_mov_b32 s9, 0xbfc55555
	s_delay_alu instid0(VALU_DEP_3) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[28:29], v[44:45], s[8:9], v[28:29]
	s_and_b32 s8, s5, 0x80000000
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[26:27], v[40:41], v[34:35], v[26:27]
	s_cmp_eq_u32 s0, 0
	s_cselect_b32 s0, -1, 0
	s_xor_b32 s5, s5, s3
	s_and_b32 s5, s5, 0x80000000
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[28:29], v[32:33], -v[28:29]
	v_add_f64 v[26:27], v[42:43], v[26:27]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v30, 0x80000000, v29
	v_cndmask_b32_e64 v29, v27, v29, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v31, v28, v26, s0
	v_cndmask_b32_e64 v28, v26, v28, s0
	v_cndmask_b32_e64 v27, v30, v27, s0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_xor_b32_e32 v29, s5, v29
	v_cndmask_b32_e32 v26, 0, v31, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v28, 0, v28, vcc_lo
	v_xor_b32_e32 v27, s8, v27
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v29, 0x7ff80000, v29, vcc_lo
	v_cndmask_b32_e32 v27, 0x7ff80000, v27, vcc_lo
.LBB0_58:                               ; %_ZL3cosd.exit514.i
                                        ; =>This Inner Loop Header: Depth=1
	v_add_f64 v[14:15], v[14:15], s[2:3]
	global_store_b64 v[8:9], v[14:15], off
	global_load_b64 v[14:15], v[16:17], off
	global_load_b64 v[30:31], v[18:19], off
	s_waitcnt vmcnt(1)
	v_mul_f64 v[32:33], v[14:15], v[26:27]
	v_mul_f64 v[14:15], v[14:15], v[28:29]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[32:33], v[30:31], v[28:29], v[32:33]
	v_fma_f64 v[14:15], v[30:31], v[26:27], -v[14:15]
	global_store_b64 v[16:17], v[32:33], off
	global_store_b64 v[18:19], v[14:15], off
	global_load_b64 v[14:15], v[8:9], off
	s_waitcnt vmcnt(0)
	v_cmp_ngt_f64_e64 s0, 0, v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s4, s0, s4
	s_and_not1_b32 exec_lo, exec_lo, s4
	s_cbranch_execnz .LBB0_58
; %bb.59:                               ; %Flow132
	s_or_b32 exec_lo, exec_lo, s4
.LBB0_60:                               ; %Flow133
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	s_mov_b32 s1, exec_lo
	v_cmpx_lt_f64_e32 s[2:3], v[14:15]
	s_cbranch_execz .LBB0_64
; %bb.61:                               ; %.lr.ph560.i
	v_cmp_le_f64_e64 s0, 0x7b000000, |s[2:3]|
	s_and_b32 s4, s3, 0x7fffffff
	s_mov_b32 s9, 0x3ff921fb
	s_mov_b32 s8, 0x54442d18
	s_mov_b32 s10, 0x6dc9c883
	s_mov_b32 s11, 0x3fe45f30
	s_mov_b32 s13, 0xbc91a626
	s_mov_b32 s12, 0x33145c00
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v25, s4, v25, s0
	v_cndmask_b32_e64 v24, s2, v24, s0
	s_mov_b32 s4, 0
	v_mul_f64 v[26:27], v[22:23], v[24:25]
	v_mul_f64 v[28:29], v[20:21], v[24:25]
	v_mul_f64 v[38:39], v[6:7], v[24:25]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[22:23], v[22:23], v[24:25], -v[26:27]
	v_fma_f64 v[20:21], v[20:21], v[24:25], -v[28:29]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[6:7], v[6:7], v[24:25], -v[38:39]
	v_add_f64 v[30:31], v[28:29], v[22:23]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[32:33], v[30:31], -v[28:29]
	v_add_f64 v[36:37], v[26:27], v[30:31]
	v_add_f64 v[34:35], v[30:31], -v[32:33]
	v_add_f64 v[22:23], v[22:23], -v[32:33]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_ldexp_f64 v[32:33], v[36:37], -2
	v_add_f64 v[26:27], v[36:37], -v[26:27]
	v_add_f64 v[28:29], v[28:29], -v[34:35]
	v_add_f64 v[34:35], v[38:39], v[20:21]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_neq_f64_e64 s0, 0x7ff00000, |v[32:33]|
	v_add_f64 v[26:27], v[30:31], -v[26:27]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[22:23], v[22:23], v[28:29]
	v_fract_f64_e32 v[28:29], v[32:33]
	v_add_f64 v[30:31], v[34:35], v[22:23]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v29, 0, v29, s0
	v_cndmask_b32_e64 v28, 0, v28, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_ldexp_f64 v[28:29], v[28:29], 2
	v_add_f64 v[32:33], v[26:27], v[30:31]
	v_add_f64 v[40:41], v[30:31], -v[34:35]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[36:37], v[32:33], v[28:29]
	v_add_f64 v[46:47], v[30:31], -v[40:41]
	v_add_f64 v[22:23], v[22:23], -v[40:41]
	v_add_f64 v[24:25], v[32:33], -v[26:27]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cmp_gt_f64_e64 s0, 0, v[36:37]
	v_add_f64 v[36:37], v[34:35], -v[38:39]
	v_add_f64 v[24:25], v[30:31], -v[24:25]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	s_and_b32 s0, s0, exec_lo
	s_cselect_b32 s5, 0x40100000, 0
	v_add_f64 v[44:45], v[34:35], -v[36:37]
	v_add_f64 v[28:29], v[28:29], s[4:5]
	v_add_f64 v[20:21], v[20:21], -v[36:37]
	v_add_f64 v[34:35], v[34:35], -v[46:47]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f64 v[36:37], v[38:39], -v[44:45]
	v_add_f64 v[42:43], v[32:33], v[28:29]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[22:23], v[22:23], v[34:35]
	v_add_f64 v[20:21], v[20:21], v[36:37]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_i32_f64_e32 v42, v[42:43]
	v_add_f64 v[20:21], v[20:21], v[22:23]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cvt_f64_i32_e32 v[40:41], v42
	v_readfirstlane_b32 s5, v42
	v_add_f64 v[6:7], v[6:7], v[20:21]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[28:29], v[28:29], -v[40:41]
	v_add_f64 v[6:7], v[24:25], v[6:7]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[22:23], v[32:33], v[28:29]
	v_add_f64 v[20:21], v[22:23], -v[28:29]
	v_cmp_le_f64_e64 s0, 0.5, v[22:23]
	v_mul_f64 v[28:29], |s[2:3]|, s[10:11]
	s_mov_b32 s10, 0x33145c07
	s_mov_b32 s11, 0x3c91a626
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[20:21], v[32:33], -v[20:21]
	s_cmp_lg_u32 s0, 0
	s_addc_u32 s14, s5, 0
	s_and_b32 s0, s0, exec_lo
	s_cselect_b32 s5, 0x3ff00000, 0
	v_cmp_gt_f64_e64 s0, 0x41d00000, |s[2:3]|
	v_add_f64 v[22:23], v[22:23], -s[4:5]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[6:7], v[6:7], v[20:21]
	v_add_f64 v[20:21], v[22:23], v[6:7]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f64 v[24:25], v[20:21], s[8:9]
	v_add_f64 v[22:23], v[20:21], -v[22:23]
	v_fma_f64 v[26:27], v[20:21], s[8:9], -v[24:25]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[6:7], v[6:7], -v[22:23]
	v_rndne_f64_e32 v[22:23], v[28:29]
	v_fma_f64 v[20:21], v[20:21], s[10:11], v[26:27]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f64 v[6:7], v[6:7], s[8:9], v[20:21]
	s_mov_b32 s9, 0xbff921fb
	v_mul_f64 v[20:21], v[22:23], s[12:13]
	v_fma_f64 v[26:27], v[22:23], s[8:9], |s[2:3]|
	s_mov_b32 s8, 0x252049c0
	s_mov_b32 s9, 0xb97b839a
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_add_f64 v[28:29], v[26:27], v[20:21]
	v_fma_f64 v[30:31], v[22:23], s[12:13], v[26:27]
	s_mov_b32 s13, s11
	s_mov_b32 s10, 0xf9a43bb8
	s_mov_b32 s11, 0x3de5e0b2
	v_add_f64 v[26:27], v[26:27], -v[28:29]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[28:29], v[28:29], -v[30:31]
	v_add_f64 v[26:27], v[26:27], v[20:21]
	v_fma_f64 v[20:21], v[22:23], s[12:13], v[20:21]
	s_mov_b32 s12, 0x46cc5e42
	s_mov_b32 s13, 0xbda907db
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[26:27], v[28:29], v[26:27]
	v_add_f64 v[20:21], v[26:27], -v[20:21]
	v_add_f64 v[26:27], v[24:25], v[6:7]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[20:21], v[22:23], s[8:9], v[20:21]
	v_add_f64 v[24:25], v[26:27], -v[24:25]
	s_mov_b32 s8, 0xb42fdfa7
	s_mov_b32 s9, 0xbe5ae600
	v_cvt_i32_f64_e32 v22, v[22:23]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[28:29], v[30:31], v[20:21]
	v_add_f64 v[6:7], v[6:7], -v[24:25]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_readfirstlane_b32 s5, v22
	v_cndmask_b32_e64 v25, v27, v29, s0
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v24, v26, v28, s0
	v_add_f64 v[30:31], v[28:29], -v[30:31]
	v_mul_f64 v[26:27], v[24:25], v[24:25]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[20:21], v[20:21], -v[30:31]
	v_fma_f64 v[28:29], v[26:27], s[10:11], s[8:9]
	s_mov_b32 s10, 0x796cde01
	s_mov_b32 s11, 0x3ec71de3
	s_mov_b32 s8, 0x9037ab78
	s_mov_b32 s9, 0x3e21eeb6
	v_mul_f64 v[32:33], v[26:27], 0.5
	v_fma_f64 v[30:31], v[26:27], s[12:13], s[8:9]
	s_mov_b32 s8, 0x19e83e5c
	s_mov_b32 s9, 0xbf2a01a0
	v_mul_f64 v[36:37], v[24:25], -v[26:27]
	v_cndmask_b32_e64 v7, v7, v21, s0
	v_cndmask_b32_e64 v6, v6, v20, s0
	s_and_b32 s0, s0, exec_lo
	s_cselect_b32 s0, s5, s14
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[10:11]
	s_mov_b32 s10, 0xa17f65f6
	s_mov_b32 s11, 0xbe927e4f
	v_mul_f64 v[20:21], v[6:7], 0.5
	s_lshl_b32 s5, s0, 30
	v_add_f64 v[34:35], -v[32:33], 1.0
	v_fma_f64 v[30:31], v[26:27], v[30:31], s[10:11]
	s_mov_b32 s10, 0x11110bb3
	s_mov_b32 s11, 0x3f811111
	s_and_b32 s0, s0, 1
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[8:9]
	s_mov_b32 s8, 0x19f4ec90
	s_mov_b32 s9, 0x3efa01a0
	v_add_f64 v[38:39], -v[34:35], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f64 v[30:31], v[26:27], v[30:31], s[8:9]
	s_mov_b32 s8, 0x16c16967
	s_mov_b32 s9, 0xbf56c16c
	v_fma_f64 v[28:29], v[26:27], v[28:29], s[10:11]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[20:21], v[36:37], v[28:29], v[20:21]
	v_fma_f64 v[28:29], v[26:27], v[30:31], s[8:9]
	v_add_f64 v[30:31], v[38:39], -v[32:33]
	s_mov_b32 s9, 0x3fa55555
	s_mov_b32 s8, 0x55555555
	v_mul_f64 v[32:33], v[26:27], v[26:27]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[20:21], v[26:27], v[20:21], -v[6:7]
	v_fma_f64 v[26:27], v[26:27], v[28:29], s[8:9]
	s_delay_alu instid0(VALU_DEP_4)
	v_fma_f64 v[6:7], v[24:25], -v[6:7], v[30:31]
	s_mov_b32 s9, 0xbfc55555
	s_delay_alu instid0(VALU_DEP_3) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[20:21], v[36:37], s[8:9], v[20:21]
	s_and_b32 s8, s5, 0x80000000
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	v_fma_f64 v[6:7], v[32:33], v[26:27], v[6:7]
	s_cmp_eq_u32 s0, 0
	s_cselect_b32 s0, -1, 0
	s_xor_b32 s5, s5, s3
	s_and_b32 s5, s5, 0x80000000
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[20:21], v[24:25], -v[20:21]
	v_add_f64 v[6:7], v[34:35], v[6:7]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v22, 0x80000000, v21
	v_cndmask_b32_e64 v21, v7, v21, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v23, v20, v6, s0
	v_cndmask_b32_e64 v20, v6, v20, s0
	v_cndmask_b32_e64 v7, v22, v7, s0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_xor_b32_e32 v21, s5, v21
	v_cndmask_b32_e32 v6, 0, v23, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v20, 0, v20, vcc_lo
	v_xor_b32_e32 v7, s8, v7
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v21, 0x7ff80000, v21, vcc_lo
	v_cndmask_b32_e32 v7, 0x7ff80000, v7, vcc_lo
.LBB0_62:                               ; %_ZL3cosd.exit534.i
                                        ; =>This Inner Loop Header: Depth=1
	v_add_f64 v[14:15], v[14:15], -s[2:3]
	global_store_b64 v[8:9], v[14:15], off
	global_load_b64 v[14:15], v[18:19], off
	global_load_b64 v[22:23], v[16:17], off
	s_waitcnt vmcnt(1)
	v_mul_f64 v[24:25], v[14:15], v[20:21]
	s_waitcnt vmcnt(0)
	v_mul_f64 v[26:27], v[22:23], v[20:21]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[22:23], v[22:23], v[6:7], -v[24:25]
	v_fma_f64 v[14:15], v[14:15], v[6:7], v[26:27]
	global_store_b64 v[16:17], v[22:23], off
	global_store_b64 v[18:19], v[14:15], off
	global_load_b64 v[14:15], v[8:9], off
	s_waitcnt vmcnt(0)
	v_cmp_nlt_f64_e32 vcc_lo, s[2:3], v[14:15]
	s_or_b32 s4, vcc_lo, s4
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s4
	s_cbranch_execnz .LBB0_62
; %bb.63:                               ; %Flow130
	s_or_b32 exec_lo, exec_lo, s4
.LBB0_64:                               ; %Flow131
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s1
	global_load_b64 v[6:7], v[2:3], off
	s_mov_b32 s1, exec_lo
                                        ; implicit-def: $vgpr20
                                        ; implicit-def: $vgpr16_vgpr17
                                        ; implicit-def: $vgpr18_vgpr19
	v_cmpx_ngt_f64_e64 0x41d00000, |v[14:15]|
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB0_66
; %bb.65:
	v_ldexp_f64 v[16:17], |v[14:15]|, 0xffffff80
	v_cmp_le_f64_e64 vcc_lo, 0x7b000000, |v[14:15]|
	v_trig_preop_f64 v[18:19], |v[14:15]|, 0
	v_and_b32_e32 v20, 0x7fffffff, v15
	v_trig_preop_f64 v[30:31], |v[14:15]|, 2
	v_mov_b32_e32 v38, 0
	s_mov_b32 s2, 0x54442d18
	s_mov_b32 s3, 0x3ff921fb
	s_mov_b32 s4, 0x33145c07
	s_mov_b32 s5, 0x3c91a626
	v_dual_cndmask_b32 v17, v20, v17 :: v_dual_cndmask_b32 v16, v14, v16
	v_trig_preop_f64 v[20:21], |v[14:15]|, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_mul_f64 v[22:23], v[18:19], v[16:17]
	v_mul_f64 v[36:37], v[30:31], v[16:17]
	v_mul_f64 v[24:25], v[20:21], v[16:17]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[18:19], v[18:19], v[16:17], -v[22:23]
	v_fma_f64 v[20:21], v[20:21], v[16:17], -v[24:25]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[16:17], v[30:31], v[16:17], -v[36:37]
	v_add_f64 v[26:27], v[24:25], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[28:29], v[26:27], -v[24:25]
	v_add_f64 v[34:35], v[22:23], v[26:27]
	v_add_f64 v[32:33], v[26:27], -v[28:29]
	v_add_f64 v[18:19], v[18:19], -v[28:29]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_ldexp_f64 v[28:29], v[34:35], -2
	v_add_f64 v[22:23], v[34:35], -v[22:23]
	v_add_f64 v[24:25], v[24:25], -v[32:33]
	v_add_f64 v[32:33], v[36:37], v[20:21]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_neq_f64_e64 vcc_lo, 0x7ff00000, |v[28:29]|
	v_add_f64 v[22:23], v[26:27], -v[22:23]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[18:19], v[18:19], v[24:25]
	v_fract_f64_e32 v[24:25], v[28:29]
	v_add_f64 v[26:27], v[32:33], v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_cndmask_b32 v25, 0, v25 :: v_dual_cndmask_b32 v24, 0, v24
	v_ldexp_f64 v[24:25], v[24:25], 2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[28:29], v[22:23], v[26:27]
	v_add_f64 v[34:35], v[28:29], v[24:25]
	v_add_f64 v[22:23], v[28:29], -v[22:23]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_cmp_gt_f64_e32 vcc_lo, 0, v[34:35]
	v_add_f64 v[34:35], v[32:33], -v[36:37]
	v_cndmask_b32_e64 v39, 0, 0x40100000, vcc_lo
	v_add_f64 v[43:44], v[32:33], -v[34:35]
	v_add_f64 v[20:21], v[20:21], -v[34:35]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_f64 v[24:25], v[24:25], v[38:39]
	v_add_f64 v[39:40], v[26:27], -v[32:33]
	v_add_f64 v[34:35], v[36:37], -v[43:44]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[41:42], v[28:29], v[24:25]
	v_add_f64 v[45:46], v[26:27], -v[39:40]
	v_add_f64 v[18:19], v[18:19], -v[39:40]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f64 v[20:21], v[20:21], v[34:35]
	v_cvt_i32_f64_e32 v41, v[41:42]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[32:33], v[32:33], -v[45:46]
	v_cvt_f64_i32_e32 v[39:40], v41
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[18:19], v[18:19], v[32:33]
	v_add_f64 v[24:25], v[24:25], -v[39:40]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[18:19], v[20:21], v[18:19]
	v_add_f64 v[20:21], v[26:27], -v[22:23]
	v_add_f64 v[30:31], v[28:29], v[24:25]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[16:17], v[16:17], v[18:19]
	v_add_f64 v[18:19], v[30:31], -v[24:25]
	v_cmp_le_f64_e32 vcc_lo, 0.5, v[30:31]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[16:17], v[20:21], v[16:17]
	v_add_f64 v[18:19], v[28:29], -v[18:19]
	v_cndmask_b32_e64 v39, 0, 0x3ff00000, vcc_lo
	v_add_co_ci_u32_e64 v20, s0, 0, v41, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[16:17], v[16:17], v[18:19]
	v_add_f64 v[18:19], v[30:31], -v[38:39]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[21:22], v[18:19], v[16:17]
	v_mul_f64 v[23:24], v[21:22], s[2:3]
	v_add_f64 v[18:19], v[21:22], -v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[25:26], v[21:22], s[2:3], -v[23:24]
	v_add_f64 v[16:17], v[16:17], -v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[18:19], v[21:22], s[4:5], v[25:26]
	v_fma_f64 v[18:19], v[16:17], s[2:3], v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[16:17], v[23:24], v[18:19]
	v_add_f64 v[21:22], v[16:17], -v[23:24]
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[18:19], v[18:19], -v[21:22]
.LBB0_66:                               ; %Flow129
	s_and_not1_saveexec_b32 s0, s1
	s_cbranch_execz .LBB0_68
; %bb.67:
	s_mov_b32 s2, 0x6dc9c883
	s_mov_b32 s3, 0x3fe45f30
	s_mov_b32 s5, 0xbc91a626
	v_mul_f64 v[16:17], |v[14:15]|, s[2:3]
	s_mov_b32 s2, 0x54442d18
	s_mov_b32 s3, 0xbff921fb
	s_mov_b32 s4, 0x33145c00
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rndne_f64_e32 v[20:21], v[16:17]
	v_fma_f64 v[16:17], v[20:21], s[2:3], |v[14:15]|
	v_mul_f64 v[18:19], v[20:21], s[4:5]
	s_mov_b32 s2, 0x252049c0
	s_mov_b32 s3, 0xb97b839a
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[24:25], v[20:21], s[4:5], v[16:17]
	v_add_f64 v[22:23], v[16:17], v[18:19]
	s_mov_b32 s5, 0x3c91a626
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[16:17], v[16:17], -v[22:23]
	v_add_f64 v[22:23], v[22:23], -v[24:25]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[16:17], v[16:17], v[18:19]
	v_fma_f64 v[18:19], v[20:21], s[4:5], v[18:19]
	v_add_f64 v[16:17], v[22:23], v[16:17]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[16:17], v[16:17], -v[18:19]
	v_fma_f64 v[18:19], v[20:21], s[2:3], v[16:17]
	v_cvt_i32_f64_e32 v20, v[20:21]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[16:17], v[24:25], v[18:19]
	v_add_f64 v[22:23], v[16:17], -v[24:25]
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[18:19], v[18:19], -v[22:23]
.LBB0_68:                               ; %_ZL3cosd.exit554.i
	s_or_b32 exec_lo, exec_lo, s0
	v_mul_f64 v[21:22], v[16:17], v[16:17]
	s_mov_b32 s0, 0xb42fdfa7
	s_mov_b32 s2, 0xf9a43bb8
	s_mov_b32 s1, 0xbe5ae600
	s_mov_b32 s3, 0x3de5e0b2
	s_mov_b32 s4, 0x796cde01
	s_mov_b32 s5, 0x3ec71de3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[31:32], v[18:19], 0.5
	v_fma_f64 v[23:24], v[21:22], s[2:3], s[0:1]
	s_mov_b32 s0, 0x9037ab78
	s_mov_b32 s2, 0x46cc5e42
	s_mov_b32 s1, 0x3e21eeb6
	s_mov_b32 s3, 0xbda907db
	v_mul_f64 v[27:28], v[21:22], 0.5
	v_fma_f64 v[25:26], v[21:22], s[2:3], s[0:1]
	s_mov_b32 s0, 0xa17f65f6
	s_mov_b32 s2, 0x19e83e5c
	s_mov_b32 s1, 0xbe927e4f
	s_mov_b32 s3, 0xbf2a01a0
	v_mul_f64 v[33:34], v[16:17], -v[21:22]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[23:24], v[21:22], v[23:24], s[4:5]
	v_add_f64 v[29:30], -v[27:28], 1.0
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f64 v[25:26], v[21:22], v[25:26], s[0:1]
	s_mov_b32 s0, 0x19f4ec90
	s_mov_b32 s1, 0x3efa01a0
	v_fma_f64 v[23:24], v[21:22], v[23:24], s[2:3]
	s_mov_b32 s2, 0x11110bb3
	s_mov_b32 s3, 0x3f811111
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[35:36], -v[29:30], 1.0
	v_fma_f64 v[25:26], v[21:22], v[25:26], s[0:1]
	s_mov_b32 s0, 0x16c16967
	s_mov_b32 s1, 0xbf56c16c
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[23:24], v[21:22], v[23:24], s[2:3]
	v_add_f64 v[27:28], v[35:36], -v[27:28]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f64 v[25:26], v[21:22], v[25:26], s[0:1]
	s_mov_b32 s1, 0x3fa55555
	s_mov_b32 s0, 0x55555555
	v_fma_f64 v[23:24], v[33:34], v[23:24], v[31:32]
	v_mul_f64 v[31:32], v[21:22], v[21:22]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[27:28], v[16:17], -v[18:19], v[27:28]
	v_fma_f64 v[25:26], v[21:22], v[25:26], s[0:1]
	s_mov_b32 s1, 0xbfc55555
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[18:19], v[21:22], v[23:24], -v[18:19]
	v_fma_f64 v[21:22], v[31:32], v[25:26], v[27:28]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_fma_f64 v[18:19], v[33:34], s[0:1], v[18:19]
	v_cmp_class_f64_e64 s0, v[14:15], 0x1f8
	s_mov_b32 s1, exec_lo
	v_lshlrev_b32_e32 v15, 30, v20
	v_add_f64 v[21:22], v[29:30], v[21:22]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_and_b32_e32 v15, 0x80000000, v15
	v_add_f64 v[16:17], v[16:17], -v[18:19]
	v_and_b32_e32 v18, 1, v20
	v_cmp_eq_u32_e32 vcc_lo, 0, v18
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v16, v16, v21, vcc_lo
	v_xor_b32_e32 v14, 0x80000000, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v14, v14, v22, vcc_lo
	v_xor_b32_e32 v15, v14, v15
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v14, 0, v16, s0
	v_cndmask_b32_e64 v15, 0x7ff80000, v15, s0
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f64 v[6:7], v[6:7], v[14:15]
                                        ; implicit-def: $vgpr14
	global_store_b64 v[12:13], v[6:7], off
	global_load_b64 v[6:7], v[8:9], off
	global_load_b64 v[2:3], v[2:3], off
                                        ; implicit-def: $vgpr8_vgpr9
                                        ; implicit-def: $vgpr12_vgpr13
	s_waitcnt vmcnt(1)
	v_cmpx_ngt_f64_e64 0x41d00000, |v[6:7]|
	s_xor_b32 s1, exec_lo, s1
	s_cbranch_execz .LBB0_70
; %bb.69:
	v_ldexp_f64 v[8:9], |v[6:7]|, 0xffffff80
	v_cmp_le_f64_e64 vcc_lo, 0x7b000000, |v[6:7]|
	v_trig_preop_f64 v[12:13], |v[6:7]|, 0
	v_and_b32_e32 v14, 0x7fffffff, v7
	v_trig_preop_f64 v[24:25], |v[6:7]|, 2
	v_mov_b32_e32 v32, 0
	s_mov_b32 s2, 0x54442d18
	s_mov_b32 s3, 0x3ff921fb
	s_mov_b32 s4, 0x33145c07
	s_mov_b32 s5, 0x3c91a626
	v_cndmask_b32_e32 v9, v14, v9, vcc_lo
	v_cndmask_b32_e32 v8, v6, v8, vcc_lo
	v_trig_preop_f64 v[14:15], |v[6:7]|, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[16:17], v[12:13], v[8:9]
	v_mul_f64 v[18:19], v[14:15], v[8:9]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[12:13], v[12:13], v[8:9], -v[16:17]
	v_fma_f64 v[14:15], v[14:15], v[8:9], -v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[20:21], v[18:19], v[12:13]
	v_add_f64 v[22:23], v[20:21], -v[18:19]
	v_add_f64 v[28:29], v[16:17], v[20:21]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[26:27], v[20:21], -v[22:23]
	v_add_f64 v[12:13], v[12:13], -v[22:23]
	v_ldexp_f64 v[22:23], v[28:29], -2
	v_add_f64 v[16:17], v[28:29], -v[16:17]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[18:19], v[18:19], -v[26:27]
	v_cmp_neq_f64_e64 vcc_lo, 0x7ff00000, |v[22:23]|
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[16:17], v[20:21], -v[16:17]
	v_add_f64 v[12:13], v[12:13], v[18:19]
	v_fract_f64_e32 v[18:19], v[22:23]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_cndmask_b32 v19, 0, v19 :: v_dual_cndmask_b32 v18, 0, v18
	v_mul_f64 v[30:31], v[24:25], v[8:9]
	v_ldexp_f64 v[18:19], v[18:19], 2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[26:27], v[30:31], v[14:15]
	v_fma_f64 v[8:9], v[24:25], v[8:9], -v[30:31]
	v_add_f64 v[20:21], v[26:27], v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[22:23], v[16:17], v[20:21]
	v_add_f64 v[28:29], v[22:23], v[18:19]
	v_add_f64 v[16:17], v[22:23], -v[16:17]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_cmp_gt_f64_e32 vcc_lo, 0, v[28:29]
	v_add_f64 v[28:29], v[26:27], -v[30:31]
	v_cndmask_b32_e64 v33, 0, 0x40100000, vcc_lo
	v_add_f64 v[37:38], v[26:27], -v[28:29]
	v_add_f64 v[14:15], v[14:15], -v[28:29]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_f64 v[18:19], v[18:19], v[32:33]
	v_add_f64 v[33:34], v[20:21], -v[26:27]
	v_add_f64 v[28:29], v[30:31], -v[37:38]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[35:36], v[22:23], v[18:19]
	v_add_f64 v[39:40], v[20:21], -v[33:34]
	v_add_f64 v[12:13], v[12:13], -v[33:34]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f64 v[14:15], v[14:15], v[28:29]
	v_cvt_i32_f64_e32 v35, v[35:36]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[26:27], v[26:27], -v[39:40]
	v_cvt_f64_i32_e32 v[33:34], v35
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[12:13], v[12:13], v[26:27]
	v_add_f64 v[18:19], v[18:19], -v[33:34]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f64 v[12:13], v[14:15], v[12:13]
	v_add_f64 v[14:15], v[20:21], -v[16:17]
	v_add_f64 v[24:25], v[22:23], v[18:19]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[8:9], v[8:9], v[12:13]
	v_add_f64 v[12:13], v[24:25], -v[18:19]
	v_cmp_le_f64_e32 vcc_lo, 0.5, v[24:25]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[8:9], v[14:15], v[8:9]
	v_add_f64 v[12:13], v[22:23], -v[12:13]
	v_cndmask_b32_e64 v33, 0, 0x3ff00000, vcc_lo
	v_add_co_ci_u32_e64 v14, s0, 0, v35, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[8:9], v[8:9], v[12:13]
	v_add_f64 v[12:13], v[24:25], -v[32:33]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[15:16], v[12:13], v[8:9]
	v_mul_f64 v[17:18], v[15:16], s[2:3]
	v_add_f64 v[12:13], v[15:16], -v[12:13]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[19:20], v[15:16], s[2:3], -v[17:18]
	v_add_f64 v[8:9], v[8:9], -v[12:13]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[12:13], v[15:16], s[4:5], v[19:20]
	v_fma_f64 v[12:13], v[8:9], s[2:3], v[12:13]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[8:9], v[17:18], v[12:13]
	v_add_f64 v[15:16], v[8:9], -v[17:18]
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[12:13], v[12:13], -v[15:16]
.LBB0_70:                               ; %Flow128
	s_and_not1_saveexec_b32 s0, s1
	s_cbranch_execz .LBB0_72
; %bb.71:
	s_mov_b32 s2, 0x6dc9c883
	s_mov_b32 s3, 0x3fe45f30
	s_mov_b32 s5, 0xbc91a626
	v_mul_f64 v[8:9], |v[6:7]|, s[2:3]
	s_mov_b32 s2, 0x54442d18
	s_mov_b32 s3, 0xbff921fb
	s_mov_b32 s4, 0x33145c00
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rndne_f64_e32 v[14:15], v[8:9]
	v_fma_f64 v[8:9], v[14:15], s[2:3], |v[6:7]|
	v_mul_f64 v[12:13], v[14:15], s[4:5]
	s_mov_b32 s2, 0x252049c0
	s_mov_b32 s3, 0xb97b839a
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[18:19], v[14:15], s[4:5], v[8:9]
	v_add_f64 v[16:17], v[8:9], v[12:13]
	s_mov_b32 s5, 0x3c91a626
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[8:9], v[8:9], -v[16:17]
	v_add_f64 v[16:17], v[16:17], -v[18:19]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f64 v[8:9], v[8:9], v[12:13]
	v_fma_f64 v[12:13], v[14:15], s[4:5], v[12:13]
	v_add_f64 v[8:9], v[16:17], v[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[8:9], v[8:9], -v[12:13]
	v_fma_f64 v[12:13], v[14:15], s[2:3], v[8:9]
	v_cvt_i32_f64_e32 v14, v[14:15]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[8:9], v[18:19], v[12:13]
	v_add_f64 v[16:17], v[8:9], -v[18:19]
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f64 v[12:13], v[12:13], -v[16:17]
.LBB0_72:                               ; %_ZL3sind.exit559.i
	s_or_b32 exec_lo, exec_lo, s0
	v_mul_f64 v[15:16], v[8:9], v[8:9]
	s_mov_b32 s0, 0xb42fdfa7
	s_mov_b32 s2, 0xf9a43bb8
	s_mov_b32 s1, 0xbe5ae600
	s_mov_b32 s3, 0x3de5e0b2
	s_mov_b32 s4, 0x796cde01
	s_mov_b32 s5, 0x3ec71de3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[25:26], v[12:13], 0.5
	v_fma_f64 v[17:18], v[15:16], s[2:3], s[0:1]
	s_mov_b32 s0, 0x9037ab78
	s_mov_b32 s2, 0x46cc5e42
	s_mov_b32 s1, 0x3e21eeb6
	s_mov_b32 s3, 0xbda907db
	v_mul_f64 v[21:22], v[15:16], 0.5
	v_fma_f64 v[19:20], v[15:16], s[2:3], s[0:1]
	s_mov_b32 s0, 0xa17f65f6
	s_mov_b32 s2, 0x19e83e5c
	s_mov_b32 s1, 0xbe927e4f
	s_mov_b32 s3, 0xbf2a01a0
	v_mul_f64 v[27:28], v[8:9], -v[15:16]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[17:18], v[15:16], v[17:18], s[4:5]
	v_add_f64 v[23:24], -v[21:22], 1.0
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f64 v[19:20], v[15:16], v[19:20], s[0:1]
	s_mov_b32 s0, 0x19f4ec90
	s_mov_b32 s1, 0x3efa01a0
	v_fma_f64 v[17:18], v[15:16], v[17:18], s[2:3]
	s_mov_b32 s2, 0x11110bb3
	s_mov_b32 s3, 0x3f811111
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f64 v[29:30], -v[23:24], 1.0
	v_fma_f64 v[19:20], v[15:16], v[19:20], s[0:1]
	s_mov_b32 s0, 0x16c16967
	s_mov_b32 s1, 0xbf56c16c
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[17:18], v[15:16], v[17:18], s[2:3]
	v_add_f64 v[21:22], v[29:30], -v[21:22]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f64 v[19:20], v[15:16], v[19:20], s[0:1]
	s_mov_b32 s1, 0x3fa55555
	s_mov_b32 s0, 0x55555555
	v_fma_f64 v[17:18], v[27:28], v[17:18], v[25:26]
	v_mul_f64 v[25:26], v[15:16], v[15:16]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f64 v[21:22], v[8:9], -v[12:13], v[21:22]
	v_fma_f64 v[19:20], v[15:16], v[19:20], s[0:1]
	s_mov_b32 s1, 0xbfc55555
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[12:13], v[15:16], v[17:18], -v[12:13]
	v_fma_f64 v[15:16], v[25:26], v[19:20], v[21:22]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_fma_f64 v[12:13], v[27:28], s[0:1], v[12:13]
	v_cmp_class_f64_e64 s0, v[6:7], 0x1f8
	v_lshlrev_b32_e32 v6, 30, v14
	v_add_f64 v[15:16], v[23:24], v[15:16]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v6, v6, v7
	v_and_b32_e32 v6, 0x80000000, v6
	v_add_f64 v[8:9], v[8:9], -v[12:13]
	v_and_b32_e32 v12, 1, v14
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_eq_u32_e32 vcc_lo, 0, v12
	v_dual_cndmask_b32 v7, v16, v9 :: v_dual_cndmask_b32 v8, v15, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v7, v7, v6
	v_cndmask_b32_e64 v6, 0, v8, s0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v7, 0x7ff80000, v7, s0
	s_mov_b32 s0, exec_lo
	s_waitcnt vmcnt(0)
	v_mul_f64 v[2:3], v[2:3], v[6:7]
	global_store_b64 v[10:11], v[2:3], off
	global_load_b64 v[2:3], v[0:1], off
	s_waitcnt vmcnt(0)
	v_cmpx_lt_f64_e32 s[6:7], v[2:3]
	s_cbranch_execz .LBB0_74
; %bb.73:
	v_fma_f64 v[2:3], s[6:7], 2.0, -v[2:3]
	global_store_b64 v[0:1], v[2:3], off
	global_load_b64 v[2:3], v[4:5], off
	s_waitcnt vmcnt(0)
	v_xor_b32_e32 v3, 0x80000000, v3
	global_store_b64 v[4:5], v[2:3], off
	global_load_b64 v[2:3], v[0:1], off
.LBB0_74:
	s_or_b32 exec_lo, exec_lo, s0
	s_waitcnt vmcnt(0)
	v_cmp_gt_f64_e32 vcc_lo, 0, v[2:3]
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_76
; %bb.75:
	v_xor_b32_e32 v3, 0x80000000, v3
	global_store_b64 v[0:1], v[2:3], off
	global_load_b64 v[0:1], v[4:5], off
	s_waitcnt vmcnt(0)
	v_xor_b32_e32 v1, 0x80000000, v1
	global_store_b64 v[4:5], v[0:1], off
.LBB0_76:                               ; %_Z3EM1PdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd.exit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 464
		.amdhsa_user_sgpr_count 14
		.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 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 88
		.amdhsa_next_free_sgpr 52
		.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	_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd, .Lfunc_end0-_Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 16512
; NumSgprs: 54
; NumVgprs: 88
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 6
; VGPRBlocks: 10
; NumSGPRsForWavesPerEU: 54
; NumVGPRsForWavesPerEU: 88
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         80
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         88
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         96
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         104
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         112
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         120
        .size:           8
        .value_kind:     global_buffer
      - .offset:         128
        .size:           4
        .value_kind:     by_value
      - .offset:         136
        .size:           8
        .value_kind:     by_value
      - .offset:         144
        .size:           8
        .value_kind:     by_value
      - .offset:         152
        .size:           8
        .value_kind:     by_value
      - .offset:         160
        .size:           4
        .value_kind:     by_value
      - .offset:         164
        .size:           4
        .value_kind:     by_value
      - .offset:         168
        .size:           8
        .value_kind:     by_value
      - .offset:         176
        .size:           8
        .value_kind:     by_value
      - .offset:         184
        .size:           8
        .value_kind:     by_value
      - .offset:         192
        .size:           8
        .value_kind:     by_value
      - .offset:         200
        .size:           8
        .value_kind:     by_value
      - .offset:         208
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         212
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         216
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         220
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         222
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         224
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         226
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         228
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         230
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         248
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         256
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         264
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         272
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 464
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd
    .private_segment_fixed_size: 0
    .sgpr_count:     54
    .sgpr_spill_count: 0
    .symbol:         _Z24processMandelbrotElementPdS_S_S_S_S_S_S_S_S_S_S_S_S_S_Piidddiiddddd.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     88
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
