/// CUDA compileable code platform specific to Linux systems. Language standard: Host: C++11; Device: CUDA C. Requires GCC and NVCC linked to CUDA 8.0+ libraries.
/// This library is released for public use as material supplementary to PhD thesis: Development of Specialized Non-Linear Inversion Algorithms, Basis Functions, Eikonal Solvers, and Their Integration for Use in Joint Seismic and Gravitational Tomographic Inversion
/// Author: Zagid Abatchev - abatchev@ucla.edu ; PI: Paul Davis. Department of Earth, Planetary, and Space Sciences; University of California, Los Angeles; Date of publication: March 2019.///
/// This library is published as is and may have unresolved bugs. Use at own discretion. Users are encouraged to adapt some or all of this code without restriction. Refer to Github page for updated source code, located at: https://github.com/abatchev/ZTM .

///GFM_GPU.cu is a standalone gravity perturbation field forward model, implemented in CUDA C. Uses system shared virtual memory for fast file access, requires NVCC and CUDA 8.0+ to compile and properly run.///
///WARNING: This is an experimental implementation, and requires properly setting inversion specific grid size and decimation factors di, dj, dk, fh, fv inside main(). ///

#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <fstream>
#include <iostream>

__global__ void GF(int bfr, int di, int dj, int dk, int fh, int fv, float *g, float *rho, float * s){                               //CUDA device kernel function GF() for forward modeling perturbation gravity fields as a function of perturbed density fields.
	int n=blockDim.x*blockIdx.x+threadIdx.x; float c=66.7f/(fh*fh*fv); int ii=di*fh, jj=dj*fh;                                      //Set raster index from CUDA kernel thread and block ID.
	float i=(n%ii)*1.f/fh, j=((n/ii)%jj)*1.f/fh, k=(n/(ii*jj))*1.f/fv;                                                              //Set gravity render field grid axial indices i,j,k.
	if(fabs(rho[n])>.001f){                                                                                                         //If absolute density perturbation is greater than .001cc, iterate over grid.
		for(int ni=bfr; ni<di-bfr; ni++){for(int nj=bfr; nj<dj-bfr; nj++){                                                          //Iterate through all internal points in grid at least "bfr" axial indices away from a boundary.
			int m=ni+nj*di; float i0=ni, j0=nj, k0=s[m];                                                                            //Compute flattened linear index and convert axial indices to floats.
			atomicAdd(&g[m], c*rho[n]*(k0-k)/powf((i0-i)*(i0-i)+(j0-j)*(j0-j)+(k0-k)*(k0-k),1.5));}}}}                              //Perform atomic addition of perturbation gravity due to the intergrated mass of a grid block approximated at far field as a point mass perturbation.

int main(){int di=61, dj=61, dk=35, fh=1, fv=2, bfr=10, ii=(di-1)*fh+1, jj=(dj-1)*fh+1, kk=(dk-1)*fv+1, N=ii*jj*kk, M=di*dj, lr,ls;  //Declare inversion specific array dimensions di,dj,dk; horizontal and vertical grid point decimation fh, fv; boundary buffer size bfr (where gravity is not rendered due to edge effects).
	float *h_r=(float *)malloc(4*N); FILE *R=fopen("/dev/shm/R.dat","r"); lr=fread(h_r,4,N,R); fclose(R);                            //Allocate host memory for imported density perturbation field h_r. Import perturbation density field from shared memory directory in R.dat
	float *h_s=(float *)malloc(4*M); FILE *S=fopen("/dev/shm/S.dat","r"); ls=fread(h_s,4,M,S); fclose(S); ls+=lr;                    //Allocate host memory for imported gravity data h_s. Import perturbation density field from shared memory directory in S.dat
	float *h_g=(float *)malloc(4*M); for(int n=0;n<M;n++){h_g[n]=0;}                                                                 //Allocate host memory for rendered gravity field h_g. Initialize with zeros.
	float *r; cudaMalloc((void **)&r,4*N); cudaMemcpy(r,h_r,4*N,cudaMemcpyHostToDevice);                                             //Allocate device memory for imported density perturbation field R.
	float *g; cudaMalloc((void **)&g,4*M); cudaMemcpy(g,h_g,4*M,cudaMemcpyHostToDevice);                                             //Allocate device memory for rendered gravity field g.
	float *s; cudaMalloc((void **)&s,4*M); cudaMemcpy(s,h_s,4*M,cudaMemcpyHostToDevice);                                             //Allocate device memory for imported gravity data S.
	GF<<<N/ii,ii>>>(bfr,di,dj,dk,fh,fv,g,r,s);                                                                                       //Execute cuda kernel on the GPU
	cudaMemcpy(h_g,g,4*M,cudaMemcpyDeviceToHost);                                                                                    //Copy rendered field at pointer g from device to host memory (as h_g)
	FILE *G=fopen("/dev/shm/G.dat","w"); fwrite(h_g,4,M,G); fclose(G);                                                               //Open gravity field render file G.dat in system shared memory, write rendered array, close file.
	cudaFree(r); cudaFree(g); cudaFree(s); free(h_r); free(h_g); free(h_s);}                                                         //Free host and device memory.

	.file	"tmpxft_0031d974_00000000-6_GFM_GPU.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3769:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3769:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z31__device_stub__Z2GFiiiiiiPfS_S_iiiiiiPfS_S_
	.type	_Z31__device_stub__Z2GFiiiiiiPfS_S_iiiiiiPfS_S_, @function
