//#include "mdCuda.h"

__device__ void cudaReposition(float PP0, float PP1, float PP2, float PA0, float PA1, float PA2, float PL0, float PL1, float PL2, float &Xi, float &Yi, float &Zi, float NA)
{
	float PAPL, H, B;

	if(PP0 > 0){
		PAPL = PA0 + PL0;
		H = (Xi-PAPL) / PL0;
		B = H - 2.0*int(H);
		Xi = B*PL0 + PAPL;
	}

	if(PP1 > 0){
		PAPL = PA1 + PL1;
		H = (Yi-PAPL) / PL1;
		B = H - 2.0*int(H);
		Yi = B*PL1 + PAPL;
	}

	if(PP2 > 0){
		PAPL = PA2 + PL2;
		H = (Zi-PAPL) / PL2;
		B = H - 2.0*int(H);
		Zi = B*PL2 + PAPL;
	}
}

__global__ void mdKernel(int NA, float* FFX, float* FFY, float* FFZ, float* EE,	float* X, float* Y, float* Z, int IPBC, 
	float PP0, float PP1, float PP2, float AL1, float AL2, float A1, float A2, float RL1, float RL2, float D21, float D22, 
	float PA0, float PA1, float PA2, float PB0, float PB1, float PB2, float PL0, float PL1, float PL2)
{
	float XIJ, YIJ, ZIJ, RIJ, RIJ2, EPP, FX2, FY2, FZ2;
	float ARG1, ARG2, EXP1, EXP2, UIJ1, UIJ2, UIJ;
	float FAC1, FAC2, FAC12, XRIJ, YRIJ, ZRIJ;

	int i = blockIdx.x * blockDim.x + threadIdx.x;

	float Xi = X[i];
	float Yi = Y[i];
	float Zi = Z[i];

	cudaReposition(PP0, PP1, PP2, PA0, PA1, PA2, PL0, PL1, PL2, Xi, Yi, Zi, NA);

	///////////////////////////////////////////////////
	// FORCE

	EPP = 0;
	// Forces that effect atoms indexed with i in all three axes
	FX2 = 0;
	FY2 = 0;
	FZ2 = 0;

	for(int j=0; j<NA; j++)
	{
		if(i == j)
			continue;

		// Apply periodic boundaries and find distances between atom I and j. RIJ2 is square of RIJ
		XIJ = X[i] - X[j];
		YIJ = Y[i] - Y[j];
		ZIJ = Z[i] - Z[j];

		double DD, ID;
	
		if(IPBC != 0){
			if(PP0 > 0){
				DD = XIJ / PP0;
				ID = int(DD);
				XIJ = XIJ - PP0*(ID+int(2.0*(DD-ID)));
			}
			if(PP1 > 0){
				DD = YIJ / PP1;
				ID = int(DD);
				YIJ = YIJ - PP1*(ID+int(2.0*(DD-ID)));
			}
			if(PP2 > 0){
				DD = ZIJ / PP2;
				ID = int(DD);
				ZIJ = ZIJ - PP2*(ID+int(2.0*(DD-ID)));
			}
		}
		RIJ2 = XIJ*XIJ + YIJ*YIJ + ZIJ*ZIJ;
		RIJ = sqrt(RIJ2);

		// Calculate potential energy U(r)
		ARG1 = AL1*RIJ2;
		ARG2 = AL2*RIJ2;
		EXP1 = exp(-ARG1);
		EXP2 = exp(-ARG2);
		UIJ1 = A1*EXP1/(pow(RIJ,RL1));
		UIJ2 = A2*EXP2/(pow(RIJ,RL2));
		UIJ = D21*UIJ1 + D22*UIJ2;
		EPP = EPP+UIJ;

		// Calculate forces
		FAC1 = -(RL1/RIJ + 2.0*AL1*RIJ);
		FAC2 = -(RL2/RIJ + 2.0*AL2*RIJ);
		FAC12 = FAC1*D21*UIJ1 + FAC2*D22*UIJ2;
		XRIJ = XIJ/RIJ;
		YRIJ = YIJ/RIJ;  
		ZRIJ = ZIJ/RIJ;
		FX2 += FAC12*XRIJ;
		FY2 += FAC12*YRIJ;
		FZ2 += FAC12*ZRIJ;
	}

	FFX[i] = -FX2;
	FFY[i] = -FY2;
	FFZ[i] = -FZ2;
	EE[i] = EPP;
}
	.file	"tmpxft_002dfea6_00000000-6_kernel.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2011:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2011:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z14cudaRepositionfffffffffRfS_S_f
	.type	_Z14cudaRepositionfffffffffRfS_S_f, @function
_Z14cudaRepositionfffffffffRfS_S_f:
.LFB2008:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2008:
	.size	_Z14cudaRepositionfffffffffRfS_S_f, .-_Z14cudaRepositionfffffffffRfS_S_f
	.globl	_Z61__device_stub__Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffffiPfS_S_S_S_S_S_iffffffffffffffffffff
	.type	_Z61__device_stub__Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffffiPfS_S_S_S_S_S_iffffffffffffffffffff, @function
