// #include <pycuda-helpers.hpp>
__global__ void copyDensity( double *cnsv, double *rho ){
	int t_j = blockIdx.x*blockDim.x + threadIdx.x;
  int t_i = blockIdx.y*blockDim.y + threadIdx.y;
  int t_k = blockIdx.z*blockDim.z + threadIdx.z;
  int tid = t_j + t_i*blockDim.x*gridDim.x + t_k*blockDim.x*gridDim.x*blockDim.y*gridDim.y;

	rho[tid] = cnsv[tid];
}

__global__ void FFT_divideK2( const double pi4,
							double *kxfft, double *kyfft, double *kzfft,
				      double *data_re, double *data_im){
  int t_j = blockIdx.x*blockDim.x + threadIdx.x;
  int t_i = blockIdx.y*blockDim.y + threadIdx.y;
  int t_k = blockIdx.z*blockDim.z + threadIdx.z;
  int tid = t_j + t_i*blockDim.x*gridDim.x + t_k*blockDim.x*gridDim.x*blockDim.y*gridDim.y;

  double kx = kxfft[t_j];
  double ky = kyfft[t_i];
  double kz = kzfft[t_k];
  double k2 = kx*kx + ky*ky + kz*kz;
	if ( abs( k2 ) < 1e-5 ) k2 = 1 ;
  data_re[tid] = -data_re[tid]*pi4/k2;
	data_im[tid] = -data_im[tid]*pi4/k2;
}

__global__ void iterPoissonStep(  const int parity,
				 const int nX, const int nY, const int nZ,
         double dx, double dy, double dz,
         const double omega, const double pi4, const double convEpsilon,
				 double* rho_all, double* phi_all, int* converged ){
  int t_j = 2*(blockIdx.x*blockDim.x + threadIdx.x);
  int t_i = blockIdx.y*blockDim.y + threadIdx.y;
  int t_k = blockIdx.z*blockDim.z + threadIdx.z;
  //Make a checkboard 3D grid
  if ( t_i%2 == 0 ){
    if ( t_k%2 == parity ) t_j +=1;
  }
  else if ( (t_k+1)%2 == parity ) t_j +=1;
  int tid = t_j + t_i*nX + t_k*nX*blockDim.y*gridDim.y;

	//Set neighbors ids
	int l_indx, r_indx, d_indx, u_indx, b_indx, t_indx;
	l_indx = t_j==0    ?    1 : t_j-1;  //Left
	r_indx = t_j==nX-1 ? nX-2 : t_j+1;  //Right
	d_indx = t_i==0    ?    1 : t_i-1;  //Down
	u_indx = t_i==nY-1 ? nY-2 : t_i+1;  //Up
	b_indx = t_k==0    ?    1 : t_k-1;  //bottom
	t_indx = t_k==nZ-1 ? nZ-2 : t_k+1;  //top

  double rho, phi_c, phi_l, phi_r, phi_d, phi_u, phi_b, phi_t, phi_new;
	rho = rho_all[ tid ];
	phi_c = phi_all[tid];
	phi_l = phi_all[ l_indx + t_i*nX + t_k*nX*blockDim.y*gridDim.y ];
	phi_r = phi_all[ r_indx + t_i*nX + t_k*nX*blockDim.y*gridDim.y ];
	phi_d = phi_all[ t_j + d_indx*nX + t_k*nX*blockDim.y*gridDim.y ];
	phi_u = phi_all[ t_j + u_indx*nX + t_k*nX*blockDim.y*gridDim.y ];
	phi_b = phi_all[ t_j + t_i*nX + b_indx*nX*blockDim.y*gridDim.y ];
	phi_t = phi_all[ t_j + t_i*nX + t_indx*nX*blockDim.y*gridDim.y ];

  phi_new = (1-omega)*phi_c + omega/6*( phi_l + phi_r + phi_d + phi_u + phi_b + phi_t - dx*dx*pi4*rho );
  phi_all[ tid ] = phi_new;


  if ( ( abs( ( phi_new - phi_c ) / phi_c ) > convEpsilon ) ) converged[0] = 0;
	// if ( ( abs( phi_new - phi_c ) > 0.0001 ) ) converged[0] = 0;
}

__global__ void getGravityForce(const int nCells,  const int nW, const int nH, const int nD,
				 double dx, double dy, double dz,
				 double* gForce_x, double* gForce_y, double* gForce_z,
				 double* cnsv, double *phi_all, double *gravWork ){
  int t_j = blockIdx.x*blockDim.x + threadIdx.x;
  int t_i = blockIdx.y*blockDim.y + threadIdx.y;
  int t_k = blockIdx.z*blockDim.z + threadIdx.z;
  int tid = t_j + t_i*blockDim.x*gridDim.x + t_k*blockDim.x*gridDim.x*blockDim.y*gridDim.y;

	//Set neighbors ids
	int l_indx, r_indx, d_indx, u_indx, b_indx, t_indx;
	l_indx = t_j==0    ?    1 : t_j-1;  //Left
	r_indx = t_j==nW-1 ? nW-2 : t_j+1;  //Right
	d_indx = t_i==0    ?    1 : t_i-1;  //Down
	u_indx = t_i==nH-1 ? nH-2 : t_i+1;  //Up
	b_indx = t_k==0    ?    1 : t_k-1;  //bottom
	t_indx = t_k==nD-1 ? nD-2 : t_k+1;  //top

  double phi_l, phi_r, phi_d, phi_u, phi_b, phi_t;
	phi_l = phi_all[ l_indx + t_i*nW + t_k*nW*nH ];
	phi_r = phi_all[ r_indx + t_i*nW + t_k*nW*nH ];
	phi_d = phi_all[ t_j + d_indx*nW + t_k*nW*nH ];
	phi_u = phi_all[ t_j + u_indx*nW + t_k*nW*nH ];
	phi_b = phi_all[ t_j + t_i*nW + b_indx*nW*nH ];
	phi_t = phi_all[ t_j + t_i*nW + t_indx*nW*nH ];

  // //Boundary conditions
  // if  ( t_j == 0 )        phi_l = phi_r;
  // if  ( t_j == nWidth-1 ) phi_r = phi_l;
  // if  ( t_i == 0 )        phi_d = phi_u;
  // if  ( t_i == nWidth-1 ) phi_u = phi_d;
  // if  ( t_k == 0 )        phi_b = phi_t;
  // if  ( t_k == nWidth-1 ) phi_t = phi_b;

  //Get partial derivatives for force
  double gField_x, gField_y, gField_z;
  gField_x = ( phi_l - phi_r ) * 0.5 / dx;
  gField_y = ( phi_d - phi_u ) * 0.5 / dy;
  gField_z = ( phi_b - phi_t ) * 0.5 / dz;
	double p_x, p_y, p_z, rho;
	rho = cnsv[ 0*nCells + tid ];
	const int factor = 1;
  gForce_x[ tid ] = gField_x * rho * factor;
  gForce_y[ tid ] = gField_y * rho * factor;
  gForce_z[ tid ] = gField_z * rho * factor;
  // gForce_x[ tid ] = gField_x;
  // gForce_y[ tid ] = gField_y;
  // gForce_z[ tid ] = gField_z;

  // //Get momentum for virtual gravitational work
	p_x = cnsv[ 1*nCells + tid ];
	p_y = cnsv[ 2*nCells + tid ];
	p_z = cnsv[ 3*nCells + tid ];
  gravWork[ tid ] = (p_x * gField_x + p_y * gField_y + p_z * gField_z )*factor;

}

	.file	"tmpxft_003309d0_00000000-6_cuda_gravity.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z33__device_stub__Z11copyDensityPdS_PdS_
	.type	_Z33__device_stub__Z11copyDensityPdS_PdS_, @function