_Z31__device_stub__Z2GFiiiiiiPfS_S_iiiiiiPfS_S_:
.LFB3791:
	.cfi_startproc
	endbr64
	subq	$200, %rsp
	.cfi_def_cfa_offset 208
	movq	208(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rax, 16(%rsp)
	movq	216(%rsp), %rax
	movl	%edx, 36(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rax, 8(%rsp)
	movq	224(%rsp), %rax
	movl	%ecx, 32(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 28(%rsp)
	movl	%r9d, 24(%rsp)
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 168(%rsp)
	movq	%rsp, %rax
	movq	%rax, 176(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	56(%rsp)
	.cfi_def_cfa_offset 216
	leaq	_Z2GFiiiiiiPfS_S_(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 224
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L2:
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$200, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3791:
	.size	_Z31__device_stub__Z2GFiiiiiiPfS_S_iiiiiiPfS_S_, .-_Z31__device_stub__Z2GFiiiiiiPfS_S_iiiiiiPfS_S_
	.globl	_Z2GFiiiiiiPfS_S_
	.type	_Z2GFiiiiiiPfS_S_, @function
_Z2GFiiiiiiPfS_S_:
.LFB3792:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z2GFiiiiiiPfS_S_iiiiiiPfS_S_
	.cfi_endproc
.LFE3792:
	.size	_Z2GFiiiiiiPfS_S_, .-_Z2GFiiiiiiPfS_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"r"
.LC1:
	.string	"/dev/shm/R.dat"
.LC2:
	.string	"/dev/shm/S.dat"
.LC3:
	.string	"w"
.LC4:
	.string	"/dev/shm/G.dat"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3766:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	$1026996, %edi
	leaq	.LC0(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	subq	$72, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	call	malloc@PLT
	movq	%r13, %rsi
	leaq	.LC1(%rip), %rdi
	movq	%rax, %r12
	call	fopen@PLT
	movl	$256749, %ecx
	movl	$4, %edx
	movq	%r12, %rdi
	movq	%rax, %r8
	movl	$1026996, %esi
	movq	%rax, %rbx
	call	__fread_chk@PLT
	movq	%rbx, %rdi
	call	fclose@PLT
	movl	$14884, %edi
	call	malloc@PLT
	movq	%r13, %rsi
	leaq	.LC2(%rip), %rdi
	movq	%rax, %rbp
	call	fopen@PLT
	movl	$4, %edx
	movl	$3721, %ecx
	movq	%rbp, %rdi
	movq	%rax, %r8
	movl	$14884, %esi
	movq	%rax, %rbx
	call	__fread_chk@PLT
	movq	%rbx, %rdi
	call	fclose@PLT
	movl	$14884, %edi
	call	malloc@PLT
	movl	$3721, %ecx
	movl	$1026996, %esi
	movq	%rax, %rbx
	xorl	%eax, %eax
	movq	%rbx, %rdi
	rep stosl
	leaq	8(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%r12, %rsi
	movl	$1026996, %edx
	call	cudaMemcpy@PLT
	leaq	16(%rsp), %rdi
	movl	$14884, %esi
	call	cudaMalloc@PLT
	movq	16(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$14884, %edx
	call	cudaMemcpy@PLT
	leaq	24(%rsp), %rdi
	movl	$14884, %esi
	call	cudaMalloc@PLT
	movq	24(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbp, %rsi
	movl	$14884, %edx
	call	cudaMemcpy@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$4294967357, %rdx
	movl	$1, %esi
	movabsq	$4294971505, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	pushq	%rax
	.cfi_def_cfa_offset 120
	movl	$35, %ecx
	movl	$61, %edx
	movl	$61, %esi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 128
	movl	$2, %r9d
	movl	$1, %r8d
	movl	$10, %edi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 136
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	call	_Z31__device_stub__Z2GFiiiiiiPfS_S_iiiiiiPfS_S_
	addq	$32, %rsp
	.cfi_def_cfa_offset 112
.L9:
	movq	16(%rsp), %rsi
	movl	$2, %ecx
	movl	$14884, %edx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	leaq	.LC3(%rip), %rsi
	leaq	.LC4(%rip), %rdi
	call	fopen@PLT
	movl	$3721, %edx
	movl	$4, %esi
	movq	%rbx, %rdi
	movq	%rax, %rcx
	movq	%rax, %r13
	call	fwrite@PLT
	movq	%r13, %rdi
	call	fclose@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$72, %rsp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3766:
	.size	main, .-main
	.section	.rodata.str1.1
.LC5:
	.string	"_Z2GFiiiiiiPfS_S_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3794:
	.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	.LC5(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z2GFiiiiiiPfS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3794:
	.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:

/// CUDA compileable code platform specific to Linux systems. Language standard: Host: C++11; Device: CUDA C. Requires GCC and NVCC linked to CUDA 8.0+ libraries.
/// This library is released for public use as material supplementary to PhD thesis: Development of Specialized Non-Linear Inversion Algorithms, Basis Functions, Eikonal Solvers, and Their Integration for Use in Joint Seismic and Gravitational Tomographic Inversion
/// Author: Zagid Abatchev - abatchev@ucla.edu ; PI: Paul Davis. Department of Earth, Planetary, and Space Sciences; University of California, Los Angeles; Date of publication: March 2019.///
/// This library is published as is and may have unresolved bugs. Use at own discretion. Users are encouraged to adapt some or all of this code without restriction. Refer to Github page for updated source code, located at: https://github.com/abatchev/ZTM .

///GFM_GPU.cu is a standalone gravity perturbation field forward model, implemented in CUDA C. Uses system shared virtual memory for fast file access, requires NVCC and CUDA 8.0+ to compile and properly run.///
///WARNING: This is an experimental implementation, and requires properly setting inversion specific grid size and decimation factors di, dj, dk, fh, fv inside main(). ///

#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <hip/hip_runtime.h>
#include <fstream>
#include <iostream>

__global__ void GF(int bfr, int di, int dj, int dk, int fh, int fv, float *g, float *rho, float * s){                               //CUDA device kernel function GF() for forward modeling perturbation gravity fields as a function of perturbed density fields.
	int n=blockDim.x*blockIdx.x+threadIdx.x; float c=66.7f/(fh*fh*fv); int ii=di*fh, jj=dj*fh;                                      //Set raster index from CUDA kernel thread and block ID.
	float i=(n%ii)*1.f/fh, j=((n/ii)%jj)*1.f/fh, k=(n/(ii*jj))*1.f/fv;                                                              //Set gravity render field grid axial indices i,j,k.
	if(fabs(rho[n])>.001f){                                                                                                         //If absolute density perturbation is greater than .001cc, iterate over grid.
		for(int ni=bfr; ni<di-bfr; ni++){for(int nj=bfr; nj<dj-bfr; nj++){                                                          //Iterate through all internal points in grid at least "bfr" axial indices away from a boundary.
			int m=ni+nj*di; float i0=ni, j0=nj, k0=s[m];                                                                            //Compute flattened linear index and convert axial indices to floats.
			atomicAdd(&g[m], c*rho[n]*(k0-k)/powf((i0-i)*(i0-i)+(j0-j)*(j0-j)+(k0-k)*(k0-k),1.5));}}}}                              //Perform atomic addition of perturbation gravity due to the intergrated mass of a grid block approximated at far field as a point mass perturbation.

int main(){int di=61, dj=61, dk=35, fh=1, fv=2, bfr=10, ii=(di-1)*fh+1, jj=(dj-1)*fh+1, kk=(dk-1)*fv+1, N=ii*jj*kk, M=di*dj, lr,ls;  //Declare inversion specific array dimensions di,dj,dk; horizontal and vertical grid point decimation fh, fv; boundary buffer size bfr (where gravity is not rendered due to edge effects).
	float *h_r=(float *)malloc(4*N); FILE *R=fopen("/dev/shm/R.dat","r"); lr=fread(h_r,4,N,R); fclose(R);                            //Allocate host memory for imported density perturbation field h_r. Import perturbation density field from shared memory directory in R.dat
	float *h_s=(float *)malloc(4*M); FILE *S=fopen("/dev/shm/S.dat","r"); ls=fread(h_s,4,M,S); fclose(S); ls+=lr;                    //Allocate host memory for imported gravity data h_s. Import perturbation density field from shared memory directory in S.dat
	float *h_g=(float *)malloc(4*M); for(int n=0;n<M;n++){h_g[n]=0;}                                                                 //Allocate host memory for rendered gravity field h_g. Initialize with zeros.
	float *r; hipMalloc((void **)&r,4*N); hipMemcpy(r,h_r,4*N,hipMemcpyHostToDevice);                                             //Allocate device memory for imported density perturbation field R.
	float *g; hipMalloc((void **)&g,4*M); hipMemcpy(g,h_g,4*M,hipMemcpyHostToDevice);                                             //Allocate device memory for rendered gravity field g.
	float *s; hipMalloc((void **)&s,4*M); hipMemcpy(s,h_s,4*M,hipMemcpyHostToDevice);                                             //Allocate device memory for imported gravity data S.
	GF<<<N/ii,ii>>>(bfr,di,dj,dk,fh,fv,g,r,s);                                                                                       //Execute cuda kernel on the GPU
	hipMemcpy(h_g,g,4*M,hipMemcpyDeviceToHost);                                                                                    //Copy rendered field at pointer g from device to host memory (as h_g)
	FILE *G=fopen("/dev/shm/G.dat","w"); fwrite(h_g,4,M,G); fclose(G);                                                               //Open gravity field render file G.dat in system shared memory, write rendered array, close file.
	hipFree(r); hipFree(g); hipFree(s); free(h_r); free(h_g); free(h_s);}                                                         //Free host and device memory.

	.text
	.file	"GFM_GPU.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z17__device_stub__GFiiiiiiPfS_S_ # -- Begin function _Z17__device_stub__GFiiiiiiPfS_S_
	.type	_Z17__device_stub__GFiiiiiiPfS_S_,@function
_Z17__device_stub__GFiiiiiiPfS_S_:      # @_Z17__device_stub__GFiiiiiiPfS_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	$160, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	28(%rsp), %rax
	movl	%edi, (%rax)
	leaq	24(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	20(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	16(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	12(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	8(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	216(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	224(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	40(%rsp), %r12
	leaq	32(%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	$_Z2GFiiiiiiPfS_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	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	_Z17__device_stub__GFiiiiiiPfS_S_, .Lfunc_end0-_Z17__device_stub__GFiiiiiiPfS_S_
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$24, %rsp
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$1026996, %edi                  # imm = 0xFABB4
	callq	malloc
	movq	%rax, %rbx
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, %r14
	movl	$4, %esi
	movl	$256749, %edx                   # imm = 0x3EAED
	movq	%rbx, %rdi
	movq	%rax, %rcx
	callq	fread
	movq	%r14, %rdi
	callq	fclose
	movl	$14884, %edi                    # imm = 0x3A24
	callq	malloc
	movq	%rax, %r14
	movl	$.L.str.2, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, %r15
	movl	$4, %esi
	movl	$3721, %edx                     # imm = 0xE89
	movq	%r14, %rdi
	movq	%rax, %rcx
	callq	fread
	movq	%r15, %rdi
	callq	fclose
	movl	$1, %edi
	movl	$14884, %esi                    # imm = 0x3A24
	callq	calloc@PLT
	movq	%rax, %r15
	leaq	16(%rsp), %r12
	movl	$1026996, %esi                  # imm = 0xFABB4
	movq	%r12, %rdi
	callq	hipMalloc
	movq	(%r12), %rdi
	movl	$1026996, %edx                  # imm = 0xFABB4
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	%rsp, %r12
	movl	$14884, %esi                    # imm = 0x3A24
	movq	%r12, %rdi
	callq	hipMalloc
	movq	(%r12), %rdi
	movl	$14884, %edx                    # imm = 0x3A24
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	8(%rsp), %r12
	movl	$14884, %esi                    # imm = 0x3A24
	movq	%r12, %rdi
	callq	hipMalloc
	movq	(%r12), %rdi
	movl	$14884, %edx                    # imm = 0x3A24
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967357, %rdx               # imm = 0x10000003D
	leaq	4148(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	$10, %edi
	movl	$61, %esi
	movl	$61, %edx
	movl	$35, %ecx
	movl	$1, %r8d
	movl	$2, %r9d
	pushq	16(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	32(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	24(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z17__device_stub__GFiiiiiiPfS_S_
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
.LBB1_2:
	movq	(%rsp), %rsi
	movl	$14884, %edx                    # imm = 0x3A24
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	callq	fopen
	movq	%rax, %r12
	movl	$4, %esi
	movl	$3721, %edx                     # imm = 0xE89
	movq	%r15, %rdi
	movq	%rax, %rcx
	callq	fwrite
	movq	%r12, %rdi
	callq	fclose
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	%rbx, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	xorl	%eax, %eax
	addq	$24, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z2GFiiiiiiPfS_S_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__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	.LBB3_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
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z2GFiiiiiiPfS_S_,@object       # @_Z2GFiiiiiiPfS_S_
	.section	.rodata,"a",@progbits
	.globl	_Z2GFiiiiiiPfS_S_
	.p2align	3, 0x0
_Z2GFiiiiiiPfS_S_:
	.quad	_Z17__device_stub__GFiiiiiiPfS_S_
	.size	_Z2GFiiiiiiPfS_S_, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"/dev/shm/R.dat"
	.size	.L.str, 15

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

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"/dev/shm/S.dat"
	.size	.L.str.2, 15

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"/dev/shm/G.dat"
	.size	.L.str.3, 15

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z2GFiiiiiiPfS_S_"
	.size	.L__unnamed_1, 18

	.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 _Z17__device_stub__GFiiiiiiPfS_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z2GFiiiiiiPfS_S_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z2GFiiiiiiPfS_S_       ; -- Begin function _Z2GFiiiiiiPfS_S_
	.globl	_Z2GFiiiiiiPfS_S_
	.p2align	8
	.type	_Z2GFiiiiiiPfS_S_,@function
_Z2GFiiiiiiPfS_S_:                      ; @_Z2GFiiiiiiPfS_S_
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x3c
	s_load_b128 s[8:11], s[0:1], 0x18
	s_mov_b32 s3, exec_lo
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[2:3], null, s15, s2, v[0:1]
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[2:3]
	v_add_co_u32 v0, vcc_lo, s10, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s11, v1, vcc_lo
	global_load_b32 v4, v[0:1], off
	s_waitcnt vmcnt(0)
	v_cmpx_lt_f32_e64 0x3a83126f, |v4|
	s_cbranch_execz .LBB0_13
; %bb.1:                                ; %.preheader59
	s_load_b128 s[12:15], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_sub_i32 s5, s13, s12
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_le_i32 s5, s12
	s_cbranch_scc1 .LBB0_13
; %bb.2:                                ; %.preheader.lr.ph
	s_load_b64 s[2:3], s[0:1], 0x10
	v_add_nc_u32_e32 v6, v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_xor_b32_e32 v6, v6, v3
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s4, s2, s13
	s_ashr_i32 s6, s4, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s7, s4, s6
	s_xor_b32 s7, s7, s6
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cvt_f32_u32_e32 v4, s7
	s_sub_i32 s10, 0, s7
	v_rcp_iflag_f32_e32 v4, v4
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v4, 0x4f7ffffe, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v4, v4
	v_mul_lo_u32 v5, s10, v4
	s_mul_i32 s10, s2, s14
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	s_mul_i32 s11, s4, s10
	s_ashr_i32 s16, s10, 31
	s_ashr_i32 s15, s11, 31
	s_add_i32 s10, s10, s16
	s_add_i32 s11, s11, s15
	v_mul_hi_u32 v5, v4, v5
	s_xor_b32 s11, s11, s15
	s_xor_b32 s10, s10, s16
	v_cvt_f32_u32_e32 v7, s11
	v_cvt_f32_u32_e32 v8, s10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_rcp_iflag_f32_e32 v7, v7
	v_add_nc_u32_e32 v4, v4, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v8, v8
	v_mul_hi_u32 v4, v6, v4
	s_waitcnt_depctr 0xfff
	v_dual_mul_f32 v7, 0x4f7ffffe, v7 :: v_dual_mul_f32 v8, 0x4f7ffffe, v8
	v_mul_lo_u32 v5, v4, s7
	v_add_nc_u32_e32 v9, 1, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_u32_f32_e32 v7, v7
	v_sub_nc_u32_e32 v5, v6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v10, s7, v5
	v_cmp_le_u32_e32 vcc_lo, s7, v5
	v_dual_cndmask_b32 v5, v5, v10 :: v_dual_cndmask_b32 v4, v4, v9
	v_xor_b32_e32 v9, s6, v3
	s_sub_i32 s6, 0, s11
	v_xor_b32_e32 v3, s15, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_cmp_le_u32_e32 vcc_lo, s7, v5
	v_cvt_u32_f32_e32 v5, v8
	v_mul_lo_u32 v8, s6, v7
	s_sub_i32 s6, 0, s10
	v_mul_hi_u32 v8, v7, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v8
	v_mul_hi_u32 v7, v6, v7
	v_add_nc_u32_e32 v10, 1, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cndmask_b32_e32 v4, v4, v10, vcc_lo
	v_mul_lo_u32 v10, s6, v5
	s_mul_i32 s6, s2, s2
	s_mul_i32 s6, s6, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v4, v4, v9
	v_sub_nc_u32_e32 v4, v4, v9
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v9, v5, v10
	v_ashrrev_i32_e32 v10, 31, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v5, v5, v9
	v_add_nc_u32_e32 v8, v4, v10
	v_mul_lo_u32 v9, v7, s11
	v_mul_lo_u32 v4, v4, s4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_xor_b32_e32 v8, v8, v10
	v_sub_nc_u32_e32 v6, v6, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_mul_hi_u32 v5, v8, v5
	v_add_nc_u32_e32 v9, 1, v7
	v_sub_nc_u32_e32 v2, v2, v4
	v_subrev_nc_u32_e32 v11, s11, v6
	v_cmp_le_u32_e32 vcc_lo, s11, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_cvt_f32_i32_e32 v2, v2
	v_mul_lo_u32 v5, v5, s10
	v_dual_cndmask_b32 v7, v7, v9 :: v_dual_cndmask_b32 v6, v6, v11
	v_cvt_f32_i32_e32 v11, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_nc_u32_e32 v9, 1, v7
	v_sub_nc_u32_e32 v5, v8, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v8, s10, v5
	v_cmp_le_u32_e32 vcc_lo, s10, v5
	v_cndmask_b32_e32 v5, v5, v8, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s11, v6
	v_cvt_f32_i32_e32 v8, s6
	s_load_b64 s[6:7], s[0:1], 0x28
	s_mov_b32 s11, 0x3e76c4e1
	v_cndmask_b32_e32 v6, v7, v9, vcc_lo
	v_subrev_nc_u32_e32 v7, s10, v5
	v_cmp_le_u32_e32 vcc_lo, s10, v5
	v_cvt_f32_i32_e32 v9, s2
	s_sub_i32 s10, s14, s12
	s_mov_b32 s14, s12
	s_cmp_gt_i32 s10, s12
	v_cndmask_b32_e32 v5, v5, v7, vcc_lo
	v_div_scale_f32 v21, s3, v2, v9, v2
	v_div_scale_f32 v7, null, v8, v8, 0x42856666
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v5, v5, v10
	v_sub_nc_u32_e32 v4, v5, v10
	v_div_scale_f32 v5, vcc_lo, 0x42856666, v8, 0x42856666
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cvt_f32_i32_e32 v12, v4
	v_div_scale_f32 v4, null, v9, v9, v2
	v_div_scale_f32 v14, null, v9, v9, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v16, v4
	s_waitcnt_depctr 0xfff
	v_fma_f32 v20, -v4, v16, 1.0
	v_fmac_f32_e32 v16, v20, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v23, v21, v16
	v_fma_f32 v25, -v4, v23, v21
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v23, v25, v16
	v_xor_b32_e32 v6, v6, v3
	v_fma_f32 v4, -v4, v23, v21
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v6, v3
	v_rcp_f32_e32 v6, v7
	v_cvt_f32_i32_e32 v3, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_div_scale_f32 v10, null, v11, v11, v3
	s_waitcnt_depctr 0xfff
	v_fma_f32 v13, -v7, v6, 1.0
	v_div_scale_f32 v18, s2, v3, v11, v3
	v_rcp_f32_e32 v15, v10
	v_fmac_f32_e32 v6, v13, v6
	v_rcp_f32_e32 v13, v14
	s_waitcnt_depctr 0xfff
	v_fma_f32 v19, -v10, v15, 1.0
	v_fma_f32 v22, -v14, v13, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v13, v22, v13
	v_fmac_f32_e32 v15, v19, v15
	v_div_scale_f32 v19, s4, v12, v9, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v24, v19, v13
	v_fma_f32 v26, -v14, v24, v19
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_fmac_f32 v24, v26, v13 :: v_dual_mul_f32 v17, v5, v6
	v_mul_f32_e32 v22, v18, v15
	v_fma_f32 v20, -v7, v17, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v17, v20, v6
	v_fma_f32 v20, -v10, v22, v18
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v5, -v7, v17, v5
	v_fmac_f32_e32 v22, v20, v15
	v_fma_f32 v7, -v14, v24, v19
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_fmas_f32 v5, v5, v6, v17
	s_mov_b32 vcc_lo, s3
	v_fma_f32 v6, -v10, v22, v18
	v_div_fmas_f32 v10, v4, v16, v23
	s_mov_b32 vcc_lo, s4
	v_div_fixup_f32 v4, v5, v8, 0x42856666
	v_div_fmas_f32 v7, v7, v13, v24
	s_mov_b32 vcc_lo, s2
	v_div_fixup_f32 v5, v10, v9, v2
	v_div_fmas_f32 v13, v6, v15, v22
	v_mov_b32_e32 v8, 0
	v_div_fixup_f32 v6, v7, v9, v12
	s_cselect_b32 s4, -1, 0
	s_delay_alu instid0(VALU_DEP_3)
	v_div_fixup_f32 v7, v13, v11, v3
.LBB0_3:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_5 Depth 2
                                        ;       Child Loop BB0_6 Depth 3
                                        ;       Child Loop BB0_9 Depth 3
	s_and_not1_b32 vcc_lo, exec_lo, s4
	s_cbranch_vccnz .LBB0_12
; %bb.4:                                ; %.lr.ph
                                        ;   in Loop: Header=BB0_3 Depth=1
	v_cvt_f32_i32_e32 v2, s14
	s_mov_b32 s15, s12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v2, v2, v5
	v_mul_f32_e32 v9, v2, v2
.LBB0_5:                                ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_6 Depth 3
                                        ;       Child Loop BB0_9 Depth 3
	s_mul_i32 s0, s15, s13
	v_cvt_f32_i32_e32 v3, s15
	s_add_i32 s0, s0, s14
	s_mov_b32 s16, exec_lo
	s_ashr_i32 s1, s0, 31
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b64 s[2:3], s[0:1], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s0, s6, s2
	s_addc_u32 s1, s7, s3
	global_load_b32 v2, v8, s[0:1]
	s_waitcnt vmcnt(0)
	v_dual_sub_f32 v3, v3, v6 :: v_dual_sub_f32 v2, v2, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v3, v3, v3, v9
	v_fmac_f32_e32 v3, v2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_frexp_mant_f32_e64 v10, |v3|
	v_cmp_gt_f32_e32 vcc_lo, 0x3f2aaaab, v10
	v_cndmask_b32_e64 v11, 0, 1, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ldexp_f32 v10, v10, v11
	v_add_f32_e32 v11, 1.0, v10
	v_add_f32_e32 v13, -1.0, v10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v15, -1.0, v11
	v_sub_f32_e32 v10, v10, v15
	v_rcp_f32_e32 v12, v11
	global_load_b32 v15, v[0:1], off
	v_mul_f32_e32 v14, v13, v12
	v_cmp_eq_f32_e64 s1, 0, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v16, v11, v14
	v_fma_f32 v11, v14, v11, -v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v11, v14, v10
	v_add_f32_e32 v10, v16, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v17, v13, v10 :: v_dual_sub_f32 v16, v10, v16
	v_sub_f32_e32 v11, v16, v11
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v13, v13, v17
	v_sub_f32_e32 v10, v13, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v10, v11, v10
	v_add_f32_e32 v10, v17, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v10, v12, v10
	v_add_f32_e32 v11, v14, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v12, v11, v14
	v_dual_mul_f32 v13, v11, v11 :: v_dual_sub_f32 v10, v10, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v12, v11, v11, -v13
	v_add_f32_e32 v14, v10, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v12, v11, v14
	v_add_f32_e32 v14, v13, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmaak_f32 v16, s11, v14, 0x3e91f4c4 :: v_dual_sub_f32 v13, v14, v13
	v_sub_f32_e32 v12, v12, v13
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmaak_f32 v16, v14, v16, 0x3ecccdef :: v_dual_mul_f32 v19, v11, v14
	v_mul_f32_e32 v17, v14, v16
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v20, v14, v11, -v19
	v_fma_f32 v13, v14, v16, -v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v13, v12, v16
	v_add_f32_e32 v16, v17, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_dual_sub_f32 v17, v16, v17 :: v_dual_add_f32 v18, 0x3f2aaaaa, v16
	v_fmac_f32_e32 v20, v14, v10
	v_ldexp_f32 v10, v10, 1
	v_sub_f32_e32 v13, v13, v17
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_dual_add_f32 v17, 0xbf2aaaaa, v18 :: v_dual_fmac_f32 v20, v12, v11
	v_ldexp_f32 v11, v11, 1
	v_add_f32_e32 v13, 0x31739010, v13
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_sub_f32_e32 v16, v16, v17
	v_add_f32_e32 v14, v19, v20
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v13, v13, v16
	v_add_f32_e32 v12, v18, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_sub_f32_e32 v16, v18, v12
	v_mul_f32_e32 v17, v14, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v18, v14, v19 :: v_dual_add_f32 v13, v13, v16
	v_fma_f32 v16, v14, v12, -v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v18, v20, v18
	v_fmac_f32_e32 v16, v14, v13
	v_frexp_exp_i32_f32_e32 v13, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v16, v18, v12
	v_subrev_co_ci_u32_e32 v12, vcc_lo, 0, v13, vcc_lo
	v_cmp_eq_f32_e32 vcc_lo, 1.0, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f32_e32 v13, v17, v16
	v_cvt_f32_i32_e32 v12, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f32_e32 v14, v11, v13
	v_dual_sub_f32 v17, v13, v17 :: v_dual_mul_f32 v18, 0x3f317218, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v11, v14, v11 :: v_dual_sub_f32 v16, v16, v17
	v_fma_f32 v17, 0x3f317218, v12, -v18
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v11, v13, v11 :: v_dual_add_f32 v10, v10, v16
	v_dual_fmac_f32 v17, 0xb102e308, v12 :: v_dual_add_f32 v10, v10, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f32_e32 v12, v14, v10
	v_add_f32_e32 v11, v18, v17
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v13, v11, v12
	v_sub_f32_e32 v16, v13, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v19, v13, v16 :: v_dual_sub_f32 v14, v12, v14
	v_sub_f32_e32 v18, v11, v18
	v_dual_sub_f32 v12, v12, v16 :: v_dual_sub_f32 v11, v11, v19
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_f32_e32 v10, v10, v14
	v_sub_f32_e32 v17, v17, v18
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_add_f32 v11, v12, v11 :: v_dual_add_f32 v14, v17, v10
	v_add_f32_e32 v11, v14, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_f32_e32 v16, v13, v11
	v_sub_f32_e32 v12, v14, v17
	v_dual_sub_f32 v13, v16, v13 :: v_dual_sub_f32 v14, v14, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v10, v10, v12 :: v_dual_sub_f32 v11, v11, v13
	v_sub_f32_e32 v12, v17, v14
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f32_e32 v10, v10, v12
	v_cndmask_b32_e64 v12, 0x3fc00000, 1.0, vcc_lo
	v_add_f32_e32 v10, v10, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v11, v16, v10
	v_dual_sub_f32 v13, v11, v16 :: v_dual_mul_f32 v14, v12, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v10, v10, v13
	v_fma_f32 v11, v12, v11, -v14
	v_cmp_class_f32_e64 vcc_lo, v14, 0x204
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v11, v12, v10
	v_add_f32_e32 v10, v14, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v13, v10, v14, vcc_lo
	v_sub_f32_e32 v10, v10, v14
	v_cmp_eq_f32_e32 vcc_lo, 0x42b17218, v13
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v10, v11, v10
	v_trunc_f32_e32 v11, v12
	v_cndmask_b32_e64 v16, 0, 0x37000000, vcc_lo
	v_cmp_neq_f32_e64 vcc_lo, 0x7f800000, |v13|
	v_dual_sub_f32 v17, v13, v16 :: v_dual_cndmask_b32 v10, 0, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_mul_f32_e32 v18, 0x3fb8aa3b, v17
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v17
	v_cmp_nlt_f32_e64 s0, 0x42b17218, v17
	v_add_f32_e32 v10, v16, v10
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f32 v19, 0x3fb8aa3b, v17, -v18
	v_rndne_f32_e32 v20, v18
	v_dual_fmac_f32 v19, 0x32a5705f, v17 :: v_dual_sub_f32 v18, v18, v20
	v_cvt_i32_f32_e32 v14, v20
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v18, v18, v19
	v_exp_f32_e32 v18, v18
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v14, v18, v14
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v18, 0.5, v12 :: v_dual_cndmask_b32 v13, 0, v14
	v_trunc_f32_e32 v14, v18
	v_cmp_eq_f32_e32 vcc_lo, v11, v12
	v_cndmask_b32_e64 v12, 0x7f800000, 0, s1
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v11, 0x7f800000, v13, s0
	v_cmp_neq_f32_e64 s0, v14, v18
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v10, v11, v10, v11
	s_and_b32 vcc_lo, vcc_lo, s0
	v_cmp_class_f32_e64 s0, v3, 0x204
	v_cndmask_b32_e32 v13, 1.0, v3, vcc_lo
	v_cndmask_b32_e32 v14, 0, v3, vcc_lo
	v_cmp_eq_f32_e32 vcc_lo, 0x7f800000, v11
	v_cndmask_b32_e32 v10, v10, v11, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_bfi_b32 v11, 0x7fffffff, v12, v14
	s_or_b32 vcc_lo, s1, s0
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v12, v4, v15
	v_bfi_b32 v10, 0x7fffffff, v10, v13
	v_mul_f32_e32 v2, v2, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v10, v10, v11, vcc_lo
	v_cmp_o_f32_e32 vcc_lo, v3, v3
	v_cndmask_b32_e32 v3, 0x7fc00000, v10, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v10, null, v3, v3, v2
	v_div_scale_f32 v13, vcc_lo, v2, v3, v2
	v_rcp_f32_e32 v11, v10
	s_waitcnt_depctr 0xfff
	v_fma_f32 v12, -v10, v11, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v11, v12, v11
	v_mul_f32_e32 v12, v13, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v14, -v10, v12, v13
	v_fmac_f32_e32 v12, v14, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v10, -v10, v12, v13
	v_div_fmas_f32 v10, v10, v11, v12
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fixup_f32 v2, v10, v3, v2
	v_bfrev_b32_e32 v10, 1
.LBB0_6:                                ; %ComputeLoop
                                        ;   Parent Loop BB0_3 Depth=1
                                        ;     Parent Loop BB0_5 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_ctz_i32_b32 s0, s16
	s_delay_alu instid0(VALU_DEP_2) | instid1(SALU_CYCLE_1)
	v_readlane_b32 s1, v2, s0
	s_lshl_b32 s0, 1, s0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s16, s16, s0
	s_cmp_lg_u32 s16, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_add_f32_e32 v10, s1, v10
	s_cbranch_scc1 .LBB0_6
; %bb.7:                                ; %ComputeEnd
                                        ;   in Loop: Header=BB0_5 Depth=2
	v_mbcnt_lo_u32_b32 v2, exec_lo, 0
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_eq_u32_e32 0, v2
	s_xor_b32 s16, exec_lo, s0
	s_cbranch_execz .LBB0_11
; %bb.8:                                ;   in Loop: Header=BB0_5 Depth=2
	s_add_u32 s0, s8, s2
	s_addc_u32 s1, s9, s3
	s_mov_b32 s2, 0
	global_load_b32 v3, v8, s[0:1]
.LBB0_9:                                ; %atomicrmw.start
                                        ;   Parent Loop BB0_3 Depth=1
                                        ;     Parent Loop BB0_5 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v2, v3, v10
	global_atomic_cmpswap_b32 v2, v8, v[2:3], s[0:1] glc
	s_waitcnt vmcnt(0)
	v_cmp_eq_u32_e32 vcc_lo, v2, v3
	v_mov_b32_e32 v3, v2
	s_or_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_9
; %bb.10:                               ; %Flow
                                        ;   in Loop: Header=BB0_5 Depth=2
	s_or_b32 exec_lo, exec_lo, s2
.LBB0_11:                               ; %Flow86
                                        ;   in Loop: Header=BB0_5 Depth=2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s16
	s_add_i32 s15, s15, 1
	s_cmp_ge_i32 s15, s10
	s_cbranch_scc0 .LBB0_5
.LBB0_12:                               ; %._crit_edge
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_add_i32 s14, s14, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_ge_i32 s14, s5
	s_cbranch_scc0 .LBB0_3
.LBB0_13:                               ; %.loopexit
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z2GFiiiiiiPfS_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 27
		.amdhsa_next_free_sgpr 17
		.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	_Z2GFiiiiiiPfS_S_, .Lfunc_end0-_Z2GFiiiiiiPfS_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 2488
; NumSgprs: 19
; NumVgprs: 27
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 3
; NumSGPRsForWavesPerEU: 19
; NumVGPRsForWavesPerEU: 27
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           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:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .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:           _Z2GFiiiiiiPfS_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     19
    .sgpr_spill_count: 0
    .symbol:         _Z2GFiiiiiiPfS_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     27
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