_Z61__device_stub__Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffffiPfS_S_S_S_S_S_iffffffffffffffffffff:
.LFB2033:
	.cfi_startproc
	endbr64
	subq	$408, %rsp
	.cfi_def_cfa_offset 416
	movq	416(%rsp), %rax
	movl	%edi, 92(%rsp)
	leaq	112(%rsp), %rdi
	movq	%rsi, 80(%rsp)
	leaq	124(%rsp), %rsi
	movq	%rax, 40(%rsp)
	movq	424(%rsp), %rax
	movq	%rdx, 72(%rsp)
	leaq	96(%rsp), %rdx
	movq	%rcx, 64(%rsp)
	leaq	104(%rsp), %rcx
	movq	%r8, 56(%rsp)
	movq	%r9, 48(%rsp)
	movq	%rax, 32(%rsp)
	movss	%xmm0, 88(%rsp)
	movss	%xmm1, 28(%rsp)
	movss	%xmm2, 24(%rsp)
	movss	%xmm3, 20(%rsp)
	movss	%xmm4, 16(%rsp)
	movss	%xmm5, 12(%rsp)
	movss	%xmm6, 8(%rsp)
	movss	%xmm7, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 392(%rsp)
	xorl	%eax, %eax
	leaq	92(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	432(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	440(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	448(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	456(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	464(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	472(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	480(%rsp), %rax
	movq	%rax, 336(%rsp)
	leaq	488(%rsp), %rax
	movq	%rax, 344(%rsp)
	leaq	496(%rsp), %rax
	movq	%rax, 352(%rsp)
	leaq	504(%rsp), %rax
	movq	%rax, 360(%rsp)
	leaq	512(%rsp), %rax
	movq	%rax, 368(%rsp)
	leaq	520(%rsp), %rax
	movq	%rax, 376(%rsp)
	leaq	528(%rsp), %rax
	movq	%rax, 384(%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	.L4
	pushq	104(%rsp)
	.cfi_def_cfa_offset 424
	leaq	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff(%rip), %rdi
	pushq	104(%rsp)
	.cfi_def_cfa_offset 432
	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 424
	popq	%rdx
	.cfi_def_cfa_offset 416
.L4:
	movq	392(%rsp), %rax
	subq	%fs:40, %rax
	je	.L6
	call	__stack_chk_fail@PLT
.L6:
	addq	$408, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2033:
	.size	_Z61__device_stub__Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffffiPfS_S_S_S_S_S_iffffffffffffffffffff, .-_Z61__device_stub__Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffffiPfS_S_S_S_S_S_iffffffffffffffffffff
	.globl	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
	.type	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff, @function
_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff:
.LFB2034:
	.cfi_startproc
	endbr64
	jmp	_Z61__device_stub__Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffffiPfS_S_S_S_S_S_iffffffffffffffffffff
	.cfi_endproc
.LFE2034:
	.size	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff, .-_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2036:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2036:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
//#include "mdCuda.h"

__device__ void cudaReposition(float PP0, float PP1, float PP2, float PA0, float PA1, float PA2, float PL0, float PL1, float PL2, float &Xi, float &Yi, float &Zi, float NA)
{
	float PAPL, H, B;

	if(PP0 > 0){
		PAPL = PA0 + PL0;
		H = (Xi-PAPL) / PL0;
		B = H - 2.0*int(H);
		Xi = B*PL0 + PAPL;
	}

	if(PP1 > 0){
		PAPL = PA1 + PL1;
		H = (Yi-PAPL) / PL1;
		B = H - 2.0*int(H);
		Yi = B*PL1 + PAPL;
	}

	if(PP2 > 0){
		PAPL = PA2 + PL2;
		H = (Zi-PAPL) / PL2;
		B = H - 2.0*int(H);
		Zi = B*PL2 + PAPL;
	}
}

__global__ void mdKernel(int NA, float* FFX, float* FFY, float* FFZ, float* EE,	float* X, float* Y, float* Z, int IPBC, 
	float PP0, float PP1, float PP2, float AL1, float AL2, float A1, float A2, float RL1, float RL2, float D21, float D22, 
	float PA0, float PA1, float PA2, float PB0, float PB1, float PB2, float PL0, float PL1, float PL2)
{
	float XIJ, YIJ, ZIJ, RIJ, RIJ2, EPP, FX2, FY2, FZ2;
	float ARG1, ARG2, EXP1, EXP2, UIJ1, UIJ2, UIJ;
	float FAC1, FAC2, FAC12, XRIJ, YRIJ, ZRIJ;

	int i = blockIdx.x * blockDim.x + threadIdx.x;

	float Xi = X[i];
	float Yi = Y[i];
	float Zi = Z[i];

	cudaReposition(PP0, PP1, PP2, PA0, PA1, PA2, PL0, PL1, PL2, Xi, Yi, Zi, NA);

	///////////////////////////////////////////////////
	// FORCE

	EPP = 0;
	// Forces that effect atoms indexed with i in all three axes
	FX2 = 0;
	FY2 = 0;
	FZ2 = 0;

	for(int j=0; j<NA; j++)
	{
		if(i == j)
			continue;

		// Apply periodic boundaries and find distances between atom I and j. RIJ2 is square of RIJ
		XIJ = X[i] - X[j];
		YIJ = Y[i] - Y[j];
		ZIJ = Z[i] - Z[j];

		double DD, ID;
	
		if(IPBC != 0){
			if(PP0 > 0){
				DD = XIJ / PP0;
				ID = int(DD);
				XIJ = XIJ - PP0*(ID+int(2.0*(DD-ID)));
			}
			if(PP1 > 0){
				DD = YIJ / PP1;
				ID = int(DD);
				YIJ = YIJ - PP1*(ID+int(2.0*(DD-ID)));
			}
			if(PP2 > 0){
				DD = ZIJ / PP2;
				ID = int(DD);
				ZIJ = ZIJ - PP2*(ID+int(2.0*(DD-ID)));
			}
		}
		RIJ2 = XIJ*XIJ + YIJ*YIJ + ZIJ*ZIJ;
		RIJ = sqrt(RIJ2);

		// Calculate potential energy U(r)
		ARG1 = AL1*RIJ2;
		ARG2 = AL2*RIJ2;
		EXP1 = exp(-ARG1);
		EXP2 = exp(-ARG2);
		UIJ1 = A1*EXP1/(pow(RIJ,RL1));
		UIJ2 = A2*EXP2/(pow(RIJ,RL2));
		UIJ = D21*UIJ1 + D22*UIJ2;
		EPP = EPP+UIJ;

		// Calculate forces
		FAC1 = -(RL1/RIJ + 2.0*AL1*RIJ);
		FAC2 = -(RL2/RIJ + 2.0*AL2*RIJ);
		FAC12 = FAC1*D21*UIJ1 + FAC2*D22*UIJ2;
		XRIJ = XIJ/RIJ;
		YRIJ = YIJ/RIJ;  
		ZRIJ = ZIJ/RIJ;
		FX2 += FAC12*XRIJ;
		FY2 += FAC12*YRIJ;
		FZ2 += FAC12*ZRIJ;
	}

	FFX[i] = -FX2;
	FFY[i] = -FY2;
	FFZ[i] = -FZ2;
	EE[i] = EPP;
}
	.text
	.file	"kernel.hip"
	.globl	_Z23__device_stub__mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff # -- Begin function _Z23__device_stub__mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
	.type	_Z23__device_stub__mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff,@function
_Z23__device_stub__mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff: # @_Z23__device_stub__mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$360, %rsp                      # imm = 0x168
	.cfi_def_cfa_offset 416
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	leaq	36(%rsp), %rax
	movl	%edi, (%rax)
	leaq	88(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	80(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	72(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	64(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	56(%rsp), %r8
	movq	%r9, (%r8)
	leaq	32(%rsp), %r9
	movss	%xmm0, (%r9)
	leaq	28(%rsp), %r10
	movss	%xmm1, (%r10)
	leaq	24(%rsp), %r11
	movss	%xmm2, (%r11)
	leaq	20(%rsp), %r14
	movss	%xmm3, (%r14)
	leaq	16(%rsp), %r15
	movss	%xmm4, (%r15)
	leaq	12(%rsp), %r12
	movss	%xmm5, (%r12)
	leaq	8(%rsp), %r13
	movss	%xmm6, (%r13)
	leaq	4(%rsp), %rbp
	movss	%xmm7, (%rbp)
	leaq	128(%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	416(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	424(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	432(%rsp), %rax
	movq	%rax, 64(%rbx)
	movq	%r9, 72(%rbx)
	movq	%r10, 80(%rbx)
	movq	%r11, 88(%rbx)
	movq	%r14, 96(%rbx)
	movq	%r15, 104(%rbx)
	movq	%r12, 112(%rbx)
	movq	%r13, 120(%rbx)
	movq	%rbp, 128(%rbx)
	leaq	440(%rsp), %rax
	movq	%rax, 136(%rbx)
	leaq	448(%rsp), %rax
	movq	%rax, 144(%rbx)
	leaq	456(%rsp), %rax
	movq	%rax, 152(%rbx)
	leaq	464(%rsp), %rax
	movq	%rax, 160(%rbx)
	leaq	472(%rsp), %rax
	movq	%rax, 168(%rbx)
	leaq	480(%rsp), %rax
	movq	%rax, 176(%rbx)
	leaq	488(%rsp), %rax
	movq	%rax, 184(%rbx)
	leaq	496(%rsp), %rax
	movq	%rax, 192(%rbx)
	leaq	504(%rsp), %rax
	movq	%rax, 200(%rbx)
	leaq	512(%rsp), %rax
	movq	%rax, 208(%rbx)
	leaq	520(%rsp), %rax
	movq	%rax, 216(%rbx)
	leaq	528(%rsp), %rax
	movq	%rax, 224(%rbx)
	leaq	112(%rsp), %r14
	leaq	96(%rsp), %r15
	leaq	48(%rsp), %r12
	leaq	40(%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	$_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$376, %rsp                      # imm = 0x178
	.cfi_adjust_cfa_offset -376
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z23__device_stub__mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff, .Lfunc_end0-_Z23__device_stub__mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB1_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB1_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end1:
	.size	__hip_module_ctor, .Lfunc_end1-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB2_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB2_2:
	retq
.Lfunc_end2:
	.size	__hip_module_dtor, .Lfunc_end2-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff,@object # @_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
	.section	.rodata,"a",@progbits
	.globl	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
	.p2align	3, 0x0
_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff:
	.quad	_Z23__device_stub__mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
	.size	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff, 8

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

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

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

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

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff ; -- Begin function _Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
	.globl	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
	.p2align	8
	.type	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff,@function
_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff: ; @_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
; %bb.0:                                ; %_Z14cudaRepositionfffffffffRfS_S_f.exit
	s_clause 0x4
	s_load_b32 s2, s[0:1], 0xa4
	s_load_b32 s33, s[0:1], 0x0
	s_load_b256 s[16:23], s[0:1], 0x8
	s_load_b128 s[24:27], s[0:1], 0x28
	s_load_b64 s[6:7], s[0:1], 0x38
	v_dual_mov_b32 v13, 0 :: v_dual_mov_b32 v18, 0
	v_mov_b32_e32 v19, 0
	v_mov_b32_e32 v17, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_cmp_lt_i32 s33, 1
	v_mad_u64_u32 v[2:3], null, s15, s2, 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], 2, v[2:3]
	s_cbranch_scc1 .LBB0_12
; %bb.1:                                ; %.lr.ph
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, s24, v0
	v_add_co_ci_u32_e32 v4, vcc_lo, s25, v1, vcc_lo
	v_add_co_u32 v5, vcc_lo, s26, v0
	v_add_co_ci_u32_e32 v6, vcc_lo, s27, v1, vcc_lo
	v_add_co_u32 v7, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v1, vcc_lo
	global_load_b32 v14, v[3:4], off
	global_load_b32 v15, v[5:6], off
	global_load_b32 v16, v[7:8], off
	s_clause 0x1
	s_load_b256 s[8:15], s[0:1], 0x40
	s_load_b128 s[28:31], s[0:1], 0x60
	v_dual_mov_b32 v13, 0 :: v_dual_mov_b32 v18, 0
	v_mov_b32_e32 v19, 0
	v_mov_b32_e32 v17, 0
	s_mov_b32 s37, 0x3e76c4e1
	s_waitcnt lgkmcnt(0)
	v_cvt_f64_f32_e32 v[9:10], s12
	v_cvt_f64_f32_e32 v[11:12], s13
	v_cvt_f64_f32_e32 v[3:4], s9
	v_cvt_f64_f32_e32 v[5:6], s10
	v_cvt_f64_f32_e32 v[7:8], s11
	v_cmp_gt_f32_e64 s34, s9, 0
	v_cmp_gt_f32_e64 s35, s10, 0
	v_cmp_gt_f32_e64 s36, s11, 0
	s_cmp_lg_u32 s8, 0
	s_cselect_b32 s8, -1, 0
	v_add_f64 v[9:10], v[9:10], v[9:10]
	v_add_f64 v[11:12], v[11:12], v[11:12]
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	s_mov_b32 s38, exec_lo
	v_cmpx_ne_u32_e32 0, v2
	s_cbranch_execz .LBB0_11
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	s_load_b32 s0, s[24:25], 0x0
	s_load_b32 s1, s[26:27], 0x0
	s_load_b32 s2, s[6:7], 0x0
	s_and_not1_b32 vcc_lo, exec_lo, s8
	s_waitcnt vmcnt(2) lgkmcnt(0)
	v_subrev_f32_e32 v22, s0, v14
	s_waitcnt vmcnt(0)
	v_dual_subrev_f32 v20, s1, v15 :: v_dual_subrev_f32 v21, s2, v16
	s_cbranch_vccnz .LBB0_10
; %bb.4:                                ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_b32 vcc_lo, exec_lo, s34
	s_cbranch_vccnz .LBB0_6
; %bb.5:                                ;   in Loop: Header=BB0_2 Depth=1
	v_div_scale_f32 v23, null, s9, s9, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v24, v23
	s_waitcnt_depctr 0xfff
	v_fma_f32 v25, -v23, v24, 1.0
	v_fmac_f32_e32 v24, v25, v24
	v_div_scale_f32 v25, vcc_lo, v22, s9, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v26, v25, v24
	v_fma_f32 v27, -v23, v26, v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v26, v27, v24
	v_cvt_f64_f32_e32 v[27:28], v22
	v_fma_f32 v23, -v23, v26, v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v23, v23, v24, v26
	v_div_fixup_f32 v23, v23, s9, v22
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cvt_i32_f32_e32 v25, v23
	v_cvt_f64_f32_e32 v[23:24], v23
	v_cvt_f64_i32_e32 v[25:26], v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[23:24], v[23:24], -v[25:26]
	v_add_f64 v[23:24], v[23:24], v[23:24]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_i32_f64_e32 v23, v[23:24]
	v_cvt_f64_i32_e32 v[23:24], v23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[22:23], v[25:26], v[23:24]
	v_fma_f64 v[22:23], -v[22:23], v[3:4], v[27:28]
	s_delay_alu instid0(VALU_DEP_1)
	v_cvt_f32_f64_e32 v22, v[22:23]
.LBB0_6:                                ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_b32 vcc_lo, exec_lo, s35
	s_cbranch_vccnz .LBB0_8
; %bb.7:                                ;   in Loop: Header=BB0_2 Depth=1
	v_div_scale_f32 v23, null, s10, s10, v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v24, v23
	s_waitcnt_depctr 0xfff
	v_fma_f32 v25, -v23, v24, 1.0
	v_fmac_f32_e32 v24, v25, v24
	v_div_scale_f32 v25, vcc_lo, v20, s10, v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v26, v25, v24
	v_fma_f32 v27, -v23, v26, v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v26, v27, v24
	v_cvt_f64_f32_e32 v[27:28], v20
	v_fma_f32 v23, -v23, v26, v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v23, v23, v24, v26
	v_div_fixup_f32 v23, v23, s10, v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cvt_i32_f32_e32 v25, v23
	v_cvt_f64_f32_e32 v[23:24], v23
	v_cvt_f64_i32_e32 v[25:26], v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[23:24], v[23:24], -v[25:26]
	v_add_f64 v[23:24], v[23:24], v[23:24]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_i32_f64_e32 v23, v[23:24]
	v_cvt_f64_i32_e32 v[23:24], v23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[23:24], v[25:26], v[23:24]
	v_fma_f64 v[23:24], -v[23:24], v[5:6], v[27:28]
	s_delay_alu instid0(VALU_DEP_1)
	v_cvt_f32_f64_e32 v20, v[23:24]
.LBB0_8:                                ;   in Loop: Header=BB0_2 Depth=1
	s_and_not1_b32 vcc_lo, exec_lo, s36
	s_cbranch_vccnz .LBB0_10
; %bb.9:                                ;   in Loop: Header=BB0_2 Depth=1
	v_div_scale_f32 v23, null, s11, s11, v21
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v24, v23
	s_waitcnt_depctr 0xfff
	v_fma_f32 v25, -v23, v24, 1.0
	v_fmac_f32_e32 v24, v25, v24
	v_div_scale_f32 v25, vcc_lo, v21, s11, v21
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v26, v25, v24
	v_fma_f32 v27, -v23, v26, v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v26, v27, v24
	v_cvt_f64_f32_e32 v[27:28], v21
	v_fma_f32 v23, -v23, v26, v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v23, v23, v24, v26
	v_div_fixup_f32 v23, v23, s11, v21
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cvt_i32_f32_e32 v25, v23
	v_cvt_f64_f32_e32 v[23:24], v23
	v_cvt_f64_i32_e32 v[25:26], v25
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[23:24], v[23:24], -v[25:26]
	v_add_f64 v[23:24], v[23:24], v[23:24]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_i32_f64_e32 v23, v[23:24]
	v_cvt_f64_i32_e32 v[23:24], v23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[23:24], v[25:26], v[23:24]
	v_fma_f64 v[23:24], -v[23:24], v[7:8], v[27:28]
	s_delay_alu instid0(VALU_DEP_1)
	v_cvt_f32_f64_e32 v21, v[23:24]
.LBB0_10:                               ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v28, v22, v22
	v_fmac_f32_e32 v28, v20, v20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v28, v21, v21
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v28
	v_mul_f32_e32 v23, 0x4f800000, v28
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v23, v28, v23, vcc_lo
	v_sqrt_f32_e32 v24, v23
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v26, 1, v24
	v_add_nc_u32_e32 v25, -1, v24
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v29, -v26, v24, v23
	v_fma_f32 v27, -v25, v24, v23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s0, 0, v27
	v_cndmask_b32_e64 v24, v24, v25, s0
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s0, 0, v29
	v_cndmask_b32_e64 v24, v24, v26, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v25, 0x37800000, v24
	v_cndmask_b32_e32 v24, v24, v25, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v23, 0x260
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v23, v24, v23, vcc_lo
	v_cmp_eq_f32_e32 vcc_lo, 1.0, v23
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v24, s28, 1.0, vcc_lo
	v_cndmask_b32_e64 v25, s29, 1.0, vcc_lo
	v_cmp_neq_f32_e32 vcc_lo, 0, v24
	v_cmp_neq_f32_e64 s5, v24, |v24|
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_cmp_neq_f32_e64 s41, v25, |v25|
	v_cndmask_b32_e32 v26, 1.0, v23, vcc_lo
	v_cmp_neq_f32_e32 vcc_lo, 0, v25
	v_frexp_mant_f32_e64 v29, |v26|
	v_cndmask_b32_e32 v27, 1.0, v23, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_gt_f32_e32 vcc_lo, 0x3f2aaaab, v29
	v_frexp_mant_f32_e64 v30, |v27|
	v_cndmask_b32_e64 v31, 0, 1, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_gt_f32_e64 s0, 0x3f2aaaab, v30
	v_ldexp_f32 v29, v29, v31
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v32, 0, 1, s0
	v_add_f32_e32 v31, 1.0, v29
	s_delay_alu instid0(VALU_DEP_2)
	v_ldexp_f32 v30, v30, v32
	v_add_f32_e32 v35, -1.0, v29
	v_cmp_lt_f32_e64 s39, |v26|, 1.0
	v_cmp_eq_f32_e64 s1, 0, v26
	v_rcp_f32_e32 v33, v31
	v_add_f32_e32 v36, -1.0, v30
	v_add_f32_e32 v32, 1.0, v30
	v_cmp_class_f32_e64 s40, v26, 0x204
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v34, v32
	s_waitcnt_depctr 0xfff
	v_dual_mul_f32 v37, v35, v33 :: v_dual_mul_f32 v38, v36, v34
	v_dual_add_f32 v39, -1.0, v31 :: v_dual_mul_f32 v40, v31, v37
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_add_f32 v41, -1.0, v32 :: v_dual_mul_f32 v42, v32, v38
	v_sub_f32_e32 v29, v29, v39
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v31, v37, v31, -v40
	v_sub_f32_e32 v30, v30, v41
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v32, v38, v32, -v42
	v_dual_fmac_f32 v31, v37, v29 :: v_dual_fmac_f32 v32, v38, v30
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_add_f32 v29, v40, v31 :: v_dual_add_f32 v30, v42, v32
	v_dual_sub_f32 v39, v35, v29 :: v_dual_sub_f32 v40, v29, v40
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v42, v30, v42
	v_sub_f32_e32 v41, v36, v30
	v_dual_sub_f32 v35, v35, v39 :: v_dual_sub_f32 v32, v42, v32
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v31, v40, v31
	v_sub_f32_e32 v29, v35, v29
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v29, v31, v29
	v_add_f32_e32 v29, v39, v29
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v29, v33, v29
	v_add_f32_e32 v31, v37, v29
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v36, v36, v41 :: v_dual_mul_f32 v35, v31, v31
	v_sub_f32_e32 v30, v36, v30
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v30, v32, v30
	v_add_f32_e32 v30, v41, v30
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v30, v34, v30
	v_add_f32_e32 v32, v38, v30
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v34, v32, v38 :: v_dual_sub_f32 v33, v31, v37
	v_mul_f32_e32 v36, v32, v32
	v_dual_sub_f32 v30, v30, v34 :: v_dual_sub_f32 v29, v29, v33
	v_mul_f32_e64 v33, -v28, s12
	v_fma_f32 v34, v31, v31, -v35
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f32 v38, v32, v32, -v36
	v_add_f32_e32 v39, v30, v30
	v_mul_f32_e64 v28, -v28, s13
	v_mul_f32_e32 v40, 0x3fb8aa3b, v33
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_add_f32 v37, v29, v29 :: v_dual_fmac_f32 v38, v32, v39
	v_fma_f32 v39, 0x3fb8aa3b, v33, -v40
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_fmac_f32 v34, v31, v37 :: v_dual_mul_f32 v37, 0x3fb8aa3b, v28
	v_add_f32_e32 v44, v36, v38
	v_rndne_f32_e32 v42, v40
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v39, 0x32a5705f, v33
	v_add_f32_e32 v41, v35, v34
	v_rndne_f32_e32 v45, v37
	v_sub_f32_e32 v36, v44, v36
	v_fmaak_f32 v47, s37, v44, 0x3e91f4c4
	v_fma_f32 v43, 0x3fb8aa3b, v28, -v37
	v_dual_fmaak_f32 v46, s37, v41, 0x3e91f4c4 :: v_dual_sub_f32 v35, v41, v35
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_dual_sub_f32 v36, v38, v36 :: v_dual_fmaak_f32 v47, v44, v47, 0x3ecccdef
	v_sub_f32_e32 v37, v37, v45
	v_fmaak_f32 v46, v41, v46, 0x3ecccdef
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_sub_f32_e32 v34, v34, v35
	v_dual_sub_f32 v40, v40, v42 :: v_dual_fmac_f32 v43, 0x32a5705f, v28
	v_dual_mul_f32 v49, v44, v47 :: v_dual_mul_f32 v48, v41, v46
	v_mul_f32_e32 v52, v32, v44
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_add_f32 v39, v40, v39 :: v_dual_mul_f32 v40, v31, v41
	v_fma_f32 v38, v44, v47, -v49
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_fma_f32 v35, v41, v46, -v48
	v_add_f32_e32 v37, v37, v43
	v_frexp_exp_i32_f32_e32 v43, v26
	v_fma_f32 v50, v41, v31, -v40
	v_dual_fmac_f32 v38, v36, v47 :: v_dual_fmac_f32 v35, v34, v46
	v_fma_f32 v54, v44, v32, -v52
	v_frexp_exp_i32_f32_e32 v47, v27
	v_cvt_i32_f32_e32 v45, v45
	s_delay_alu instid0(VALU_DEP_4)
	v_dual_add_f32 v51, v49, v38 :: v_dual_add_f32 v46, v48, v35
	v_fmac_f32_e32 v50, v41, v29
	v_ldexp_f32 v29, v29, 1
	v_exp_f32_e32 v39, v39
	v_cvt_i32_f32_e32 v42, v42
	v_dual_sub_f32 v48, v46, v48 :: v_dual_add_f32 v53, 0x3f2aaaaa, v46
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_dual_sub_f32 v35, v35, v48 :: v_dual_add_f32 v48, 0xbf2aaaaa, v53
	v_sub_f32_e32 v41, v51, v49
	v_add_f32_e32 v49, 0x3f2aaaaa, v51
	v_fmac_f32_e32 v50, v34, v31
	v_dual_add_f32 v35, 0x31739010, v35 :: v_dual_sub_f32 v46, v46, v48
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_sub_f32_e32 v38, v38, v41
	v_add_f32_e32 v41, 0xbf2aaaaa, v49
	v_ldexp_f32 v31, v31, 1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f32_e32 v34, 0x31739010, v38
	v_sub_f32_e32 v38, v51, v41
	v_subrev_co_ci_u32_e32 v41, vcc_lo, 0, v43, vcc_lo
	v_subrev_co_ci_u32_e64 v47, vcc_lo, 0, v47, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_f32_e32 v34, v34, v38
	v_cvt_f32_i32_e32 v41, v41
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v33
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_add_f32_e32 v43, v49, v34
	v_fmac_f32_e32 v54, v44, v30
	v_add_f32_e32 v35, v35, v46
	v_ldexp_f32 v30, v30, 1
	v_dual_sub_f32 v49, v49, v43 :: v_dual_fmac_f32 v54, v36, v32
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f32_e32 v38, v53, v35
	v_ldexp_f32 v32, v32, 1
	v_add_f32_e32 v44, v52, v54
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v46, v53, v38
	v_sub_f32_e32 v52, v44, v52
	v_add_f32_e32 v36, v40, v50
	v_mul_f32_e32 v51, v44, v43
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_add_f32 v35, v35, v46 :: v_dual_add_f32 v34, v34, v49
	v_mul_f32_e32 v48, v36, v38
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v49, v44, v43, -v51
	v_fma_f32 v46, v36, v38, -v48
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_fmac_f32 v49, v44, v34 :: v_dual_mul_f32 v34, 0x3f317218, v41
	v_dual_fmac_f32 v46, v36, v35 :: v_dual_sub_f32 v35, v54, v52
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_dual_fmac_f32 v49, v35, v43 :: v_dual_sub_f32 v40, v36, v40
	v_fma_f32 v35, 0x3f317218, v41, -v34
	v_exp_f32_e32 v36, v37
	v_cvt_f32_i32_e32 v37, v47
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v40, v50, v40 :: v_dual_fmac_f32 v35, 0xb102e308, v41
	v_dual_mul_f32 v43, 0x3f317218, v37 :: v_dual_fmac_f32 v46, v40, v38
	v_add_f32_e32 v40, v51, v49
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f32_e32 v38, v48, v46
	v_sub_f32_e32 v50, v40, v51
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v44, v31, v38
	v_sub_f32_e32 v31, v44, v31
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v31, v38, v31
	v_sub_f32_e32 v47, v38, v48
	v_ldexp_f32 v38, v39, v42
	v_dual_sub_f32 v46, v46, v47 :: v_dual_sub_f32 v47, v49, v50
	v_fma_f32 v49, 0x3f317218, v37, -v43
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_add_f32 v29, v29, v46 :: v_dual_add_f32 v30, v30, v47
	v_fmac_f32_e32 v49, 0xb102e308, v37
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_f32_e32 v29, v29, v31
	v_ldexp_f32 v31, v36, v45
	v_add_f32_e32 v36, v43, v49
	v_cndmask_b32_e32 v38, 0, v38, vcc_lo
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v28
	s_delay_alu instid0(VALU_DEP_3)
	v_sub_f32_e32 v41, v36, v43
	v_add_f32_e32 v48, v32, v40
	v_add_f32_e32 v37, v44, v29
	v_cndmask_b32_e32 v31, 0, v31, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v33
	v_sub_f32_e32 v41, v49, v41
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v43, v37, v44
	v_dual_sub_f32 v32, v48, v32 :: v_dual_cndmask_b32 v33, 0x7f800000, v38
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v28
	v_dual_sub_f32 v29, v29, v43 :: v_dual_sub_f32 v32, v40, v32
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_mul_f32_e32 v33, s14, v33
	v_cndmask_b32_e32 v28, 0x7f800000, v31, vcc_lo
	v_cmp_gt_f32_e32 vcc_lo, 0, v24
	v_add_f32_e32 v30, v30, v32
	v_add_f32_e32 v32, v34, v35
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v34, v32, v34
	v_sub_f32_e32 v34, v35, v34
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_f32_e32 v46, v34, v29
	v_add_f32_e32 v40, v32, v37
	v_sub_f32_e32 v35, v40, v32
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v45, v40, v35
	v_add_f32_e32 v39, v48, v30
	v_sub_f32_e32 v35, v37, v35
	v_sub_f32_e32 v43, v39, v48
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v30, v30, v43
	v_dual_add_f32 v42, v36, v39 :: v_dual_add_f32 v43, v41, v30
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v44, v42, v36
	v_sub_f32_e32 v37, v42, v44
	v_dual_sub_f32 v39, v39, v44 :: v_dual_sub_f32 v32, v32, v45
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v36, v36, v37
	v_add_f32_e32 v32, v35, v32
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v35, v39, v36
	v_add_f32_e32 v35, v43, v35
	v_sub_f32_e32 v37, v46, v34
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v38, v46, v37
	v_dual_sub_f32 v29, v29, v37 :: v_dual_sub_f32 v34, v34, v38
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_dual_add_f32 v29, v29, v34 :: v_dual_add_f32 v32, v46, v32
	v_sub_f32_e32 v36, v43, v41
	v_mul_f32_e32 v34, s15, v28
	v_add_f32_e32 v39, v40, v32
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_sub_f32_e32 v37, v43, v36
	v_add_f32_e32 v43, v42, v35
	v_sub_f32_e32 v30, v30, v36
	v_dual_sub_f32 v36, v41, v37 :: v_dual_sub_f32 v37, v39, v40
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v38, v43, v42
	v_dual_add_f32 v30, v30, v36 :: v_dual_sub_f32 v31, v32, v37
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_sub_f32 v32, v35, v38 :: v_dual_mul_f32 v35, 0.5, v24
	v_dual_add_f32 v28, v29, v31 :: v_dual_add_f32 v29, v30, v32
	v_trunc_f32_e32 v31, v24
	v_div_scale_f32 v30, null, v23, v23, s28
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_f32_e32 v32, v39, v28
	v_add_f32_e32 v36, v43, v29
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cmp_eq_f32_e64 s0, v31, v24
	v_sub_f32_e32 v37, v32, v39
	v_mul_f32_e32 v38, v24, v32
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_sub_f32 v39, v36, v43 :: v_dual_mul_f32 v40, v25, v36
	v_fma_f32 v31, v24, v32, -v38
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_sub_f32_e32 v29, v29, v39
	v_fma_f32 v32, v25, v36, -v40
	v_cmp_class_f32_e64 s2, v38, 0x204
	v_trunc_f32_e32 v36, v35
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v32, v25, v29
	v_cmp_neq_f32_e64 s3, v36, v35
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_add_f32_e32 v39, v40, v32
	v_sub_f32_e32 v28, v28, v37
	v_mul_f32_e32 v37, 0.5, v25
	s_and_b32 s0, s0, s3
	v_cmp_eq_f32_e64 s3, |v26|, 1.0
	v_cndmask_b32_e64 v53, 1.0, v26, s0
	v_fmac_f32_e32 v31, v24, v28
	v_trunc_f32_e32 v28, v25
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v29, v38, v31
	v_cndmask_b32_e64 v41, v29, v38, s2
	v_cmp_class_f32_e64 s2, v40, 0x204
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_eq_f32_e64 s4, 0x42b17218, v41
	v_cndmask_b32_e64 v42, v39, v40, s2
	v_cmp_gt_f32_e64 s2, 0, v25
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v35, 0, 0x37000000, s4
	v_cmp_eq_f32_e64 s4, 0x42b17218, v42
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v44, v41, v35
	v_cndmask_b32_e64 v36, 0, 0x37000000, s4
	s_xor_b32 s4, s5, s39
	v_cmp_lt_f32_e64 s5, |v27|, 1.0
	v_cndmask_b32_e64 v43, |v24|, 0, s4
	v_mul_f32_e32 v47, 0x3fb8aa3b, v44
	s_xor_b32 s4, vcc_lo, s1
	v_cmp_eq_f32_e32 vcc_lo, v28, v25
	v_trunc_f32_e32 v28, v37
	s_xor_b32 s5, s41, s5
	v_rndne_f32_e32 v50, v47
	v_fma_f32 v49, 0x3fb8aa3b, v44, -v47
	v_cndmask_b32_e64 v43, v43, |v26|, s3
	v_cmp_neq_f32_e64 s3, v28, v37
	v_cndmask_b32_e64 v45, 0x7f800000, 0, s4
	v_dual_sub_f32 v47, v47, v50 :: v_dual_sub_f32 v46, v42, v36
	v_sub_f32_e32 v29, v29, v38
	v_cndmask_b32_e64 v38, |v25|, 0, s5
	v_cmp_neq_f32_e64 s5, 0x7f800000, |v41|
	v_cmp_eq_f32_e64 s4, 0, v27
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_dual_mul_f32 v48, 0x3fb8aa3b, v46 :: v_dual_sub_f32 v29, v31, v29
	v_fmac_f32_e32 v49, 0x32a5705f, v44
	s_and_b32 vcc_lo, vcc_lo, s3
	s_xor_b32 s2, s2, s4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_4)
	v_fma_f32 v51, 0x3fb8aa3b, v46, -v48
	v_rndne_f32_e32 v52, v48
	v_add_f32_e32 v28, v47, v49
	v_cndmask_b32_e64 v29, 0, v29, s5
	v_cndmask_b32_e64 v41, 0x7f800000, 0, s2
	v_dual_fmac_f32 v51, 0x32a5705f, v46 :: v_dual_sub_f32 v48, v48, v52
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_exp_f32_e32 v28, v28
	v_add_f32_e32 v29, v35, v29
	v_cmp_ngt_f32_e64 s2, 0xc2ce8ed0, v44
	v_cndmask_b32_e64 v35, 0, v26, s0
	v_add_f32_e32 v37, v48, v51
	v_sub_f32_e32 v31, v39, v40
	v_cvt_i32_f32_e32 v39, v50
	v_cvt_i32_f32_e32 v40, v52
	v_cmp_eq_f32_e64 s0, |v27|, 1.0
	v_exp_f32_e32 v37, v37
	v_sub_f32_e32 v31, v32, v31
	v_ldexp_f32 v28, v28, v39
	v_cmp_class_f32_e64 s39, v27, 0x204
	s_delay_alu instid0(VALU_DEP_2)
	v_cndmask_b32_e64 v28, 0, v28, s2
	v_cmp_neq_f32_e64 s2, 0x7f800000, |v42|
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v32, v37, v40
	v_cndmask_b32_e64 v31, 0, v31, s2
	v_cmp_ngt_f32_e64 s2, 0xc2ce8ed0, v46
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f32_e32 v31, v36, v31
	v_cndmask_b32_e64 v32, 0, v32, s2
	v_cmp_nlt_f32_e64 s2, 0x42b17218, v44
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v28, 0x7f800000, v28, s2
	v_cmp_nlt_f32_e64 s2, 0x42b17218, v46
	v_fma_f32 v29, v28, v29, v28
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e64 v32, 0x7f800000, v32, s2
	v_cndmask_b32_e32 v37, 0, v27, vcc_lo
	v_fma_f32 v31, v32, v31, v32
	v_cndmask_b32_e32 v36, 1.0, v27, vcc_lo
	v_cmp_eq_f32_e32 vcc_lo, 0x7f800000, v28
	v_cndmask_b32_e32 v28, v29, v28, vcc_lo
	v_cmp_eq_f32_e32 vcc_lo, 0x7f800000, v32
	s_delay_alu instid0(VALU_DEP_2)
	v_bfi_b32 v28, 0x7fffffff, v28, v53
	v_cndmask_b32_e32 v29, v31, v32, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v24, 0x204
	v_cndmask_b32_e64 v31, v38, |v27|, s0
	v_bfi_b32 v32, 0x7fffffff, v45, v35
	v_bfi_b32 v35, 0x7fffffff, v41, v37
	v_bfi_b32 v29, 0x7fffffff, v29, v36
	v_cndmask_b32_e32 v28, v28, v43, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v25, 0x204
	v_div_scale_f32 v37, s0, s28, v23, s28
	s_delay_alu instid0(VALU_DEP_4)
	v_cndmask_b32_e32 v29, v29, v31, vcc_lo
	s_or_b32 vcc_lo, s1, s40
	v_div_scale_f32 v31, null, v23, v23, s29
	v_cndmask_b32_e32 v28, v28, v32, vcc_lo
	s_or_b32 vcc_lo, s4, s39
	v_rcp_f32_e32 v32, v30
	v_cndmask_b32_e32 v29, v29, v35, vcc_lo
	v_cmp_o_f32_e32 vcc_lo, v26, v24
	v_rcp_f32_e32 v24, v31
	v_div_scale_f32 v39, s1, s29, v23, s29
	v_cndmask_b32_e32 v35, 0x7fc00000, v28, vcc_lo
	v_cmp_o_f32_e32 vcc_lo, v27, v25
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_div_scale_f32 v25, null, v35, v35, v33
	s_waitcnt_depctr 0xfff
	v_fma_f32 v38, -v31, v24, 1.0
	v_rcp_f32_e32 v27, v25
	v_fmac_f32_e32 v24, v38, v24
	v_cndmask_b32_e32 v36, 0x7fc00000, v29, vcc_lo
	v_fma_f32 v29, -v30, v32, 1.0
	v_div_scale_f32 v40, vcc_lo, v33, v35, v33
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v42, v39, v24
	v_fmac_f32_e32 v32, v29, v32
	s_delay_alu instid0(TRANS32_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v29, -v25, v27, 1.0
	v_fma_f32 v46, -v31, v42, v39
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v27, v29, v27
	v_div_scale_f32 v26, null, v36, v36, v34
	v_div_scale_f32 v29, s2, v34, v36, v34
	v_mul_f32_e32 v41, v40, v27
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_rcp_f32_e32 v28, v26
	v_fmac_f32_e32 v42, v46, v24
	v_fma_f32 v45, -v25, v41, v40
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v31, -v31, v42, v39
	v_fmac_f32_e32 v41, v45, v27
	s_waitcnt_depctr 0xfff
	v_fma_f32 v38, -v26, v28, 1.0
	v_fma_f32 v25, -v25, v41, v40
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v28, v38, v28
	v_mul_f32_e32 v38, v37, v32
	v_mul_f32_e32 v43, v29, v28
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v44, -v30, v38, v37
	v_fma_f32 v47, -v26, v43, v29
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v38, v44, v32
	v_fmac_f32_e32 v43, v47, v28
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_fma_f32 v30, -v30, v38, v37
	v_div_fmas_f32 v37, v25, v27, v41
	s_mov_b32 vcc_lo, s2
	v_fma_f32 v26, -v26, v43, v29
	s_delay_alu instid0(VALU_DEP_1)
	v_div_fmas_f32 v39, v26, v28, v43
	s_mov_b32 vcc_lo, s0
	v_div_scale_f32 v40, s0, v20, v23, v20
	v_div_fmas_f32 v28, v30, v32, v38
	s_mov_b32 vcc_lo, s1
	v_div_scale_f32 v30, null, v23, v23, v21
	v_div_fmas_f32 v24, v31, v24, v42
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_div_fixup_f32 v28, v28, v23, s28
	v_div_scale_f32 v43, s1, v21, v23, v21
	v_rcp_f32_e32 v38, v30
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_div_fixup_f32 v26, v24, v23, s29
	v_cvt_f64_f32_e32 v[24:25], v23
	v_cvt_f64_f32_e32 v[28:29], v28
	v_cvt_f64_f32_e32 v[26:27], v26
	s_waitcnt_depctr 0xfff
	v_fma_f32 v42, -v30, v38, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v38, v42, v38
	v_mul_f32_e32 v42, v43, v38
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v46, -v30, v42, v43
	v_fmac_f32_e32 v42, v46, v38
	v_fma_f64 v[26:27], v[11:12], v[24:25], v[26:27]
	v_fma_f64 v[24:25], v[9:10], v[24:25], v[28:29]
	v_div_scale_f32 v28, null, v23, v23, v22
	v_div_scale_f32 v29, null, v23, v23, v20
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v31, v28
	v_rcp_f32_e32 v32, v29
	s_waitcnt_depctr 0xfff
	v_fma_f32 v41, -v28, v31, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v31, v41, v31
	v_cvt_f32_f64_e32 v26, v[26:27]
	v_cvt_f32_f64_e64 v24, -v[24:25]
	v_fma_f32 v25, -v29, v32, 1.0
	v_div_scale_f32 v27, vcc_lo, v22, v23, v22
	v_dual_fmac_f32 v32, v25, v32 :: v_dual_mul_f32 v25, v27, v31
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v41, v40, v32
	v_fma_f32 v44, -v28, v25, v27
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v45, -v29, v41, v40
	v_fmac_f32_e32 v25, v44, v31
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v41, v45, v32
	v_fma_f32 v27, -v28, v25, v27
	s_delay_alu instid0(VALU_DEP_2)
	v_fma_f32 v28, -v29, v41, v40
	v_fma_f32 v29, -v30, v42, v43
	v_div_fixup_f32 v30, v37, v35, v33
	v_div_fixup_f32 v33, v39, v36, v34
	v_div_fmas_f32 v25, v27, v31, v25
	s_mov_b32 vcc_lo, s0
	v_div_fmas_f32 v27, v28, v32, v41
	s_mov_b32 vcc_lo, s1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_4)
	v_div_fixup_f32 v22, v25, v23, v22
	v_div_fmas_f32 v29, v29, v38, v42
	v_mul_f32_e32 v26, s31, v26
	v_mul_f32_e32 v24, s30, v24
	v_div_fixup_f32 v20, v27, v23, v20
	v_div_fixup_f32 v21, v29, v23, v21
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v26, v26, v33
	v_fma_f32 v24, v24, v30, -v26
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mul_f32 v28, s31, v33 :: v_dual_fmac_f32 v19, v22, v24
	v_dual_fmac_f32 v28, s30, v30 :: v_dual_fmac_f32 v17, v21, v24
	v_fmac_f32_e32 v18, v20, v24
	s_delay_alu instid0(VALU_DEP_2)
	v_add_f32_e32 v13, v13, v28
.LBB0_11:                               ; %Flow226
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s38
	s_add_i32 s33, s33, -1
	s_add_u32 s24, s24, 4
	s_addc_u32 s25, s25, 0
	s_add_u32 s26, s26, 4
	v_add_nc_u32_e32 v2, -1, v2
	s_addc_u32 s27, s27, 0
	s_add_u32 s6, s6, 4
	s_addc_u32 s7, s7, 0
	s_cmp_eq_u32 s33, 0
	s_cbranch_scc0 .LBB0_2
.LBB0_12:                               ; %Flow228
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s16, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s17, v1, vcc_lo
	v_add_co_u32 v4, vcc_lo, s18, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, s19, v1, vcc_lo
	v_add_co_u32 v6, vcc_lo, s20, v0
	v_xor_b32_e32 v8, 0x80000000, v19
	v_add_co_ci_u32_e32 v7, vcc_lo, s21, v1, vcc_lo
	v_xor_b32_e32 v9, 0x80000000, v18
	v_add_co_u32 v0, vcc_lo, s22, v0
	v_xor_b32_e32 v10, 0x80000000, v17
	v_add_co_ci_u32_e32 v1, vcc_lo, s23, v1, vcc_lo
	global_store_b32 v[2:3], v8, off
	global_store_b32 v[4:5], v9, off
	global_store_b32 v[6:7], v10, off
	global_store_b32 v[0:1], v13, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 408
		.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 55
		.amdhsa_next_free_sgpr 42
		.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	_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff, .Lfunc_end0-_Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4276
; NumSgprs: 44
; NumVgprs: 55
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 5
; VGPRBlocks: 6
; NumSGPRsForWavesPerEU: 44
; NumVGPRsForWavesPerEU: 55
; 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
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         40
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .offset:         64
        .size:           4
        .value_kind:     by_value
      - .offset:         68
        .size:           4
        .value_kind:     by_value
      - .offset:         72
        .size:           4
        .value_kind:     by_value
      - .offset:         76
        .size:           4
        .value_kind:     by_value
      - .offset:         80
        .size:           4
        .value_kind:     by_value
      - .offset:         84
        .size:           4
        .value_kind:     by_value
      - .offset:         88
        .size:           4
        .value_kind:     by_value
      - .offset:         92
        .size:           4
        .value_kind:     by_value
      - .offset:         96
        .size:           4
        .value_kind:     by_value
      - .offset:         100
        .size:           4
        .value_kind:     by_value
      - .offset:         104
        .size:           4
        .value_kind:     by_value
      - .offset:         108
        .size:           4
        .value_kind:     by_value
      - .offset:         112
        .size:           4
        .value_kind:     by_value
      - .offset:         116
        .size:           4
        .value_kind:     by_value
      - .offset:         120
        .size:           4
        .value_kind:     by_value
      - .offset:         124
        .size:           4
        .value_kind:     by_value
      - .offset:         128
        .size:           4
        .value_kind:     by_value
      - .offset:         132
        .size:           4
        .value_kind:     by_value
      - .offset:         136
        .size:           4
        .value_kind:     by_value
      - .offset:         140
        .size:           4
        .value_kind:     by_value
      - .offset:         144
        .size:           4
        .value_kind:     by_value
      - .offset:         152
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         156
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         160
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         164
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         166
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         168
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         170
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         172
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         174
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         192
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         200
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         208
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         216
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 408
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff
    .private_segment_fixed_size: 0
    .sgpr_count:     44
    .sgpr_spill_count: 0
    .symbol:         _Z8mdKerneliPfS_S_S_S_S_S_iffffffffffffffffffff.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     55
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