_Z33__device_stub__Z11copyDensityPdS_PdS_:
.LFB2032:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z11copyDensityPdS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L2:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2032:
	.size	_Z33__device_stub__Z11copyDensityPdS_PdS_, .-_Z33__device_stub__Z11copyDensityPdS_PdS_
	.globl	_Z11copyDensityPdS_
	.type	_Z11copyDensityPdS_, @function
_Z11copyDensityPdS_:
.LFB2033:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z11copyDensityPdS_PdS_
	.cfi_endproc
.LFE2033:
	.size	_Z11copyDensityPdS_, .-_Z11copyDensityPdS_
	.globl	_Z41__device_stub__Z12FFT_divideK2dPdS_S_S_S_dPdS_S_S_S_
	.type	_Z41__device_stub__Z12FFT_divideK2dPdS_S_S_S_dPdS_S_S_S_, @function
_Z41__device_stub__Z12FFT_divideK2dPdS_S_S_S_dPdS_S_S_S_:
.LFB2034:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 32(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 24(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, (%rsp)
	movsd	%xmm0, 40(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 152(%rsp)
	movq	%rsp, %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z12FFT_divideK2dPdS_S_S_S_(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L8:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2034:
	.size	_Z41__device_stub__Z12FFT_divideK2dPdS_S_S_S_dPdS_S_S_S_, .-_Z41__device_stub__Z12FFT_divideK2dPdS_S_S_S_dPdS_S_S_S_
	.globl	_Z12FFT_divideK2dPdS_S_S_S_
	.type	_Z12FFT_divideK2dPdS_S_S_S_, @function
_Z12FFT_divideK2dPdS_S_S_S_:
.LFB2035:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z12FFT_divideK2dPdS_S_S_S_dPdS_S_S_S_
	.cfi_endproc
.LFE2035:
	.size	_Z12FFT_divideK2dPdS_S_S_S_, .-_Z12FFT_divideK2dPdS_S_S_S_
	.globl	_Z49__device_stub__Z15iterPoissonStepiiiiddddddPdS_PiiiiiddddddPdS_Pi
	.type	_Z49__device_stub__Z15iterPoissonStepiiiiddddddPdS_PiiiiiddddddPdS_Pi, @function
_Z49__device_stub__Z15iterPoissonStepiiiiddddddPdS_PiiiiiddddddPdS_Pi:
.LFB2036:
	.cfi_startproc
	endbr64
	subq	$280, %rsp
	.cfi_def_cfa_offset 288
	movq	288(%rsp), %rax
	movl	%edi, 92(%rsp)
	leaq	112(%rsp), %rdi
	movl	%esi, 88(%rsp)
	leaq	124(%rsp), %rsi
	movl	%edx, 84(%rsp)
	leaq	96(%rsp), %rdx
	movl	%ecx, 80(%rsp)
	leaq	104(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movsd	%xmm0, 72(%rsp)
	movsd	%xmm1, 64(%rsp)
	movsd	%xmm2, 56(%rsp)
	movsd	%xmm3, 48(%rsp)
	movsd	%xmm4, 40(%rsp)
	movsd	%xmm5, 32(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 264(%rsp)
	xorl	%eax, %eax
	leaq	92(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	84(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 256(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 112(%rsp)
	movl	$1, 120(%rsp)
	movq	%rax, 124(%rsp)
	movl	$1, 132(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	104(%rsp)
	.cfi_def_cfa_offset 296
	leaq	_Z15iterPoissonStepiiiiddddddPdS_Pi(%rip), %rdi
	pushq	104(%rsp)
	.cfi_def_cfa_offset 304
	movq	140(%rsp), %rcx
	movl	148(%rsp), %r8d
	movq	128(%rsp), %rsi
	movl	136(%rsp), %edx
	leaq	176(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 296
	popq	%rdx
	.cfi_def_cfa_offset 288
.L13:
	movq	264(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$280, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2036:
	.size	_Z49__device_stub__Z15iterPoissonStepiiiiddddddPdS_PiiiiiddddddPdS_Pi, .-_Z49__device_stub__Z15iterPoissonStepiiiiddddddPdS_PiiiiiddddddPdS_Pi
	.globl	_Z15iterPoissonStepiiiiddddddPdS_Pi
	.type	_Z15iterPoissonStepiiiiddddddPdS_Pi, @function
_Z15iterPoissonStepiiiiddddddPdS_Pi:
.LFB2037:
	.cfi_startproc
	endbr64
	jmp	_Z49__device_stub__Z15iterPoissonStepiiiiddddddPdS_PiiiiiddddddPdS_Pi
	.cfi_endproc
.LFE2037:
	.size	_Z15iterPoissonStepiiiiddddddPdS_Pi, .-_Z15iterPoissonStepiiiiddddddPdS_Pi
	.globl	_Z52__device_stub__Z15getGravityForceiiiidddPdS_S_S_S_S_iiiidddPdS_S_S_S_S_
	.type	_Z52__device_stub__Z15getGravityForceiiiidddPdS_S_S_S_S_iiiidddPdS_S_S_S_S_, @function
_Z52__device_stub__Z15getGravityForceiiiidddPdS_S_S_S_S_iiiidddPdS_S_S_S_S_:
.LFB2038:
	.cfi_startproc
	endbr64
	subq	$280, %rsp
	.cfi_def_cfa_offset 288
	movq	288(%rsp), %rax
	movl	%edi, 92(%rsp)
	leaq	112(%rsp), %rdi
	movl	%esi, 88(%rsp)
	leaq	124(%rsp), %rsi
	movq	%rax, 32(%rsp)
	movq	296(%rsp), %rax
	movl	%edx, 84(%rsp)
	leaq	96(%rsp), %rdx
	movq	%rax, 24(%rsp)
	movq	304(%rsp), %rax
	movl	%ecx, 80(%rsp)
	leaq	104(%rsp), %rcx
	movq	%rax, 16(%rsp)
	movq	312(%rsp), %rax
	movq	%r8, 48(%rsp)
	movq	%r9, 40(%rsp)
	movq	%rax, 8(%rsp)
	movsd	%xmm0, 72(%rsp)
	movsd	%xmm1, 64(%rsp)
	movsd	%xmm2, 56(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 264(%rsp)
	xorl	%eax, %eax
	leaq	92(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	84(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 256(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 112(%rsp)
	movl	$1, 120(%rsp)
	movq	%rax, 124(%rsp)
	movl	$1, 132(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L18
	pushq	104(%rsp)
	.cfi_def_cfa_offset 296
	leaq	_Z15getGravityForceiiiidddPdS_S_S_S_S_(%rip), %rdi
	pushq	104(%rsp)
	.cfi_def_cfa_offset 304
	movq	140(%rsp), %rcx
	movl	148(%rsp), %r8d
	movq	128(%rsp), %rsi
	movl	136(%rsp), %edx
	leaq	176(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 296
	popq	%rdx
	.cfi_def_cfa_offset 288
.L18:
	movq	264(%rsp), %rax
	subq	%fs:40, %rax
	je	.L20
	call	__stack_chk_fail@PLT
.L20:
	addq	$280, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2038:
	.size	_Z52__device_stub__Z15getGravityForceiiiidddPdS_S_S_S_S_iiiidddPdS_S_S_S_S_, .-_Z52__device_stub__Z15getGravityForceiiiidddPdS_S_S_S_S_iiiidddPdS_S_S_S_S_
	.globl	_Z15getGravityForceiiiidddPdS_S_S_S_S_
	.type	_Z15getGravityForceiiiidddPdS_S_S_S_S_, @function
_Z15getGravityForceiiiidddPdS_S_S_S_S_:
.LFB2039:
	.cfi_startproc
	endbr64
	jmp	_Z52__device_stub__Z15getGravityForceiiiidddPdS_S_S_S_S_iiiidddPdS_S_S_S_S_
	.cfi_endproc
.LFE2039:
	.size	_Z15getGravityForceiiiidddPdS_S_S_S_S_, .-_Z15getGravityForceiiiidddPdS_S_S_S_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z15getGravityForceiiiidddPdS_S_S_S_S_"
.LC1:
	.string	"_Z15iterPoissonStepiiiiddddddPdS_Pi"
.LC2:
	.string	"_Z12FFT_divideK2dPdS_S_S_S_"
.LC3:
	.string	"_Z11copyDensityPdS_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2041:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z15getGravityForceiiiidddPdS_S_S_S_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z15iterPoissonStepiiiiddddddPdS_Pi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z12FFT_divideK2dPdS_S_S_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z11copyDensityPdS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2041:
	.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>
// #include <pycuda-helpers.hpp>
__global__ void copyDensity( double *cnsv, double *rho ){
	int t_j = blockIdx.x*blockDim.x + threadIdx.x;
  int t_i = blockIdx.y*blockDim.y + threadIdx.y;
  int t_k = blockIdx.z*blockDim.z + threadIdx.z;
  int tid = t_j + t_i*blockDim.x*gridDim.x + t_k*blockDim.x*gridDim.x*blockDim.y*gridDim.y;

	rho[tid] = cnsv[tid];
}

__global__ void FFT_divideK2( const double pi4,
							double *kxfft, double *kyfft, double *kzfft,
				      double *data_re, double *data_im){
  int t_j = blockIdx.x*blockDim.x + threadIdx.x;
  int t_i = blockIdx.y*blockDim.y + threadIdx.y;
  int t_k = blockIdx.z*blockDim.z + threadIdx.z;
  int tid = t_j + t_i*blockDim.x*gridDim.x + t_k*blockDim.x*gridDim.x*blockDim.y*gridDim.y;

  double kx = kxfft[t_j];
  double ky = kyfft[t_i];
  double kz = kzfft[t_k];
  double k2 = kx*kx + ky*ky + kz*kz;
	if ( abs( k2 ) < 1e-5 ) k2 = 1 ;
  data_re[tid] = -data_re[tid]*pi4/k2;
	data_im[tid] = -data_im[tid]*pi4/k2;
}

__global__ void iterPoissonStep(  const int parity,
				 const int nX, const int nY, const int nZ,
         double dx, double dy, double dz,
         const double omega, const double pi4, const double convEpsilon,
				 double* rho_all, double* phi_all, int* converged ){
  int t_j = 2*(blockIdx.x*blockDim.x + threadIdx.x);
  int t_i = blockIdx.y*blockDim.y + threadIdx.y;
  int t_k = blockIdx.z*blockDim.z + threadIdx.z;
  //Make a checkboard 3D grid
  if ( t_i%2 == 0 ){
    if ( t_k%2 == parity ) t_j +=1;
  }
  else if ( (t_k+1)%2 == parity ) t_j +=1;
  int tid = t_j + t_i*nX + t_k*nX*blockDim.y*gridDim.y;

	//Set neighbors ids
	int l_indx, r_indx, d_indx, u_indx, b_indx, t_indx;
	l_indx = t_j==0    ?    1 : t_j-1;  //Left
	r_indx = t_j==nX-1 ? nX-2 : t_j+1;  //Right
	d_indx = t_i==0    ?    1 : t_i-1;  //Down
	u_indx = t_i==nY-1 ? nY-2 : t_i+1;  //Up
	b_indx = t_k==0    ?    1 : t_k-1;  //bottom
	t_indx = t_k==nZ-1 ? nZ-2 : t_k+1;  //top

  double rho, phi_c, phi_l, phi_r, phi_d, phi_u, phi_b, phi_t, phi_new;
	rho = rho_all[ tid ];
	phi_c = phi_all[tid];
	phi_l = phi_all[ l_indx + t_i*nX + t_k*nX*blockDim.y*gridDim.y ];
	phi_r = phi_all[ r_indx + t_i*nX + t_k*nX*blockDim.y*gridDim.y ];
	phi_d = phi_all[ t_j + d_indx*nX + t_k*nX*blockDim.y*gridDim.y ];
	phi_u = phi_all[ t_j + u_indx*nX + t_k*nX*blockDim.y*gridDim.y ];
	phi_b = phi_all[ t_j + t_i*nX + b_indx*nX*blockDim.y*gridDim.y ];
	phi_t = phi_all[ t_j + t_i*nX + t_indx*nX*blockDim.y*gridDim.y ];

  phi_new = (1-omega)*phi_c + omega/6*( phi_l + phi_r + phi_d + phi_u + phi_b + phi_t - dx*dx*pi4*rho );
  phi_all[ tid ] = phi_new;


  if ( ( abs( ( phi_new - phi_c ) / phi_c ) > convEpsilon ) ) converged[0] = 0;
	// if ( ( abs( phi_new - phi_c ) > 0.0001 ) ) converged[0] = 0;
}

__global__ void getGravityForce(const int nCells,  const int nW, const int nH, const int nD,
				 double dx, double dy, double dz,
				 double* gForce_x, double* gForce_y, double* gForce_z,
				 double* cnsv, double *phi_all, double *gravWork ){
  int t_j = blockIdx.x*blockDim.x + threadIdx.x;
  int t_i = blockIdx.y*blockDim.y + threadIdx.y;
  int t_k = blockIdx.z*blockDim.z + threadIdx.z;
  int tid = t_j + t_i*blockDim.x*gridDim.x + t_k*blockDim.x*gridDim.x*blockDim.y*gridDim.y;

	//Set neighbors ids
	int l_indx, r_indx, d_indx, u_indx, b_indx, t_indx;
	l_indx = t_j==0    ?    1 : t_j-1;  //Left
	r_indx = t_j==nW-1 ? nW-2 : t_j+1;  //Right
	d_indx = t_i==0    ?    1 : t_i-1;  //Down
	u_indx = t_i==nH-1 ? nH-2 : t_i+1;  //Up
	b_indx = t_k==0    ?    1 : t_k-1;  //bottom
	t_indx = t_k==nD-1 ? nD-2 : t_k+1;  //top

  double phi_l, phi_r, phi_d, phi_u, phi_b, phi_t;
	phi_l = phi_all[ l_indx + t_i*nW + t_k*nW*nH ];
	phi_r = phi_all[ r_indx + t_i*nW + t_k*nW*nH ];
	phi_d = phi_all[ t_j + d_indx*nW + t_k*nW*nH ];
	phi_u = phi_all[ t_j + u_indx*nW + t_k*nW*nH ];
	phi_b = phi_all[ t_j + t_i*nW + b_indx*nW*nH ];
	phi_t = phi_all[ t_j + t_i*nW + t_indx*nW*nH ];

  // //Boundary conditions
  // if  ( t_j == 0 )        phi_l = phi_r;
  // if  ( t_j == nWidth-1 ) phi_r = phi_l;
  // if  ( t_i == 0 )        phi_d = phi_u;
  // if  ( t_i == nWidth-1 ) phi_u = phi_d;
  // if  ( t_k == 0 )        phi_b = phi_t;
  // if  ( t_k == nWidth-1 ) phi_t = phi_b;

  //Get partial derivatives for force
  double gField_x, gField_y, gField_z;
  gField_x = ( phi_l - phi_r ) * 0.5 / dx;
  gField_y = ( phi_d - phi_u ) * 0.5 / dy;
  gField_z = ( phi_b - phi_t ) * 0.5 / dz;
	double p_x, p_y, p_z, rho;
	rho = cnsv[ 0*nCells + tid ];
	const int factor = 1;
  gForce_x[ tid ] = gField_x * rho * factor;
  gForce_y[ tid ] = gField_y * rho * factor;
  gForce_z[ tid ] = gField_z * rho * factor;
  // gForce_x[ tid ] = gField_x;
  // gForce_y[ tid ] = gField_y;
  // gForce_z[ tid ] = gField_z;

  // //Get momentum for virtual gravitational work
	p_x = cnsv[ 1*nCells + tid ];
	p_y = cnsv[ 2*nCells + tid ];
	p_z = cnsv[ 3*nCells + tid ];
  gravWork[ tid ] = (p_x * gField_x + p_y * gField_y + p_z * gField_z )*factor;

}

	.text
	.file	"cuda_gravity.hip"
	.globl	_Z26__device_stub__copyDensityPdS_ # -- Begin function _Z26__device_stub__copyDensityPdS_
	.type	_Z26__device_stub__copyDensityPdS_,@function
_Z26__device_stub__copyDensityPdS_:     # @_Z26__device_stub__copyDensityPdS_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%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	$_Z11copyDensityPdS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z26__device_stub__copyDensityPdS_, .Lfunc_end0-_Z26__device_stub__copyDensityPdS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z27__device_stub__FFT_divideK2dPdS_S_S_S_ # -- Begin function _Z27__device_stub__FFT_divideK2dPdS_S_S_S_
	.type	_Z27__device_stub__FFT_divideK2dPdS_S_S_S_,@function
_Z27__device_stub__FFT_divideK2dPdS_S_S_S_: # @_Z27__device_stub__FFT_divideK2dPdS_S_S_S_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movsd	%xmm0, (%rax)
	leaq	48(%rsp), %r9
	movq	%rdi, (%r9)
	leaq	40(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	32(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	24(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	16(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%r9, 8(%rbx)
	movq	%rdi, 16(%rbx)
	movq	%rsi, 24(%rbx)
	movq	%rdx, 32(%rbx)
	movq	%rcx, 40(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rsp, %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z12FFT_divideK2dPdS_S_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z27__device_stub__FFT_divideK2dPdS_S_S_S_, .Lfunc_end1-_Z27__device_stub__FFT_divideK2dPdS_S_S_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z30__device_stub__iterPoissonStepiiiiddddddPdS_Pi # -- Begin function _Z30__device_stub__iterPoissonStepiiiiddddddPdS_Pi
	.type	_Z30__device_stub__iterPoissonStepiiiiddddddPdS_Pi,@function
_Z30__device_stub__iterPoissonStepiiiiddddddPdS_Pi: # @_Z30__device_stub__iterPoissonStepiiiiddddddPdS_Pi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$240, %rsp
	.cfi_def_cfa_offset 288
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	4(%rsp), %rsi
	movl	%edx, (%rsi)
	movq	%rsp, %rdx
	movl	%ecx, (%rdx)
	leaq	88(%rsp), %rcx
	movsd	%xmm0, (%rcx)
	leaq	80(%rsp), %r10
	movsd	%xmm1, (%r10)
	leaq	72(%rsp), %r11
	movsd	%xmm2, (%r11)
	leaq	64(%rsp), %r14
	movsd	%xmm3, (%r14)
	leaq	56(%rsp), %r15
	movsd	%xmm4, (%r15)
	leaq	48(%rsp), %r12
	movsd	%xmm5, (%r12)
	leaq	40(%rsp), %r13
	movq	%r8, (%r13)
	leaq	32(%rsp), %r8
	movq	%r9, (%r8)
	leaq	128(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r10, 40(%rbx)
	movq	%r11, 48(%rbx)
	movq	%r14, 56(%rbx)
	movq	%r15, 64(%rbx)
	movq	%r12, 72(%rbx)
	movq	%r13, 80(%rbx)
	movq	%r8, 88(%rbx)
	leaq	288(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	112(%rsp), %r14
	leaq	96(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z15iterPoissonStepiiiiddddddPdS_Pi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$256, %rsp                      # imm = 0x100
	.cfi_adjust_cfa_offset -256
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z30__device_stub__iterPoissonStepiiiiddddddPdS_Pi, .Lfunc_end2-_Z30__device_stub__iterPoissonStepiiiiddddddPdS_Pi
	.cfi_endproc
                                        # -- End function
	.globl	_Z30__device_stub__getGravityForceiiiidddPdS_S_S_S_S_ # -- Begin function _Z30__device_stub__getGravityForceiiiidddPdS_S_S_S_S_
	.type	_Z30__device_stub__getGravityForceiiiidddPdS_S_S_S_S_,@function
_Z30__device_stub__getGravityForceiiiidddPdS_S_S_S_S_: # @_Z30__device_stub__getGravityForceiiiidddPdS_S_S_S_S_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$224, %rsp
	.cfi_def_cfa_offset 272
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	20(%rsp), %rax
	movl	%edi, (%rax)
	leaq	16(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	8(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	72(%rsp), %rcx
	movsd	%xmm0, (%rcx)
	leaq	64(%rsp), %r10
	movsd	%xmm1, (%r10)
	leaq	56(%rsp), %r11
	movsd	%xmm2, (%r11)
	leaq	48(%rsp), %r14
	movq	%r8, (%r14)
	leaq	40(%rsp), %r8
	movq	%r9, (%r8)
	leaq	112(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r10, 40(%rbx)
	movq	%r11, 48(%rbx)
	movq	%r14, 56(%rbx)
	movq	%r8, 64(%rbx)
	leaq	272(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	280(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	288(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	296(%rsp), %rax
	movq	%rax, 96(%rbx)
	leaq	96(%rsp), %r14
	leaq	80(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%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	$_Z15getGravityForceiiiidddPdS_S_S_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$240, %rsp
	.cfi_adjust_cfa_offset -240
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z30__device_stub__getGravityForceiiiidddPdS_S_S_S_S_, .Lfunc_end3-_Z30__device_stub__getGravityForceiiiidddPdS_S_S_S_S_
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z11copyDensityPdS_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z12FFT_divideK2dPdS_S_S_S_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z15iterPoissonStepiiiiddddddPdS_Pi, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z15getGravityForceiiiidddPdS_S_S_S_S_, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__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	.LBB5_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
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z11copyDensityPdS_,@object     # @_Z11copyDensityPdS_
	.section	.rodata,"a",@progbits
	.globl	_Z11copyDensityPdS_
	.p2align	3, 0x0
_Z11copyDensityPdS_:
	.quad	_Z26__device_stub__copyDensityPdS_
	.size	_Z11copyDensityPdS_, 8

	.type	_Z12FFT_divideK2dPdS_S_S_S_,@object # @_Z12FFT_divideK2dPdS_S_S_S_
	.globl	_Z12FFT_divideK2dPdS_S_S_S_
	.p2align	3, 0x0
_Z12FFT_divideK2dPdS_S_S_S_:
	.quad	_Z27__device_stub__FFT_divideK2dPdS_S_S_S_
	.size	_Z12FFT_divideK2dPdS_S_S_S_, 8

	.type	_Z15iterPoissonStepiiiiddddddPdS_Pi,@object # @_Z15iterPoissonStepiiiiddddddPdS_Pi
	.globl	_Z15iterPoissonStepiiiiddddddPdS_Pi
	.p2align	3, 0x0
_Z15iterPoissonStepiiiiddddddPdS_Pi:
	.quad	_Z30__device_stub__iterPoissonStepiiiiddddddPdS_Pi
	.size	_Z15iterPoissonStepiiiiddddddPdS_Pi, 8

	.type	_Z15getGravityForceiiiidddPdS_S_S_S_S_,@object # @_Z15getGravityForceiiiidddPdS_S_S_S_S_
	.globl	_Z15getGravityForceiiiidddPdS_S_S_S_S_
	.p2align	3, 0x0
_Z15getGravityForceiiiidddPdS_S_S_S_S_:
	.quad	_Z30__device_stub__getGravityForceiiiidddPdS_S_S_S_S_
	.size	_Z15getGravityForceiiiidddPdS_S_S_S_S_, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z12FFT_divideK2dPdS_S_S_S_"
	.size	.L__unnamed_2, 28

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z15iterPoissonStepiiiiddddddPdS_Pi"
	.size	.L__unnamed_3, 36

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z15getGravityForceiiiidddPdS_S_S_S_S_"
	.size	.L__unnamed_4, 39

	.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 _Z26__device_stub__copyDensityPdS_
	.addrsig_sym _Z27__device_stub__FFT_divideK2dPdS_S_S_S_
	.addrsig_sym _Z30__device_stub__iterPoissonStepiiiiddddddPdS_Pi
	.addrsig_sym _Z30__device_stub__getGravityForceiiiidddPdS_S_S_S_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z11copyDensityPdS_
	.addrsig_sym _Z12FFT_divideK2dPdS_S_S_S_
	.addrsig_sym _Z15iterPoissonStepiiiiddddddPdS_Pi
	.addrsig_sym _Z15getGravityForceiiiidddPdS_S_S_S_S_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z11copyDensityPdS_     ; -- Begin function _Z11copyDensityPdS_
	.globl	_Z11copyDensityPdS_
	.p2align	8
	.type	_Z11copyDensityPdS_,@function
_Z11copyDensityPdS_:                    ; @_Z11copyDensityPdS_
; %bb.0:
	s_clause 0x1
	s_load_b64 s[4:5], s[0:1], 0x1c
	s_load_b64 s[6:7], s[0:1], 0x10
	v_bfe_u32 v1, v0, 20, 10
	s_mov_b32 s2, s13
	s_waitcnt lgkmcnt(0)
	s_and_b32 s3, s5, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s15, s3, v[1:2]
	v_bfe_u32 v1, v0, 10, 10
	s_lshr_b32 s3, s4, 16
	v_and_b32_e32 v0, 0x3ff, v0
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[3:4], null, v2, s7, s[14:15]
	v_mad_u64_u32 v[4:5], null, v3, s3, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, v4, s6, s[2:3]
	s_load_b128 s[0:3], s[0:1], 0x0
	v_mad_u64_u32 v[2:3], null, v1, s4, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[0:1], 3, v[2:3]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s0, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s1, v1, vcc_lo
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_load_b64 v[2:3], v[2:3], off
	s_waitcnt vmcnt(0)
	global_store_b64 v[0:1], v[2:3], off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z11copyDensityPdS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 272
		.amdhsa_user_sgpr_count 13
		.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 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 6
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z11copyDensityPdS_, .Lfunc_end0-_Z11copyDensityPdS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 212
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; 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: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.protected	_Z12FFT_divideK2dPdS_S_S_S_ ; -- Begin function _Z12FFT_divideK2dPdS_S_S_S_
	.globl	_Z12FFT_divideK2dPdS_S_S_S_
	.p2align	8
	.type	_Z12FFT_divideK2dPdS_S_S_S_,@function
_Z12FFT_divideK2dPdS_S_S_S_:            ; @_Z12FFT_divideK2dPdS_S_S_S_
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x3c
	v_bfe_u32 v1, v0, 10, 10
	s_load_b256 s[4:11], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s12, s2, 16
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s14, s12, v[1:2]
	v_and_b32_e32 v1, 0x3ff, v0
	v_bfe_u32 v0, v0, 20, 10
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[3:4], 3, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, s8, v3
	s_and_b32 s8, s2, 0xffff
	v_add_co_ci_u32_e32 v4, vcc_lo, s9, v4, vcc_lo
	v_mad_u64_u32 v[5:6], null, s13, s8, v[1:2]
	s_and_b32 s2, s3, 0xffff
	global_load_b64 v[3:4], v[3:4], off
	v_mad_u64_u32 v[8:9], null, s15, s2, v[0:1]
	v_ashrrev_i32_e32 v6, 31, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v9, 31, v8
	v_lshlrev_b64 v[6:7], 3, v[5:6]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[0:1], 3, v[8:9]
	v_add_co_u32 v6, vcc_lo, s6, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v7, vcc_lo, s7, v7, vcc_lo
	s_load_b64 s[6:7], s[0:1], 0x30
	v_add_co_u32 v0, vcc_lo, s10, v0
	global_load_b64 v[6:7], v[6:7], off
	v_add_co_ci_u32_e32 v1, vcc_lo, s11, v1, vcc_lo
	global_load_b64 v[0:1], v[0:1], off
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s2, s7, s12
	s_mul_i32 s6, s6, s8
	s_waitcnt vmcnt(2)
	v_mad_u64_u32 v[9:10], null, s2, v8, v[2:3]
	s_load_b128 s[0:3], s[0:1], 0x20
	v_mul_f64 v[2:3], v[3:4], v[3:4]
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[10:11], null, s6, v9, v[5:6]
	v_ashrrev_i32_e32 v11, 31, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[8:9], 3, v[10:11]
	v_fma_f64 v[2:3], v[6:7], v[6:7], v[2:3]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v10, vcc_lo, s0, v8
	v_add_co_ci_u32_e32 v11, vcc_lo, s1, v9, vcc_lo
	s_mov_b32 s0, 0x88e368f1
	s_mov_b32 s1, 0x3ee4f8b5
	global_load_b64 v[12:13], v[10:11], off
	s_waitcnt vmcnt(1)
	v_fma_f64 v[0:1], v[0:1], v[0:1], v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_cmp_nlt_f64_e64 vcc_lo, |v[0:1]|, s[0:1]
	v_cndmask_b32_e32 v1, 0x3ff00000, v1, vcc_lo
	v_cndmask_b32_e32 v0, 0, v0, vcc_lo
	s_waitcnt vmcnt(0)
	v_mul_f64 v[2:3], -v[12:13], s[4:5]
	v_div_scale_f64 v[4:5], null, v[0:1], v[0:1], v[2:3]
	v_div_scale_f64 v[14:15], vcc_lo, v[2:3], v[0:1], v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[6:7], v[4:5]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[12:13], -v[4:5], v[6:7], 1.0
	v_fma_f64 v[6:7], v[6:7], v[12:13], v[6:7]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[12:13], -v[4:5], v[6:7], 1.0
	v_fma_f64 v[6:7], v[6:7], v[12:13], v[6:7]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[12:13], v[14:15], v[6:7]
	v_fma_f64 v[4:5], -v[4:5], v[12:13], v[14:15]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[4:5], v[4:5], v[6:7], v[12:13]
	v_div_fixup_f64 v[2:3], v[4:5], v[0:1], v[2:3]
	v_add_co_u32 v4, vcc_lo, s2, v8
	v_add_co_ci_u32_e32 v5, vcc_lo, s3, v9, vcc_lo
	global_store_b64 v[10:11], v[2:3], off
	global_load_b64 v[2:3], v[4:5], off
	s_waitcnt vmcnt(0)
	v_mul_f64 v[2:3], -v[2:3], s[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_scale_f64 v[6:7], null, v[0:1], v[0:1], v[2:3]
	v_rcp_f64_e32 v[8:9], v[6:7]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[10:11], -v[6:7], v[8:9], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[8:9], v[8:9], v[10:11], v[8:9]
	v_fma_f64 v[10:11], -v[6:7], v[8:9], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f64 v[8:9], v[8:9], v[10:11], v[8:9]
	v_div_scale_f64 v[10:11], vcc_lo, v[2:3], v[0:1], v[2:3]
	v_mul_f64 v[12:13], v[10:11], v[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[6:7], -v[6:7], v[12:13], v[10:11]
	v_div_fmas_f64 v[6:7], v[6:7], v[8:9], v[12:13]
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fixup_f64 v[0:1], v[6:7], v[0:1], v[2:3]
	global_store_b64 v[4:5], v[0:1], off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z12FFT_divideK2dPdS_S_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.amdhsa_user_sgpr_count 13
		.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 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 16
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z12FFT_divideK2dPdS_S_S_S_, .Lfunc_end1-_Z12FFT_divideK2dPdS_S_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 660
; NumSgprs: 18
; NumVgprs: 16
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 16
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; 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: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.protected	_Z15iterPoissonStepiiiiddddddPdS_Pi ; -- Begin function _Z15iterPoissonStepiiiiddddddPdS_Pi
	.globl	_Z15iterPoissonStepiiiiddddddPdS_Pi
	.p2align	8
	.type	_Z15iterPoissonStepiiiiddddddPdS_Pi,@function
_Z15iterPoissonStepiiiiddddddPdS_Pi:    ; @_Z15iterPoissonStepiiiiddddddPdS_Pi
; %bb.0:
	s_load_b64 s[2:3], s[0:1], 0x64
	v_bfe_u32 v1, v0, 10, 10
	v_bfe_u32 v2, v0, 20, 10
	v_and_b32_e32 v0, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s8, s2, 16
	s_and_b32 s3, s3, 0xffff
	v_mad_u64_u32 v[3:4], null, s14, s8, v[1:2]
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(VALU_DEP_1)
	v_mad_u64_u32 v[4:5], null, s15, s3, v[2:3]
	v_and_b32_e32 v1, 1, v3
	s_clause 0x2
	s_load_b32 s3, s[0:1], 0x5c
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[14:15], s[0:1], 0x10
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v1, v4, v1
	v_add_nc_u32_e32 v13, -1, v4
	v_lshrrev_b32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v2, v1, v2
	s_waitcnt lgkmcnt(0)
	v_mul_lo_u32 v11, v3, s5
	v_and_b32_e32 v2, -2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sub_nc_u32_e32 v5, v1, v2
	v_mad_u64_u32 v[1:2], null, s13, s2, v[0:1]
	s_mul_i32 s2, s3, s5
	s_add_i32 s3, s5, -1
	v_cmp_eq_u32_e32 vcc_lo, s4, v5
	s_mul_i32 s2, s2, s8
	s_load_b128 s[8:11], s[0:1], 0x48
	v_mul_lo_u32 v17, s2, v4
	s_add_i32 s4, s5, -2
	v_cndmask_b32_e64 v0, 0, 1, vcc_lo
	s_mov_b32 s12, 0
	s_mov_b32 s13, 0x40180000
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshl_or_b32 v12, v1, 1, v0
	v_dual_mov_b32 v1, 0 :: v_dual_add_nc_u32 v2, v17, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_add_nc_u32_e32 v0, -1, v12
	v_cmp_ne_u32_e32 vcc_lo, 0, v12
	v_add_nc_u32_e32 v5, 1, v12
	v_cndmask_b32_e32 v0, 1, v0, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, s3, v12
	s_add_i32 s3, s6, -1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v0, v2, v0
	v_cndmask_b32_e64 v7, v5, s4, vcc_lo
	s_add_i32 s4, s6, -2
	v_lshlrev_b64 v[5:6], 3, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v2, v7
	v_lshlrev_b64 v[7:8], 3, v[0:1]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s8, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s9, v6, vcc_lo
	v_add_nc_u32_e32 v0, -1, v3
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v7, vcc_lo, s8, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s9, v8, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 0, v3
	s_clause 0x1
	global_load_b64 v[5:6], v[5:6], off
	global_load_b64 v[7:8], v[7:8], off
	v_cndmask_b32_e32 v0, 1, v0, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, v0, s5
	v_add3_u32 v0, v17, v0, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[9:10], 3, v[0:1]
	v_add_nc_u32_e32 v0, 1, v3
	v_add_co_u32 v9, vcc_lo, s8, v9
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v10, vcc_lo, s9, v10, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, s3, v3
	s_add_i32 s3, s7, -1
	global_load_b64 v[9:10], v[9:10], off
	v_cndmask_b32_e64 v0, v0, s4, vcc_lo
	s_add_i32 s4, s7, -2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, v0, s5
	v_add3_u32 v0, v17, v0, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[2:3], 3, v[0:1]
	v_add_nc_u32_e32 v0, v12, v11
	v_add_co_u32 v2, vcc_lo, s8, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v3, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 0, v4
	v_add_nc_u32_e32 v17, v0, v17
	v_cndmask_b32_e32 v13, 1, v13, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v18, 31, v17
	v_mad_u64_u32 v[11:12], null, s2, v13, v[0:1]
	v_dual_mov_b32 v12, v1 :: v_dual_add_nc_u32 v13, 1, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[17:18], 3, v[17:18]
	v_lshlrev_b64 v[11:12], 3, v[11:12]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v11, vcc_lo, s8, v11
	v_add_co_ci_u32_e32 v12, vcc_lo, s9, v12, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, s3, v4
	global_load_b64 v[11:12], v[11:12], off
	v_cndmask_b32_e64 v4, v13, s4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[13:14], null, s2, v4, v[0:1]
	v_mov_b32_e32 v14, v1
	s_load_b256 s[0:7], s[0:1], 0x28
	v_lshlrev_b64 v[13:14], 3, v[13:14]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v13, vcc_lo, s8, v13
	v_add_co_ci_u32_e32 v14, vcc_lo, s9, v14, vcc_lo
	s_clause 0x1
	global_load_b64 v[13:14], v[13:14], off
	global_load_b64 v[2:3], v[2:3], off
	s_waitcnt lgkmcnt(0)
	v_div_scale_f64 v[15:16], null, s[12:13], s[12:13], s[0:1]
	v_add_co_u32 v19, vcc_lo, s6, v17
	v_add_co_ci_u32_e32 v20, vcc_lo, s7, v18, vcc_lo
	v_add_co_u32 v17, vcc_lo, s8, v17
	v_add_co_ci_u32_e32 v18, vcc_lo, s9, v18, vcc_lo
	global_load_b64 v[19:20], v[19:20], off
	global_load_b64 v[25:26], v[17:18], off
	v_rcp_f64_e32 v[21:22], v[15:16]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[23:24], -v[15:16], 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[15:16], v[21:22], 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_fma_f64 v[21:22], v[21:22], v[23:24], v[21:22]
	s_waitcnt vmcnt(6)
	v_add_f64 v[4:5], v[5:6], v[7:8]
	v_div_scale_f64 v[6:7], vcc_lo, s[0:1], 0x40180000, s[0:1]
	s_waitcnt vmcnt(5)
	v_add_f64 v[4:5], v[4:5], v[9:10]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[8:9], v[6:7], v[21:22]
	v_fma_f64 v[6:7], -v[15:16], v[8:9], v[6:7]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_div_fmas_f64 v[6:7], v[6:7], v[21:22], v[8:9]
	s_waitcnt vmcnt(2)
	v_add_f64 v[2:3], v[4:5], v[2:3]
	v_mul_f64 v[4:5], s[14:15], s[14:15]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_div_fixup_f64 v[6:7], v[6:7], 0x40180000, s[0:1]
	v_add_f64 v[2:3], v[2:3], v[11:12]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[4:5], v[4:5], s[2:3]
	v_add_f64 v[2:3], v[2:3], v[13:14]
	s_waitcnt vmcnt(1)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fma_f64 v[2:3], -v[4:5], v[19:20], v[2:3]
	v_add_f64 v[4:5], -s[0:1], 1.0
	s_mov_b32 s1, exec_lo
	v_mul_f64 v[2:3], v[6:7], v[2:3]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[2:3], v[4:5], v[25:26], v[2:3]
	v_add_f64 v[4:5], v[2:3], -v[25:26]
	global_store_b64 v[17:18], v[2:3], off
	v_div_scale_f64 v[6:7], null, v[25:26], v[25:26], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[8:9], v[6:7]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[10:11], -v[6:7], v[8:9], 1.0
	v_fma_f64 v[8:9], v[8:9], v[10:11], v[8:9]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[10:11], -v[6:7], v[8:9], 1.0
	v_fma_f64 v[8:9], v[8:9], v[10:11], v[8:9]
	v_div_scale_f64 v[10:11], vcc_lo, v[4:5], v[25:26], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[12:13], v[10:11], v[8:9]
	v_fma_f64 v[6:7], -v[6:7], v[12:13], v[10:11]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[6:7], v[6:7], v[8:9], v[12:13]
	v_div_fixup_f64 v[4:5], v[6:7], v[25:26], v[4:5]
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_f64_e64 |v[4:5]|, s[4:5]
	s_cbranch_execz .LBB2_2
; %bb.1:
	global_store_b32 v1, v1, s[10:11]
.LBB2_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z15iterPoissonStepiiiiddddddPdS_Pi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 344
		.amdhsa_user_sgpr_count 13
		.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 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 27
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end2:
	.size	_Z15iterPoissonStepiiiiddddddPdS_Pi, .Lfunc_end2-_Z15iterPoissonStepiiiiddddddPdS_Pi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1112
; NumSgprs: 18
; NumVgprs: 27
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 3
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 27
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; 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: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.protected	_Z15getGravityForceiiiidddPdS_S_S_S_S_ ; -- Begin function _Z15getGravityForceiiiidddPdS_S_S_S_S_
	.globl	_Z15getGravityForceiiiidddPdS_S_S_S_S_
	.p2align	8
	.type	_Z15getGravityForceiiiidddPdS_S_S_S_S_,@function
_Z15getGravityForceiiiidddPdS_S_S_S_S_: ; @_Z15getGravityForceiiiidddPdS_S_S_S_S_
; %bb.0:
	s_clause 0x2
	s_load_b64 s[2:3], s[0:1], 0x64
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b512 s[16:31], s[0:1], 0x10
	v_bfe_u32 v2, v0, 10, 10
	v_and_b32_e32 v1, 0x3ff, v0
	v_bfe_u32 v0, v0, 20, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s12, s2, 16
	s_and_b32 s3, s3, 0xffff
	s_and_b32 s2, s2, 0xffff
	v_mad_u64_u32 v[3:4], null, s14, s12, v[2:3]
	v_mad_u64_u32 v[4:5], null, s15, s3, v[0:1]
	v_mad_u64_u32 v[5:6], null, s13, s2, v[1:2]
	s_add_i32 s8, s5, -1
	s_add_i32 s9, s5, -2
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v6, -1, v3
	s_add_i32 s3, s6, -1
	s_add_i32 s10, s6, -2
	s_mul_i32 s6, s6, s5
	v_add_nc_u32_e32 v0, -1, v5
	v_cmp_ne_u32_e32 vcc_lo, 0, v5
	v_add_nc_u32_e32 v1, 1, v5
	v_mul_lo_u32 v12, v3, s5
	v_mul_lo_u32 v2, s6, v4
	v_dual_cndmask_b32 v0, 1, v0 :: v_dual_add_nc_u32 v7, 1, v3
	v_cmp_eq_u32_e32 vcc_lo, s8, v5
	v_add_nc_u32_e32 v9, -1, v4
	s_add_i32 s11, s7, -1
	v_add_nc_u32_e32 v10, 1, v4
	s_add_i32 s7, s7, -2
	v_cndmask_b32_e64 v1, v1, s9, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 0, v3
	v_add_nc_u32_e32 v8, v2, v12
	v_dual_cndmask_b32 v11, 1, v6 :: v_dual_add_nc_u32 v2, v2, v5
	v_cmp_eq_u32_e32 vcc_lo, s3, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_add_nc_u32_e32 v0, v8, v0
	v_add_nc_u32_e32 v6, v8, v1
	v_cndmask_b32_e64 v13, v7, s10, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 0, v4
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_4)
	v_ashrrev_i32_e32 v7, 31, v6
	v_cndmask_b32_e32 v14, 1, v9, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, s11, v4
	v_mad_u64_u32 v[8:9], null, v11, s5, v[2:3]
	v_lshlrev_b64 v[0:1], 3, v[0:1]
	v_lshlrev_b64 v[6:7], 3, v[6:7]
	v_cndmask_b32_e64 v16, v10, s7, vcc_lo
	v_mad_u64_u32 v[10:11], null, v13, s5, v[2:3]
	v_add_nc_u32_e32 v2, v12, v5
	v_ashrrev_i32_e32 v9, 31, v8
	v_add_co_u32 v0, vcc_lo, s30, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s31, v1, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_mad_u64_u32 v[12:13], null, s6, v14, v[2:3]
	v_mad_u64_u32 v[14:15], null, s6, v16, v[2:3]
	v_ashrrev_i32_e32 v11, 31, v10
	v_lshlrev_b64 v[8:9], 3, v[8:9]
	v_add_co_u32 v6, vcc_lo, s30, v6
	v_ashrrev_i32_e32 v13, 31, v12
	s_delay_alu instid0(VALU_DEP_4)
	v_lshlrev_b64 v[10:11], 3, v[10:11]
	v_ashrrev_i32_e32 v15, 31, v14
	v_add_co_ci_u32_e32 v7, vcc_lo, s31, v7, vcc_lo
	v_add_co_u32 v8, vcc_lo, s30, v8
	v_lshlrev_b64 v[12:13], 3, v[12:13]
	v_add_co_ci_u32_e32 v9, vcc_lo, s31, v9, vcc_lo
	v_add_co_u32 v10, vcc_lo, s30, v10
	v_lshlrev_b64 v[14:15], 3, v[14:15]
	v_add_co_ci_u32_e32 v11, vcc_lo, s31, v11, vcc_lo
	v_add_co_u32 v12, vcc_lo, s30, v12
	v_add_co_ci_u32_e32 v13, vcc_lo, s31, v13, vcc_lo
	s_clause 0x1
	global_load_b64 v[0:1], v[0:1], off
	global_load_b64 v[6:7], v[6:7], off
	v_add_co_u32 v14, vcc_lo, s30, v14
	v_add_co_ci_u32_e32 v15, vcc_lo, s31, v15, vcc_lo
	s_clause 0x3
	global_load_b64 v[8:9], v[8:9], off
	global_load_b64 v[10:11], v[10:11], off
	global_load_b64 v[12:13], v[12:13], off
	global_load_b64 v[14:15], v[14:15], off
	s_load_b128 s[8:11], s[0:1], 0x50
	s_ashr_i32 s5, s4, 31
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s0, s11, s12
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mad_u64_u32 v[16:17], null, s0, v4, v[3:4]
	s_mul_i32 s0, s10, s2
	s_waitcnt vmcnt(4)
	v_add_f64 v[0:1], v[0:1], -v[6:7]
	s_waitcnt vmcnt(2)
	v_add_f64 v[6:7], v[8:9], -v[10:11]
	s_waitcnt vmcnt(0)
	v_add_f64 v[8:9], v[12:13], -v[14:15]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f64 v[0:1], v[0:1], 0.5
	v_mul_f64 v[6:7], v[6:7], 0.5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f64 v[8:9], v[8:9], 0.5
	v_div_scale_f64 v[10:11], null, s[16:17], s[16:17], v[0:1]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_div_scale_f64 v[12:13], null, s[18:19], s[18:19], v[6:7]
	v_div_scale_f64 v[14:15], null, s[20:21], s[20:21], v[8:9]
	v_mad_u64_u32 v[17:18], null, s0, v16, v[5:6]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v18, 31, v17
	v_lshlrev_b64 v[21:22], 3, v[17:18]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v23, vcc_lo, s28, v21
	v_add_co_ci_u32_e32 v24, vcc_lo, s29, v22, vcc_lo
	v_div_scale_f64 v[33:34], vcc_lo, v[0:1], s[16:17], v[0:1]
	global_load_b64 v[25:26], v[23:24], off
	v_rcp_f64_e32 v[2:3], v[10:11]
	v_rcp_f64_e32 v[4:5], v[12:13]
	v_rcp_f64_e32 v[19:20], v[14:15]
	s_delay_alu instid0(TRANS32_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_3)
	v_fma_f64 v[27:28], -v[10:11], v[2:3], 1.0
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[29:30], -v[12:13], v[4:5], 1.0
	v_fma_f64 v[31:32], -v[14:15], v[19:20], 1.0
	v_fma_f64 v[2:3], v[2:3], v[27:28], v[2:3]
	v_fma_f64 v[4:5], v[4:5], v[29:30], v[4:5]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[18:19], v[19:20], v[31:32], v[19:20]
	v_fma_f64 v[27:28], -v[10:11], v[2:3], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[29:30], -v[12:13], v[4:5], 1.0
	v_fma_f64 v[31:32], -v[14:15], v[18:19], 1.0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_fma_f64 v[2:3], v[2:3], v[27:28], v[2:3]
	v_div_scale_f64 v[27:28], s0, v[6:7], s[18:19], v[6:7]
	v_fma_f64 v[4:5], v[4:5], v[29:30], v[4:5]
	v_div_scale_f64 v[29:30], s1, v[8:9], s[20:21], v[8:9]
	v_fma_f64 v[18:19], v[18:19], v[31:32], v[18:19]
	v_mul_f64 v[31:32], v[33:34], v[2:3]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f64 v[35:36], v[27:28], v[4:5]
	v_mul_f64 v[37:38], v[29:30], v[18:19]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[10:11], -v[10:11], v[31:32], v[33:34]
	v_fma_f64 v[12:13], -v[12:13], v[35:36], v[27:28]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[14:15], -v[14:15], v[37:38], v[29:30]
	v_div_fmas_f64 v[2:3], v[10:11], v[2:3], v[31:32]
	s_mov_b32 vcc_lo, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_div_fmas_f64 v[4:5], v[12:13], v[4:5], v[35:36]
	s_mov_b32 vcc_lo, s1
	v_add3_u32 v12, s4, s4, v17
	v_div_fmas_f64 v[10:11], v[14:15], v[18:19], v[37:38]
	v_add_co_u32 v14, vcc_lo, s22, v21
	v_add_co_ci_u32_e32 v15, vcc_lo, s23, v22, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v13, 31, v12
	v_add_co_u32 v16, vcc_lo, s24, v21
	v_add_co_ci_u32_e32 v17, vcc_lo, s25, v22, vcc_lo
	v_lshlrev_b64 v[18:19], 3, v[12:13]
	s_lshl_b64 s[0:1], s[4:5], 3
	v_div_fixup_f64 v[0:1], v[2:3], s[16:17], v[0:1]
	v_div_fixup_f64 v[2:3], v[4:5], s[18:19], v[6:7]
	v_div_fixup_f64 v[4:5], v[10:11], s[20:21], v[8:9]
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f64 v[6:7], v[0:1], v[25:26]
	v_mul_f64 v[8:9], v[2:3], v[25:26]
	s_delay_alu instid0(VALU_DEP_3)
	v_mul_f64 v[10:11], v[25:26], v[4:5]
	v_add_co_u32 v25, vcc_lo, s26, v21
	v_add_co_ci_u32_e32 v26, vcc_lo, s27, v22, vcc_lo
	v_add_co_u32 v18, vcc_lo, s28, v18
	v_add_co_ci_u32_e32 v19, vcc_lo, s29, v19, vcc_lo
	global_store_b64 v[14:15], v[6:7], off
	global_store_b64 v[16:17], v[8:9], off
	global_store_b64 v[25:26], v[10:11], off
	global_load_b64 v[6:7], v[18:19], off
	v_add_co_u32 v8, vcc_lo, v23, s0
	v_add_co_ci_u32_e32 v9, vcc_lo, s1, v24, vcc_lo
	v_add_nc_u32_e32 v10, s4, v12
	global_load_b64 v[8:9], v[8:9], off
	v_ashrrev_i32_e32 v11, 31, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[10:11], 3, v[10:11]
	v_add_co_u32 v10, vcc_lo, s28, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v11, vcc_lo, s29, v11, vcc_lo
	global_load_b64 v[10:11], v[10:11], off
	s_waitcnt vmcnt(2)
	v_mul_f64 v[2:3], v[2:3], v[6:7]
	s_waitcnt vmcnt(1)
	v_fma_f64 v[0:1], v[0:1], v[8:9], v[2:3]
	v_add_co_u32 v2, vcc_lo, s8, v21
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v22, vcc_lo
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_3)
	v_fma_f64 v[0:1], v[4:5], v[10:11], v[0:1]
	global_store_b64 v[2:3], v[0:1], off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z15getGravityForceiiiidddPdS_S_S_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 344
		.amdhsa_user_sgpr_count 13
		.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 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 39
		.amdhsa_next_free_sgpr 32
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end3:
	.size	_Z15getGravityForceiiiidddPdS_S_S_S_S_, .Lfunc_end3-_Z15getGravityForceiiiidddPdS_S_S_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1232
; NumSgprs: 34
; NumVgprs: 39
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 4
; VGPRBlocks: 4
; NumSGPRsForWavesPerEU: 34
; NumVGPRsForWavesPerEU: 39
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; 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: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.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
      - .offset:         16
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         20
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         28
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         30
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         32
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         34
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         36
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         38
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         80
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 272
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z11copyDensityPdS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z11copyDensityPdS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           8
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z12FFT_divideK2dPdS_S_S_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z12FFT_divideK2dPdS_S_S_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     16
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         12
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           8
        .value_kind:     by_value
      - .offset:         24
        .size:           8
        .value_kind:     by_value
      - .offset:         32
        .size:           8
        .value_kind:     by_value
      - .offset:         40
        .size:           8
        .value_kind:     by_value
      - .offset:         48
        .size:           8
        .value_kind:     by_value
      - .offset:         56
        .size:           8
        .value_kind:     by_value
      - .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
      - .offset:         88
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         92
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         96
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         100
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         102
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         104
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         106
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         108
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         110
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         136
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         144
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         152
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 344
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z15iterPoissonStepiiiiddddddPdS_Pi
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z15iterPoissonStepiiiiddddddPdS_Pi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     27
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         12
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           8
        .value_kind:     by_value
      - .offset:         24
        .size:           8
        .value_kind:     by_value
      - .offset:         32
        .size:           8
        .value_kind:     by_value
      - .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
      - .offset:         88
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         92
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         96
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         100
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         102
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         104
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         106
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         108
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         110
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         128
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         136
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         144
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         152
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 344
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z15getGravityForceiiiidddPdS_S_S_S_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     34
    .sgpr_spill_count: 0
    .symbol:         _Z15getGravityForceiiiidddPdS_S_S_S_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     39
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
