#include "includes.h"
__device__ static float disp_absolute_residual(float Xd, float Yd, float Zd, float Xm, float Ym, float Zm, float nx, float ny, float nz, float T0, float T1, float T2, float R0, float R1, float R2, float fx, float b) {
float r = -Xd * nx + Xm * nx - Yd * ny + Ym * ny - Zd * nz + Zm * nz +
nx * T0 + ny * T1 + nz * T2 + Xm * ny * R2 - Xm * nz * R1 -
Ym * nx * R2 + Ym * nz * R0 + Zm * nx * R1 - Zm * ny * R0;

// weight to convert distance units to pixels
r *= fx * b / (Zm * Zm);

return fabsf(r);
}
__global__ void normal_eqs_disparity_weighted_GPU( float *d_CD, const float *d_disparity_compact, const float4 *d_Zbuffer_normals_compact, const int *d_ind_disparity_Zbuffer, float fx, float fy, float ox, float oy, float b, int n_cols, const int *d_n_values_disparity, const int *d_start_ind_disparity, const float *d_abs_res_scales, float w_disp, const float *d_dTR) {

int n_val_accum =
gridDim.x * blockDim.x; // n_val_accum may not be multiple of blocksize

int n_disparity = d_n_values_disparity[blockIdx.y];
int n_accum = (int)ceilf((float)n_disparity / (float)n_val_accum);
int start_ind = d_start_ind_disparity[blockIdx.y];

// initialize accumulators

float A0 = 0.0f, A1 = 0.0f, A2 = 0.0f, A3 = 0.0f, A4 = 0.0f, A5 = 0.0f,
A6 = 0.0f, A7 = 0.0f, A8 = 0.0f, A9 = 0.0f, A10 = 0.0f, A11 = 0.0f,
A12 = 0.0f, A13 = 0.0f, A14 = 0.0f, A15 = 0.0f, A16 = 0.0f, A17 = 0.0f,
A18 = 0.0f, A19 = 0.0f, A20 = 0.0f, A21 = 0.0f, A22 = 0.0f, A23 = 0.0f,
A24 = 0.0f, A25 = 0.0f, A26 = 0.0f;

for (int in_ind = blockDim.x * blockIdx.x * n_accum + threadIdx.x;
in_ind < blockDim.x * (blockIdx.x + 1) * n_accum; in_ind += blockDim.x) {

if (in_ind < n_disparity) { // is this a valid sample?

// fetch disparity, Zbuffer and normal from global memory
float disp = d_disparity_compact[in_ind + start_ind];
float4 tmp = d_Zbuffer_normals_compact[in_ind + start_ind];
float Zbuffer = tmp.x;
float nx = tmp.y;
float ny = tmp.z;
float nz = tmp.w;

// compute coordinates
int pixel_ind = d_ind_disparity_Zbuffer[in_ind + start_ind];

float y = floorf(__fdividef((float)pixel_ind, n_cols));
float x = (float)pixel_ind - y * n_cols;

x = __fdividef((x - ox), fx);
y = __fdividef((y - oy), fy);

// reconstruct 3D point from disparity

float Zd = -(fx * b) / disp; // arbitrary use of fx
float Xd = x * Zd;
float Yd = y * Zd;

// reconstruct 3D point from model

float Zm = Zbuffer;
float Xm = x * Zm;
float Ym = y * Zm;

// determine M-estimation weight
// disparity residual weighed by rel. importance disp vs flow
int s6 = blockIdx.y * 6;
float w = w_disp * disp_absolute_residual(
Xd, Yd, Zd, Xm, Ym, Zm, nx, ny, nz, d_dTR[s6],
d_dTR[s6 + 1], d_dTR[s6 + 2], d_dTR[s6 + 3],
d_dTR[s6 + 4], d_dTR[s6 + 5], fx, b);
w /= d_abs_res_scales[blockIdx.y];
w = (w > 1) ? 0 : (1.0f - 2.0f * w * w + w * w * w * w);

// multiply m estimation weight with distance->pixel conversion weight
// (squared)
w *= (fx * fx * b * b) / (Zm * Zm * Zm * Zm);

/************************/
/* evaluate constraints */
/************************/

// unique values A-matrix

A0 += w * (nx * nx);
A1 += w * (nx * ny);
A2 += w * (nx * nz);
A3 += w * (Ym * nx * nz - Zm * nx * ny);
A4 += w * (Zm * (nx * nx) - Xm * nx * nz);
A5 += w * (-Ym * (nx * nx) + Xm * nx * ny);

A6 += w * (ny * ny);
A7 += w * (ny * nz);
A8 += w * (-Zm * (ny * ny) + Ym * ny * nz);
A9 += w * (-Xm * ny * nz + Zm * nx * ny);
A10 += w * (Xm * (ny * ny) - Ym * nx * ny);

A11 += w * (nz * nz);
A12 += w * (Ym * (nz * nz) - Zm * ny * nz);
A13 += w * (-Xm * (nz * nz) + Zm * nx * nz);
A14 += w * (Xm * ny * nz - Ym * nx * nz);

A15 += w * ((Ym * Ym) * (nz * nz) + (Zm * Zm) * (ny * ny) -
Ym * Zm * ny * nz * 2.0f);
A16 += w * (-Xm * Ym * (nz * nz) - (Zm * Zm) * nx * ny +
Xm * Zm * ny * nz + Ym * Zm * nx * nz);
A17 += w * (-Xm * Zm * (ny * ny) - (Ym * Ym) * nx * nz +
Xm * Ym * ny * nz + Ym * Zm * nx * ny);

A18 += w * ((Xm * Xm) * (nz * nz) + (Zm * Zm) * (nx * nx) -
Xm * Zm * nx * nz * 2.0f);
A19 += w * (-Ym * Zm * (nx * nx) - (Xm * Xm) * ny * nz +
Xm * Ym * nx * nz + Xm * Zm * nx * ny);

A20 += w * ((Xm * Xm) * (ny * ny) + (Ym * Ym) * (nx * nx) -
Xm * Ym * nx * ny * 2.0f);

// B-vector

A21 += w * (Xd * (nx * nx) - Xm * (nx * nx) + Yd * nx * ny -
Ym * nx * ny + Zd * nx * nz - Zm * nx * nz);
A22 += w * (Yd * (ny * ny) - Ym * (ny * ny) + Xd * nx * ny -
Xm * nx * ny + Zd * ny * nz - Zm * ny * nz);
A23 += w * (Zd * (nz * nz) - Zm * (nz * nz) + Xd * nx * nz -
Xm * nx * nz + Yd * ny * nz - Ym * ny * nz);
A24 += w *
(-Yd * Zm * (ny * ny) + Ym * Zd * (nz * nz) + Ym * Zm * (ny * ny) -
Ym * Zm * (nz * nz) - (Ym * Ym) * ny * nz + (Zm * Zm) * ny * nz +
Xd * Ym * nx * nz - Xm * Ym * nx * nz - Xd * Zm * nx * ny +
Yd * Ym * ny * nz + Xm * Zm * nx * ny - Zd * Zm * ny * nz);
A25 +=
w * (Xd * Zm * (nx * nx) - Xm * Zd * (nz * nz) - Xm * Zm * (nx * nx) +
Xm * Zm * (nz * nz) + (Xm * Xm) * nx * nz - (Zm * Zm) * nx * nz -
Xd * Xm * nx * nz - Xm * Yd * ny * nz + Xm * Ym * ny * nz +
Yd * Zm * nx * ny - Ym * Zm * nx * ny + Zd * Zm * nx * nz);
A26 += w *
(-Xd * Ym * (nx * nx) + Xm * Yd * (ny * ny) + Xm * Ym * (nx * nx) -
Xm * Ym * (ny * ny) - (Xm * Xm) * nx * ny + (Ym * Ym) * nx * ny +
Xd * Xm * nx * ny - Yd * Ym * nx * ny + Xm * Zd * ny * nz -
Xm * Zm * ny * nz - Ym * Zd * nx * nz + Ym * Zm * nx * nz);
}
}

/**************************/
/* write out accumulators */
/**************************/

int out_ind =
27 * n_val_accum * blockIdx.y + blockDim.x * blockIdx.x + threadIdx.x;

w_disp *= w_disp; // weight relative to flow

d_CD[out_ind] = w_disp * A0;
d_CD[out_ind + n_val_accum] = w_disp * A1;
d_CD[out_ind + 2 * n_val_accum] = w_disp * A2;
d_CD[out_ind + 3 * n_val_accum] = w_disp * A3;
d_CD[out_ind + 4 * n_val_accum] = w_disp * A4;
d_CD[out_ind + 5 * n_val_accum] = w_disp * A5;
d_CD[out_ind + 6 * n_val_accum] = w_disp * A6;
d_CD[out_ind + 7 * n_val_accum] = w_disp * A7;
d_CD[out_ind + 8 * n_val_accum] = w_disp * A8;
d_CD[out_ind + 9 * n_val_accum] = w_disp * A9;
d_CD[out_ind + 10 * n_val_accum] = w_disp * A10;
d_CD[out_ind + 11 * n_val_accum] = w_disp * A11;
d_CD[out_ind + 12 * n_val_accum] = w_disp * A12;
d_CD[out_ind + 13 * n_val_accum] = w_disp * A13;
d_CD[out_ind + 14 * n_val_accum] = w_disp * A14;
d_CD[out_ind + 15 * n_val_accum] = w_disp * A15;
d_CD[out_ind + 16 * n_val_accum] = w_disp * A16;
d_CD[out_ind + 17 * n_val_accum] = w_disp * A17;
d_CD[out_ind + 18 * n_val_accum] = w_disp * A18;
d_CD[out_ind + 19 * n_val_accum] = w_disp * A19;
d_CD[out_ind + 20 * n_val_accum] = w_disp * A20;
d_CD[out_ind + 21 * n_val_accum] = w_disp * A21;
d_CD[out_ind + 22 * n_val_accum] = w_disp * A22;
d_CD[out_ind + 23 * n_val_accum] = w_disp * A23;
d_CD[out_ind + 24 * n_val_accum] = w_disp * A24;
d_CD[out_ind + 25 * n_val_accum] = w_disp * A25;
d_CD[out_ind + 26 * n_val_accum] = w_disp * A26;
}
	.file	"tmpxft_0033ebde_00000000-6_normal_eqs_disparity_weighted_GPU.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	_Z87__device_stub__Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_PfPKfPK6float4PKifffffiS6_S6_S1_fS1_
	.type	_Z87__device_stub__Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_PfPKfPK6float4PKifffffiS6_S6_S1_fS1_, @function
_Z87__device_stub__Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_PfPKfPK6float4PKifffffiS6_S6_S1_fS1_:
.LFB2033:
	.cfi_startproc
	endbr64
	subq	$296, %rsp
	.cfi_def_cfa_offset 304
	movq	304(%rsp), %rax
	movq	%rdi, 88(%rsp)
	leaq	112(%rsp), %rdi
	movq	%rsi, 80(%rsp)
	leaq	124(%rsp), %rsi
	movq	%rax, 24(%rsp)
	movq	312(%rsp), %rax
	movq	%rdx, 72(%rsp)
	leaq	96(%rsp), %rdx
	movq	%rax, 16(%rsp)
	movq	320(%rsp), %rax
	movq	%rcx, 64(%rsp)
	leaq	104(%rsp), %rcx
	movl	%r8d, 40(%rsp)
	movq	%r9, 32(%rsp)
	movq	%rax, (%rsp)
	movss	%xmm0, 60(%rsp)
	movss	%xmm1, 56(%rsp)
	movss	%xmm2, 52(%rsp)
	movss	%xmm3, 48(%rsp)
	movss	%xmm4, 44(%rsp)
	movss	%xmm5, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 280(%rsp)
	xorl	%eax, %eax
	leaq	88(%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	60(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	52(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	44(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 264(%rsp)
	movq	%rsp, %rax
	movq	%rax, 272(%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	.L2
	pushq	104(%rsp)
	.cfi_def_cfa_offset 312
	leaq	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_(%rip), %rdi
	pushq	104(%rsp)
	.cfi_def_cfa_offset 320
	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 312
	popq	%rdx
	.cfi_def_cfa_offset 304
.L2:
	movq	280(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$296, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2033:
	.size	_Z87__device_stub__Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_PfPKfPK6float4PKifffffiS6_S6_S1_fS1_, .-_Z87__device_stub__Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_PfPKfPK6float4PKifffffiS6_S6_S1_fS1_
	.globl	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_
	.type	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_, @function
_Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_:
.LFB2034:
	.cfi_startproc
	endbr64
	jmp	_Z87__device_stub__Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_PfPKfPK6float4PKifffffiS6_S6_S1_fS1_
	.cfi_endproc
.LFE2034:
	.size	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_, .-_Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_"
	.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	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK6float4PKifffffiS6_S6_S1_fS1_(%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 "includes.h"
__device__ static float disp_absolute_residual(float Xd, float Yd, float Zd, float Xm, float Ym, float Zm, float nx, float ny, float nz, float T0, float T1, float T2, float R0, float R1, float R2, float fx, float b) {
float r = -Xd * nx + Xm * nx - Yd * ny + Ym * ny - Zd * nz + Zm * nz +
nx * T0 + ny * T1 + nz * T2 + Xm * ny * R2 - Xm * nz * R1 -
Ym * nx * R2 + Ym * nz * R0 + Zm * nx * R1 - Zm * ny * R0;

// weight to convert distance units to pixels
r *= fx * b / (Zm * Zm);

return fabsf(r);
}
__global__ void normal_eqs_disparity_weighted_GPU( float *d_CD, const float *d_disparity_compact, const float4 *d_Zbuffer_normals_compact, const int *d_ind_disparity_Zbuffer, float fx, float fy, float ox, float oy, float b, int n_cols, const int *d_n_values_disparity, const int *d_start_ind_disparity, const float *d_abs_res_scales, float w_disp, const float *d_dTR) {

int n_val_accum =
gridDim.x * blockDim.x; // n_val_accum may not be multiple of blocksize

int n_disparity = d_n_values_disparity[blockIdx.y];
int n_accum = (int)ceilf((float)n_disparity / (float)n_val_accum);
int start_ind = d_start_ind_disparity[blockIdx.y];

// initialize accumulators

float A0 = 0.0f, A1 = 0.0f, A2 = 0.0f, A3 = 0.0f, A4 = 0.0f, A5 = 0.0f,
A6 = 0.0f, A7 = 0.0f, A8 = 0.0f, A9 = 0.0f, A10 = 0.0f, A11 = 0.0f,
A12 = 0.0f, A13 = 0.0f, A14 = 0.0f, A15 = 0.0f, A16 = 0.0f, A17 = 0.0f,
A18 = 0.0f, A19 = 0.0f, A20 = 0.0f, A21 = 0.0f, A22 = 0.0f, A23 = 0.0f,
A24 = 0.0f, A25 = 0.0f, A26 = 0.0f;

for (int in_ind = blockDim.x * blockIdx.x * n_accum + threadIdx.x;
in_ind < blockDim.x * (blockIdx.x + 1) * n_accum; in_ind += blockDim.x) {

if (in_ind < n_disparity) { // is this a valid sample?

// fetch disparity, Zbuffer and normal from global memory
float disp = d_disparity_compact[in_ind + start_ind];
float4 tmp = d_Zbuffer_normals_compact[in_ind + start_ind];
float Zbuffer = tmp.x;
float nx = tmp.y;
float ny = tmp.z;
float nz = tmp.w;

// compute coordinates
int pixel_ind = d_ind_disparity_Zbuffer[in_ind + start_ind];

float y = floorf(__fdividef((float)pixel_ind, n_cols));
float x = (float)pixel_ind - y * n_cols;

x = __fdividef((x - ox), fx);
y = __fdividef((y - oy), fy);

// reconstruct 3D point from disparity

float Zd = -(fx * b) / disp; // arbitrary use of fx
float Xd = x * Zd;
float Yd = y * Zd;

// reconstruct 3D point from model

float Zm = Zbuffer;
float Xm = x * Zm;
float Ym = y * Zm;

// determine M-estimation weight
// disparity residual weighed by rel. importance disp vs flow
int s6 = blockIdx.y * 6;
float w = w_disp * disp_absolute_residual(
Xd, Yd, Zd, Xm, Ym, Zm, nx, ny, nz, d_dTR[s6],
d_dTR[s6 + 1], d_dTR[s6 + 2], d_dTR[s6 + 3],
d_dTR[s6 + 4], d_dTR[s6 + 5], fx, b);
w /= d_abs_res_scales[blockIdx.y];
w = (w > 1) ? 0 : (1.0f - 2.0f * w * w + w * w * w * w);

// multiply m estimation weight with distance->pixel conversion weight
// (squared)
w *= (fx * fx * b * b) / (Zm * Zm * Zm * Zm);

/************************/
/* evaluate constraints */
/************************/

// unique values A-matrix

A0 += w * (nx * nx);
A1 += w * (nx * ny);
A2 += w * (nx * nz);
A3 += w * (Ym * nx * nz - Zm * nx * ny);
A4 += w * (Zm * (nx * nx) - Xm * nx * nz);
A5 += w * (-Ym * (nx * nx) + Xm * nx * ny);

A6 += w * (ny * ny);
A7 += w * (ny * nz);
A8 += w * (-Zm * (ny * ny) + Ym * ny * nz);
A9 += w * (-Xm * ny * nz + Zm * nx * ny);
A10 += w * (Xm * (ny * ny) - Ym * nx * ny);

A11 += w * (nz * nz);
A12 += w * (Ym * (nz * nz) - Zm * ny * nz);
A13 += w * (-Xm * (nz * nz) + Zm * nx * nz);
A14 += w * (Xm * ny * nz - Ym * nx * nz);

A15 += w * ((Ym * Ym) * (nz * nz) + (Zm * Zm) * (ny * ny) -
Ym * Zm * ny * nz * 2.0f);
A16 += w * (-Xm * Ym * (nz * nz) - (Zm * Zm) * nx * ny +
Xm * Zm * ny * nz + Ym * Zm * nx * nz);
A17 += w * (-Xm * Zm * (ny * ny) - (Ym * Ym) * nx * nz +
Xm * Ym * ny * nz + Ym * Zm * nx * ny);

A18 += w * ((Xm * Xm) * (nz * nz) + (Zm * Zm) * (nx * nx) -
Xm * Zm * nx * nz * 2.0f);
A19 += w * (-Ym * Zm * (nx * nx) - (Xm * Xm) * ny * nz +
Xm * Ym * nx * nz + Xm * Zm * nx * ny);

A20 += w * ((Xm * Xm) * (ny * ny) + (Ym * Ym) * (nx * nx) -
Xm * Ym * nx * ny * 2.0f);

// B-vector

A21 += w * (Xd * (nx * nx) - Xm * (nx * nx) + Yd * nx * ny -
Ym * nx * ny + Zd * nx * nz - Zm * nx * nz);
A22 += w * (Yd * (ny * ny) - Ym * (ny * ny) + Xd * nx * ny -
Xm * nx * ny + Zd * ny * nz - Zm * ny * nz);
A23 += w * (Zd * (nz * nz) - Zm * (nz * nz) + Xd * nx * nz -
Xm * nx * nz + Yd * ny * nz - Ym * ny * nz);
A24 += w *
(-Yd * Zm * (ny * ny) + Ym * Zd * (nz * nz) + Ym * Zm * (ny * ny) -
Ym * Zm * (nz * nz) - (Ym * Ym) * ny * nz + (Zm * Zm) * ny * nz +
Xd * Ym * nx * nz - Xm * Ym * nx * nz - Xd * Zm * nx * ny +
Yd * Ym * ny * nz + Xm * Zm * nx * ny - Zd * Zm * ny * nz);
A25 +=
w * (Xd * Zm * (nx * nx) - Xm * Zd * (nz * nz) - Xm * Zm * (nx * nx) +
Xm * Zm * (nz * nz) + (Xm * Xm) * nx * nz - (Zm * Zm) * nx * nz -
Xd * Xm * nx * nz - Xm * Yd * ny * nz + Xm * Ym * ny * nz +
Yd * Zm * nx * ny - Ym * Zm * nx * ny + Zd * Zm * nx * nz);
A26 += w *
(-Xd * Ym * (nx * nx) + Xm * Yd * (ny * ny) + Xm * Ym * (nx * nx) -
Xm * Ym * (ny * ny) - (Xm * Xm) * nx * ny + (Ym * Ym) * nx * ny +
Xd * Xm * nx * ny - Yd * Ym * nx * ny + Xm * Zd * ny * nz -
Xm * Zm * ny * nz - Ym * Zd * nx * nz + Ym * Zm * nx * nz);
}
}

/**************************/
/* write out accumulators */
/**************************/

int out_ind =
27 * n_val_accum * blockIdx.y + blockDim.x * blockIdx.x + threadIdx.x;

w_disp *= w_disp; // weight relative to flow

d_CD[out_ind] = w_disp * A0;
d_CD[out_ind + n_val_accum] = w_disp * A1;
d_CD[out_ind + 2 * n_val_accum] = w_disp * A2;
d_CD[out_ind + 3 * n_val_accum] = w_disp * A3;
d_CD[out_ind + 4 * n_val_accum] = w_disp * A4;
d_CD[out_ind + 5 * n_val_accum] = w_disp * A5;
d_CD[out_ind + 6 * n_val_accum] = w_disp * A6;
d_CD[out_ind + 7 * n_val_accum] = w_disp * A7;
d_CD[out_ind + 8 * n_val_accum] = w_disp * A8;
d_CD[out_ind + 9 * n_val_accum] = w_disp * A9;
d_CD[out_ind + 10 * n_val_accum] = w_disp * A10;
d_CD[out_ind + 11 * n_val_accum] = w_disp * A11;
d_CD[out_ind + 12 * n_val_accum] = w_disp * A12;
d_CD[out_ind + 13 * n_val_accum] = w_disp * A13;
d_CD[out_ind + 14 * n_val_accum] = w_disp * A14;
d_CD[out_ind + 15 * n_val_accum] = w_disp * A15;
d_CD[out_ind + 16 * n_val_accum] = w_disp * A16;
d_CD[out_ind + 17 * n_val_accum] = w_disp * A17;
d_CD[out_ind + 18 * n_val_accum] = w_disp * A18;
d_CD[out_ind + 19 * n_val_accum] = w_disp * A19;
d_CD[out_ind + 20 * n_val_accum] = w_disp * A20;
d_CD[out_ind + 21 * n_val_accum] = w_disp * A21;
d_CD[out_ind + 22 * n_val_accum] = w_disp * A22;
d_CD[out_ind + 23 * n_val_accum] = w_disp * A23;
d_CD[out_ind + 24 * n_val_accum] = w_disp * A24;
d_CD[out_ind + 25 * n_val_accum] = w_disp * A25;
d_CD[out_ind + 26 * n_val_accum] = w_disp * A26;
}
	.text
	.file	"normal_eqs_disparity_weighted_GPU.hip"
	.globl	_Z48__device_stub__normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_ # -- Begin function _Z48__device_stub__normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
	.type	_Z48__device_stub__normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_,@function
_Z48__device_stub__normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_: # @_Z48__device_stub__normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
	.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	$256, %rsp                      # imm = 0x100
	.cfi_def_cfa_offset 304
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	88(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	80(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	72(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	64(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	36(%rsp), %rcx
	movss	%xmm0, (%rcx)
	leaq	32(%rsp), %r10
	movss	%xmm1, (%r10)
	leaq	28(%rsp), %r11
	movss	%xmm2, (%r11)
	leaq	24(%rsp), %r14
	movss	%xmm3, (%r14)
	leaq	20(%rsp), %r15
	movss	%xmm4, (%r15)
	leaq	16(%rsp), %r12
	movl	%r8d, (%r12)
	leaq	56(%rsp), %r8
	movq	%r9, (%r8)
	leaq	12(%rsp), %r9
	movss	%xmm5, (%r9)
	leaq	128(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r10, 40(%rbx)
	movq	%r11, 48(%rbx)
	movq	%r14, 56(%rbx)
	movq	%r15, 64(%rbx)
	movq	%r12, 72(%rbx)
	movq	%r8, 80(%rbx)
	leaq	304(%rsp), %rax
	movq	%rax, 88(%rbx)
	leaq	312(%rsp), %rax
	movq	%rax, 96(%rbx)
	movq	%r9, 104(%rbx)
	leaq	320(%rsp), %rax
	movq	%rax, 112(%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	$_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$272, %rsp                      # imm = 0x110
	.cfi_adjust_cfa_offset -272
	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	_Z48__device_stub__normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_, .Lfunc_end0-_Z48__device_stub__normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
	.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	$_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_, %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	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_,@object # @_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
	.section	.rodata,"a",@progbits
	.globl	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
	.p2align	3, 0x0
_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_:
	.quad	_Z48__device_stub__normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
	.size	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_, 8

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

	.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 _Z48__device_stub__normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_ ; -- Begin function _Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
	.globl	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
	.p2align	8
	.type	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_,@function
_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_: ; @_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x6c
	s_load_b128 s[16:19], s[0:1], 0x38
	s_mov_b32 s22, s15
	s_mov_b32 s23, 0
	s_clause 0x1
	s_load_b32 s21, s[0:1], 0x50
	s_load_b32 s6, s[0:1], 0x60
	s_lshl_b64 s[2:3], s[22:23], 2
	v_dual_mov_b32 v29, 0 :: v_dual_mov_b32 v28, 0
	v_dual_mov_b32 v27, 0 :: v_dual_mov_b32 v26, 0
	v_dual_mov_b32 v25, 0 :: v_dual_mov_b32 v24, 0
	v_dual_mov_b32 v23, 0 :: v_dual_mov_b32 v22, 0
	v_dual_mov_b32 v21, 0 :: v_dual_mov_b32 v20, 0
	v_dual_mov_b32 v19, 0 :: v_dual_mov_b32 v18, 0
	v_dual_mov_b32 v17, 0 :: v_dual_mov_b32 v16, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s31, s4, 0xffff
	s_add_u32 s4, s16, s2
	s_addc_u32 s5, s17, s3
	s_mul_i32 s30, s14, s31
	s_load_b32 s33, s[4:5], 0x0
	s_mul_i32 s20, s6, s31
	s_load_b256 s[4:11], s[0:1], 0x0
	v_cvt_f32_i32_e32 v1, s20
	s_add_i32 s12, s30, s31
	v_dual_mov_b32 v15, 0 :: v_dual_mov_b32 v14, 0
	v_dual_mov_b32 v8, 0 :: v_dual_mov_b32 v9, 0
	v_dual_mov_b32 v10, 0 :: v_dual_mov_b32 v11, 0
	v_dual_mov_b32 v12, 0 :: v_dual_mov_b32 v13, 0
	s_mov_b32 s34, exec_lo
	v_dual_mov_b32 v31, 0 :: v_dual_mov_b32 v30, 0
	v_mov_b32_e32 v32, 0
	s_waitcnt lgkmcnt(0)
	v_cvt_f32_i32_e32 v2, s33
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v3, null, v1, v1, v2
	v_div_scale_f32 v6, vcc_lo, v2, v1, v2
	v_rcp_f32_e32 v4, v3
	s_waitcnt_depctr 0xfff
	v_fma_f32 v5, -v3, v4, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v4, v5, v4
	v_mul_f32_e32 v5, v6, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v7, -v3, v5, v6
	v_fmac_f32_e32 v5, v7, v4
	v_mov_b32_e32 v7, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v3, -v3, v5, v6
	v_div_fmas_f32 v3, v3, v4, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v1, v3, v1, v2
	v_ceil_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_i32_f32_e32 v1, v1
	v_mad_u64_u32 v[5:6], null, s30, v1, v[0:1]
	v_mul_lo_u32 v33, s12, v1
	v_mov_b32_e32 v6, 0
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_lt_u32_e64 v5, v33
	s_cbranch_execz .LBB0_6
; %bb.1:                                ; %.lr.ph
	s_clause 0x3
	s_load_b128 s[12:15], s[0:1], 0x20
	s_load_b64 s[16:17], s[0:1], 0x30
	s_load_b64 s[26:27], s[0:1], 0x48
	s_load_b64 s[0:1], s[0:1], 0x58
	s_add_u32 s18, s18, s2
	s_addc_u32 s19, s19, s3
	s_mul_i32 s24, s22, 6
	s_load_b32 s35, s[18:19], 0x0
	s_ashr_i32 s25, s24, 31
	v_dual_mov_b32 v30, 0 :: v_dual_mov_b32 v31, 0
	s_lshl_b64 s[18:19], s[24:25], 2
	v_dual_mov_b32 v32, 0 :: v_dual_mov_b32 v29, 0
	v_dual_mov_b32 v28, 0 :: v_dual_mov_b32 v27, 0
	v_dual_mov_b32 v26, 0 :: v_dual_mov_b32 v25, 0
	v_dual_mov_b32 v24, 0 :: v_dual_mov_b32 v23, 0
	s_waitcnt lgkmcnt(0)
	v_mul_f32_e64 v1, s12, s12
	v_dual_mov_b32 v22, 0 :: v_dual_mov_b32 v21, 0
	v_dual_mov_b32 v20, 0 :: v_dual_mov_b32 v19, 0
	s_delay_alu instid0(VALU_DEP_3)
	v_dual_mul_f32 v1, s16, v1 :: v_dual_mov_b32 v18, 0
	v_cvt_f32_i32_e32 v34, s17
	v_mul_f32_e64 v35, -s12, s16
	s_add_u32 s24, s0, s18
	v_mul_f32_e64 v36, s12, s16
	v_dual_mul_f32 v37, s16, v1 :: v_dual_mov_b32 v16, 0
	v_dual_mov_b32 v17, 0 :: v_dual_mov_b32 v14, 0
	v_dual_mov_b32 v15, 0 :: v_dual_mov_b32 v6, 0
	v_dual_mov_b32 v7, 0 :: v_dual_mov_b32 v8, 0
	v_dual_mov_b32 v9, 0 :: v_dual_mov_b32 v10, 0
	v_dual_mov_b32 v11, 0 :: v_dual_mov_b32 v12, 0
	v_mov_b32_e32 v13, 0
	s_addc_u32 s25, s1, s19
	s_add_u32 s26, s26, s2
	s_addc_u32 s27, s27, s3
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	s_mov_b32 s36, exec_lo
	v_cmpx_gt_i32_e64 s33, v5
	s_cbranch_execz .LBB0_4
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	v_add_nc_u32_e32 v1, s35, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v2, 31, v1
	v_lshlrev_b64 v[3:4], 2, v[1:2]
	v_lshlrev_b64 v[1:2], 4, v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v38, vcc_lo, s10, v3
	v_add_co_ci_u32_e32 v39, vcc_lo, s11, v4, vcc_lo
	v_add_co_u32 v3, vcc_lo, s6, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s7, v4, vcc_lo
	global_load_b32 v38, v[38:39], off
	v_add_co_u32 v1, vcc_lo, s8, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s9, v2, vcc_lo
	global_load_b32 v46, v[3:4], off
	global_load_b128 v[1:4], v[1:2], off
	s_clause 0x1
	s_load_b128 s[16:19], s[24:25], 0x0
	s_load_b64 s[28:29], s[24:25], 0x10
	s_load_b32 s37, s[26:27], 0x0
	s_waitcnt vmcnt(2)
	v_cvt_f32_i32_e32 v47, v38
	s_waitcnt vmcnt(1)
	v_div_scale_f32 v48, null, v46, v46, v35
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_scale_f32 v38, null, v34, v34, v47
	v_div_scale_f32 v41, vcc_lo, v47, v34, v47
	v_rcp_f32_e32 v52, v48
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_rcp_f32_e32 v39, v38
	v_div_scale_f32 v49, s1, v35, v46, v35
	s_waitcnt vmcnt(0)
	v_dual_mul_f32 v50, v1, v1 :: v_dual_mul_f32 v43, v3, v3
	v_mul_f32_e32 v44, v3, v4
	v_div_scale_f32 v53, null, v50, v50, v36
	s_waitcnt_depctr 0xfff
	v_fma_f32 v40, -v38, v39, 1.0
	v_fma_f32 v66, -v48, v52, 1.0
	v_mul_f32_e32 v55, v1, v50
	v_rcp_f32_e32 v62, v53
	v_div_scale_f32 v54, s2, v36, v50, v36
	v_fmac_f32_e32 v39, v40, v39
	v_mul_f32_e32 v59, v50, v2
	v_fmac_f32_e32 v52, v66, v52
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v40, v41, v39
	s_waitcnt_depctr 0xfff
	v_fma_f32 v66, -v53, v62, 1.0
	v_mul_f32_e32 v55, v1, v55
	v_fma_f32 v42, -v38, v40, v41
	v_fmac_f32_e32 v62, v66, v62
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_div_scale_f32 v67, null, v55, v55, v37
	v_fmac_f32_e32 v40, v42, v39
	v_mul_f32_e32 v42, v2, v4
	v_div_scale_f32 v68, s0, v37, v55, v37
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_rcp_f32_e32 v70, v67
	v_fma_f32 v38, -v38, v40, v41
	v_mul_f32_e32 v76, v54, v62
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(TRANS32_DEP_1)
	v_div_fmas_f32 v41, v38, v39, v40
	v_mul_f32_e32 v38, v1, v3
	v_mul_f32_e32 v39, v1, v2
	v_mul_f32_e32 v40, v2, v2
	v_fma_f32 v73, -v67, v70, 1.0
	v_div_fixup_f32 v45, v41, v34, v47
	v_mul_f32_e32 v41, v2, v3
	v_mul_f32_e32 v56, v39, v3
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v70, v73, v70
	v_floor_f32_e32 v51, v45
	v_mul_f32_e32 v45, v4, v4
	v_dual_mul_f32 v57, v1, v43 :: v_dual_mul_f32 v58, v38, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v47, -v51, v34, v47
	v_dual_subrev_f32 v51, s15, v51 :: v_dual_mul_f32 v64, v1, v45
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_subrev_f32_e32 v47, s14, v47
	v_div_scale_f32 v61, null, s13, s13, v51
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_scale_f32 v60, null, s12, s12, v47
	v_rcp_f32_e32 v65, v61
	v_div_scale_f32 v72, vcc_lo, v47, s12, v47
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v63, v60
	s_waitcnt_depctr 0xfff
	v_fma_f32 v71, -v61, v65, 1.0
	v_fma_f32 v69, -v60, v63, 1.0
	v_fmac_f32_e32 v63, v69, v63
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v65, v71, v65
	v_mul_f32_e32 v71, v49, v52
	v_div_scale_f32 v69, s3, v51, s13, v51
	v_mul_f32_e32 v66, v72, v63
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v75, -v48, v71, v49
	v_mul_f32_e32 v74, v69, v65
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v77, -v60, v66, v72
	v_fmac_f32_e32 v71, v75, v52
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_fma_f32 v73, -v61, v74, v69
	v_fma_f32 v75, -v53, v76, v54
	v_fmac_f32_e32 v66, v77, v63
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v48, -v48, v71, v49
	v_fmac_f32_e32 v76, v75, v62
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v49, -v60, v66, v72
	v_div_fmas_f32 v49, v49, v63, v66
	s_mov_b32 vcc_lo, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v47, v49, s12, v47
	v_mul_f32_e32 v49, v50, v3
	v_mul_f32_e32 v77, v68, v70
	v_fma_f32 v60, -v67, v77, v68
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v74, v73, v65 :: v_dual_fmac_f32 v77, v60, v70
	v_fma_f32 v61, -v61, v74, v69
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v60, -v67, v77, v68
	v_div_fmas_f32 v61, v61, v65, v74
	s_mov_b32 vcc_lo, s1
	v_div_fmas_f32 v48, v48, v52, v71
	v_fma_f32 v52, -v53, v76, v54
	s_delay_alu instid0(VALU_DEP_3)
	v_div_fixup_f32 v51, v61, s13, v51
	s_mov_b32 vcc_lo, s2
	v_mul_f32_e32 v54, v59, v3
	v_div_fixup_f32 v46, v48, v46, v35
	v_mul_f32_e32 v48, v1, v47
	v_div_fmas_f32 v52, v52, v62, v76
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v63, v46, v2
	v_dual_mul_f32 v47, v46, v47 :: v_dual_mul_f32 v68, v4, v48
	v_mul_f32_e32 v53, v2, v48
	v_dual_mul_f32 v61, v46, v51 :: v_dual_mul_f32 v72, v48, v48
	v_dual_mul_f32 v51, v1, v51 :: v_dual_mul_f32 v66, v46, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_fma_f32 v62, -v2, v47, v53
	v_mul_f32_e32 v73, v40, v48
	v_mul_f32_e32 v91, v45, v72
	s_delay_alu instid0(VALU_DEP_4)
	v_mul_f32_e32 v80, v1, v51
	v_mul_f32_e32 v84, v43, v51
	v_fma_f32 v62, -v3, v61, v62
	v_mul_f32_e32 v69, v1, v48
	v_dual_mul_f32 v67, v3, v48 :: v_dual_mul_f32 v76, v2, v51
	v_dual_mul_f32 v78, v4, v51 :: v_dual_mul_f32 v87, v2, v61
	s_delay_alu instid0(VALU_DEP_4)
	v_fmac_f32_e32 v62, v3, v51
	v_mul_f32_e64 v81, v51, -v48
	v_fma_f32 v73, v40, v47, -v73
	v_fma_f32 v64, v46, v45, -v64
	v_dual_mul_f32 v74, v46, v48 :: v_dual_mul_f32 v75, v3, v51
	v_fma_f32 v62, -v46, v4, v62
	v_mul_f32_e32 v65, v46, v3
	v_fmac_f32_e32 v91, v50, v40
	v_mul_f32_e32 v46, v46, v51
	v_fma_f32 v54, v45, v81, -v54
	v_mul_f32_e32 v81, v3, v72
	v_dual_fmac_f32 v62, v1, v4 :: v_dual_mul_f32 v79, v51, v51
	v_fmac_f32_e32 v73, v3, v87
	v_div_fixup_f32 v52, v52, v50, v36
	v_dual_mul_f32 v82, v51, v48 :: v_dual_mul_f32 v85, v2, v47
	s_waitcnt lgkmcnt(0)
	v_fmac_f32_e32 v62, s16, v2
	v_dual_mul_f32 v86, v3, v61 :: v_dual_mul_f32 v87, v45, v79
	v_mul_f32_e32 v81, v4, v81
	v_mul_f32_e64 v83, v1, -v51
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v62, s17, v3
	v_dual_mul_f32 v88, v1, v61 :: v_dual_mul_f32 v89, v51, v47
	v_fmac_f32_e32 v87, v50, v43
	v_fma_f32 v84, v43, v61, -v84
	v_fmac_f32_e32 v62, s18, v4
	v_fmac_f32_e32 v64, v4, v85
	v_fma_f32 v81, v40, v83, -v81
	v_mul_f32_e64 v71, v1, -v48
	v_fmac_f32_e32 v84, v3, v85
	v_fmac_f32_e32 v62, s29, v67
	v_fma_f32 v73, -v3, v76, v73
	v_fma_f32 v57, v4, v75, -v57
	v_fma_f32 v58, v45, v51, -v58
	s_delay_alu instid0(VALU_DEP_4)
	v_fma_f32 v62, -s28, v68, v62
	v_mul_f32_e32 v68, v43, v72
	v_fmac_f32_e32 v73, v63, v4
	v_mul_f32_e32 v63, v40, v89
	v_mul_f32_e32 v89, v2, v89
	v_fma_f32 v50, -s29, v76, v62
	v_fmac_f32_e32 v68, v40, v79
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v50, s19, v78
	v_mul_f32_e32 v78, v2, v79
	v_dual_mul_f32 v92, v2, v69 :: v_dual_mul_f32 v79, v3, v79
	v_dual_mul_f32 v83, v4, v78 :: v_dual_mul_f32 v90, v1, v47
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_mul_f32_e32 v85, v4, v92
	v_mul_f32_e32 v62, v3, v80
	v_mul_f32_e32 v47, v47, v48
	v_fma_f32 v71, v43, v71, -v83
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_4) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v91, -2.0, v85
	v_mul_f32_e32 v85, v2, v82
	v_mul_f32_e32 v62, v4, v62
	v_fma_f32 v64, -v4, v53, v64
	v_mul_f32_e32 v47, v2, v47
	v_dual_mul_f32 v83, v3, v85 :: v_dual_fmac_f32 v50, s28, v39
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_fmac_f32_e32 v87, -2.0, v62
	v_fma_f32 v84, -v3, v53, v84
	v_fmac_f32_e32 v64, v4, v86
	v_fmac_f32_e32 v68, -2.0, v83
	v_fma_f32 v50, -v38, s19, v50
	v_mul_f32_e32 v83, v43, v88
	v_mul_f32_e32 v62, v61, v48
	v_fmac_f32_e32 v84, v65, v4
	v_mul_f32_e32 v61, v61, v51
	v_mul_f32_e32 v51, v40, v51
	v_mul_f32_e32 v50, v52, v50
	v_fma_f32 v83, v45, v46, -v83
	v_mul_f32_e32 v52, v45, v74
	v_fma_f32 v63, v43, v62, -v63
	v_mul_f32_e32 v72, v2, v72
	v_mul_f32_e64 v50, |v50|, s21
	v_fmac_f32_e32 v83, v43, v80
	v_fma_f32 v64, -v4, v75, v64
	v_fmac_f32_e32 v63, v40, v82
	v_fma_f32 v51, v3, v53, -v51
	v_div_scale_f32 v86, null, s37, s37, v50
	v_fma_f32 v83, -v45, v80, v83
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_fma_f32 v63, -v43, v82, v63
	v_mul_f32_e32 v53, v4, v53
	v_rcp_f32_e32 v65, v86
	v_fma_f32 v38, -v38, v4, v84
	v_fma_f32 v79, -v4, v79, v83
	v_mul_f32_e32 v83, v4, v76
	v_fma_f32 v52, v40, v90, -v52
	v_fma_f32 v63, -v3, v72, v63
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_dual_mul_f32 v74, v3, v74 :: v_dual_fmac_f32 v79, v49, v4
	v_fma_f32 v83, v4, v67, -v83
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(TRANS32_DEP_1)
	v_fma_f32 v52, -v40, v69, v52
	v_fma_f32 v75, -v86, v65, 1.0
	v_fma_f32 v67, -v4, v67, v56
	v_fma_f32 v56, v4, v76, -v56
	v_mul_f32_e32 v76, v3, v76
	v_fmac_f32_e32 v52, v45, v69
	v_fmac_f32_e32 v65, v75, v65
	v_div_scale_f32 v75, vcc_lo, v50, s37, v50
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_fma_f32 v76, v43, v48, -v76
	v_dual_fmac_f32 v52, v4, v72 :: v_dual_mul_f32 v69, v3, v69
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_mul_f32_e32 v72, v75, v65
	v_mul_f32_e32 v48, v45, v48
	v_dual_mul_f32 v82, v3, v82 :: v_dual_fmac_f32 v79, v4, v89
	v_fma_f32 v52, -v59, v4, v52
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_fma_f32 v59, -v86, v72, v75
	v_dual_mul_f32 v80, v2, v80 :: v_dual_fmac_f32 v63, v3, v78
	v_fmac_f32_e32 v71, v4, v82
	v_mul_f32_e32 v90, v2, v90
	s_delay_alu instid0(VALU_DEP_4)
	v_fmac_f32_e32 v72, v59, v65
	v_fma_f32 v52, -v4, v47, v52
	v_fmac_f32_e32 v63, v3, v47
	v_fma_f32 v1, v1, v40, -v53
	v_mul_f32_e32 v59, v66, v3
	v_fma_f32 v75, -v86, v72, v75
	v_dual_mul_f32 v66, v66, v2 :: v_dual_fmac_f32 v81, v4, v85
	v_mul_f32_e32 v46, v2, v46
	v_fma_f32 v48, v39, v4, -v48
	s_delay_alu instid0(VALU_DEP_4)
	v_div_fmas_f32 v49, v75, v65, v72
	s_mov_b32 vcc_lo, s0
	v_fmac_f32_e32 v81, v3, v92
	v_div_fmas_f32 v60, v60, v70, v77
	v_fma_f32 v39, -v39, v4, v73
	v_div_fixup_f32 v47, v49, s37, v50
	v_fmac_f32_e32 v54, v4, v69
	v_fma_f32 v50, -v4, v85, v79
	v_fmac_f32_e32 v71, v3, v80
	s_delay_alu instid0(VALU_DEP_4)
	v_mul_f32_e32 v53, v47, v47
	v_cmp_nlt_f32_e32 vcc_lo, 1.0, v47
	v_mul_f32_e32 v62, v3, v62
	v_mul_f32_e32 v88, v2, v88
	v_fma_f32 v50, -v3, v90, v50
	v_mul_f32_e32 v53, v47, v53
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v49, -v4, v62, v52
	v_dual_mul_f32 v86, v3, v61 :: v_dual_fmac_f32 v49, v4, v82
	v_mul_f32_e32 v2, v2, v61
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_fmac_f32 v50, v4, v86 :: v_dual_fmac_f32 v49, v3, v88
	v_fma_f32 v2, -v3, v2, v63
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v50, v3, v92
	v_fma_f32 v49, -v3, v80, v49
	v_add_f32_e32 v52, v47, v47
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_dual_fmac_f32 v2, v4, v74 :: v_dual_fmac_f32 v49, v66, v4
	v_fma_f32 v52, -v47, v52, 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v2, -v4, v69, v2
	v_fmac_f32_e32 v52, v47, v53
	v_div_fixup_f32 v53, v60, v55, v37
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f32 v2, -v4, v46, v2
	v_cndmask_b32_e32 v47, 0, v52, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v54, v4, v80 :: v_dual_mul_f32 v3, v53, v47
	v_dual_fmac_f32 v2, v4, v80 :: v_dual_fmac_f32 v23, v67, v3
	v_fmac_f32_e32 v29, v56, v3
	v_fma_f32 v46, -v59, v4, v50
	v_fmac_f32_e32 v30, v40, v3
	v_fmac_f32_e32 v31, v41, v3
	v_fmac_f32_e32 v32, v42, v3
	v_fmac_f32_e32 v28, v1, v3
	v_fmac_f32_e32 v27, v51, v3
	v_fmac_f32_e32 v26, v43, v3
	v_fmac_f32_e32 v25, v44, v3
	v_fmac_f32_e32 v24, v57, v3
	v_fmac_f32_e32 v22, v76, v3
	v_fmac_f32_e32 v21, v45, v3
	v_fmac_f32_e32 v20, v58, v3
	v_fmac_f32_e32 v19, v48, v3
	v_fmac_f32_e32 v18, v83, v3
	v_fmac_f32_e32 v17, v87, v3
	v_fmac_f32_e32 v16, v54, v3
	v_fmac_f32_e32 v15, v71, v3
	v_fmac_f32_e32 v14, v91, v3
	v_fmac_f32_e32 v13, v81, v3
	v_fmac_f32_e32 v12, v68, v3
	v_fmac_f32_e32 v11, v39, v3
	v_fmac_f32_e32 v10, v38, v3
	v_fmac_f32_e32 v9, v64, v3
	v_fmac_f32_e32 v8, v46, v3
	v_fmac_f32_e32 v7, v49, v3
	v_fmac_f32_e32 v6, v2, v3
.LBB0_4:                                ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s36
	v_add_nc_u32_e32 v5, s31, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_ge_u32_e32 vcc_lo, v5, v33
	s_or_b32 s23, vcc_lo, s23
	s_and_not1_b32 exec_lo, exec_lo, s23
	s_cbranch_execnz .LBB0_2
; %bb.5:                                ; %Flow
	s_or_b32 exec_lo, exec_lo, s23
.LBB0_6:                                ; %Flow788
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s34
	s_mul_i32 s0, s22, s20
	s_mul_i32 s0, s0, 27
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_add3_u32 v1, s0, s30, v0
	v_mul_f32_e64 v0, s21, s21
	s_ashr_i32 s21, s20, 31
	s_lshl_b64 s[0:1], s[20:21], 2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_4)
	v_ashrrev_i32_e32 v2, 31, v1
	v_add3_u32 v3, s20, s20, v1
	v_mul_f32_e32 v5, v0, v30
	v_mul_f32_e32 v12, v0, v12
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v30, s20, v3
	v_ashrrev_i32_e32 v4, 31, v3
	v_dual_mul_f32 v35, v0, v31 :: v_dual_add_nc_u32 v34, s20, v30
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_mul_f32_e32 v36, v0, v32
	v_add_co_u32 v32, vcc_lo, v1, s0
	v_lshlrev_b64 v[3:4], 2, v[3:4]
	v_ashrrev_i32_e32 v31, 31, v30
	v_add_co_ci_u32_e32 v33, vcc_lo, s1, v2, vcc_lo
	s_clause 0x1
	global_store_b32 v[1:2], v5, off
	global_store_b32 v[32:33], v35, off
	v_add_co_u32 v1, vcc_lo, s4, v3
	v_add_nc_u32_e32 v3, s20, v34
	v_lshlrev_b64 v[30:31], 2, v[30:31]
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v4, vcc_lo
	v_ashrrev_i32_e32 v35, 31, v34
	v_mul_f32_e32 v5, v0, v29
	v_ashrrev_i32_e32 v4, 31, v3
	v_add_co_u32 v29, vcc_lo, s4, v30
	global_store_b32 v[1:2], v36, off
	v_add_nc_u32_e32 v1, s20, v3
	v_add_co_ci_u32_e32 v30, vcc_lo, s5, v31, vcc_lo
	v_lshlrev_b64 v[31:32], 2, v[34:35]
	v_lshlrev_b64 v[3:4], 2, v[3:4]
	s_delay_alu instid0(VALU_DEP_4)
	v_ashrrev_i32_e32 v2, 31, v1
	global_store_b32 v[29:30], v5, off
	v_mul_f32_e32 v5, v0, v28
	v_mul_f32_e32 v33, v0, v26
	v_add_co_u32 v28, vcc_lo, s4, v31
	v_add_co_ci_u32_e32 v29, vcc_lo, s5, v32, vcc_lo
	v_mul_f32_e32 v32, v0, v27
	v_lshlrev_b64 v[30:31], 2, v[1:2]
	v_add_nc_u32_e32 v1, s20, v1
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_3)
	v_ashrrev_i32_e32 v2, 31, v1
	v_add_nc_u32_e32 v26, s20, v1
	v_add_co_u32 v30, vcc_lo, s4, v30
	global_store_b32 v[28:29], v5, off
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_ashrrev_i32_e32 v27, 31, v26
	v_mul_f32_e32 v5, v0, v25
	v_add_nc_u32_e32 v25, s20, v26
	v_add_co_ci_u32_e32 v31, vcc_lo, s5, v31, vcc_lo
	v_add_co_u32 v1, vcc_lo, s4, v1
	global_store_b32 v[3:4], v32, off
	v_lshlrev_b64 v[3:4], 2, v[26:27]
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_add_nc_u32_e32 v27, s20, v25
	v_ashrrev_i32_e32 v26, 31, v25
	s_clause 0x1
	global_store_b32 v[30:31], v33, off
	global_store_b32 v[1:2], v5, off
	v_mul_f32_e32 v29, v0, v24
	v_add_nc_u32_e32 v1, s20, v27
	v_lshlrev_b64 v[24:25], 2, v[25:26]
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_ashrrev_i32_e32 v2, 31, v1
	v_mul_f32_e32 v5, v0, v23
	v_add_co_u32 v23, vcc_lo, s4, v24
	v_ashrrev_i32_e32 v28, 31, v27
	v_add_co_ci_u32_e32 v24, vcc_lo, s5, v25, vcc_lo
	v_lshlrev_b64 v[25:26], 2, v[1:2]
	v_add_nc_u32_e32 v1, s20, v1
	global_store_b32 v[3:4], v29, off
	v_lshlrev_b64 v[3:4], 2, v[27:28]
	v_mul_f32_e32 v28, v0, v21
	v_mul_f32_e32 v27, v0, v22
	v_ashrrev_i32_e32 v2, 31, v1
	v_add_nc_u32_e32 v21, s20, v1
	global_store_b32 v[23:24], v5, off
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	v_add_co_u32 v25, vcc_lo, s4, v25
	v_ashrrev_i32_e32 v22, 31, v21
	v_dual_mul_f32 v5, v0, v20 :: v_dual_add_nc_u32 v20, s20, v21
	v_add_co_ci_u32_e32 v26, vcc_lo, s5, v26, vcc_lo
	v_add_co_u32 v1, vcc_lo, s4, v1
	global_store_b32 v[3:4], v27, off
	v_lshlrev_b64 v[3:4], 2, v[21:22]
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_add_nc_u32_e32 v22, s20, v20
	v_ashrrev_i32_e32 v21, 31, v20
	s_clause 0x1
	global_store_b32 v[25:26], v28, off
	global_store_b32 v[1:2], v5, off
	v_dual_mul_f32 v24, v0, v19 :: v_dual_add_nc_u32 v1, s20, v22
	v_lshlrev_b64 v[19:20], 2, v[20:21]
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_ashrrev_i32_e32 v2, 31, v1
	v_mul_f32_e32 v5, v0, v18
	v_add_co_u32 v18, vcc_lo, s4, v19
	v_ashrrev_i32_e32 v23, 31, v22
	v_add_co_ci_u32_e32 v19, vcc_lo, s5, v20, vcc_lo
	v_lshlrev_b64 v[20:21], 2, v[1:2]
	v_add_nc_u32_e32 v1, s20, v1
	global_store_b32 v[3:4], v24, off
	v_lshlrev_b64 v[3:4], 2, v[22:23]
	v_mul_f32_e32 v23, v0, v16
	v_mul_f32_e32 v22, v0, v17
	v_ashrrev_i32_e32 v2, 31, v1
	v_add_nc_u32_e32 v16, s20, v1
	global_store_b32 v[18:19], v5, off
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	v_add_co_u32 v20, vcc_lo, s4, v20
	v_ashrrev_i32_e32 v17, 31, v16
	v_mul_f32_e32 v5, v0, v15
	v_add_nc_u32_e32 v15, s20, v16
	v_add_co_ci_u32_e32 v21, vcc_lo, s5, v21, vcc_lo
	v_add_co_u32 v1, vcc_lo, s4, v1
	global_store_b32 v[3:4], v22, off
	v_lshlrev_b64 v[3:4], 2, v[16:17]
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_add_nc_u32_e32 v17, s20, v15
	v_ashrrev_i32_e32 v16, 31, v15
	s_clause 0x1
	global_store_b32 v[20:21], v23, off
	global_store_b32 v[1:2], v5, off
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_add_nc_u32_e32 v1, s20, v17
	v_mul_f32_e32 v19, v0, v14
	v_lshlrev_b64 v[14:15], 2, v[15:16]
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	v_ashrrev_i32_e32 v18, 31, v17
	v_ashrrev_i32_e32 v2, 31, v1
	v_mul_f32_e32 v5, v0, v13
	global_store_b32 v[3:4], v19, off
	v_add_co_u32 v13, vcc_lo, s4, v14
	v_lshlrev_b64 v[3:4], 2, v[17:18]
	v_add_co_ci_u32_e32 v14, vcc_lo, s5, v15, vcc_lo
	v_lshlrev_b64 v[15:16], 2, v[1:2]
	v_add_nc_u32_e32 v1, s20, v1
	v_mul_f32_e32 v17, v0, v11
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_add_nc_u32_e32 v11, s20, v1
	v_add_co_u32 v15, vcc_lo, s4, v15
	v_add_co_ci_u32_e32 v16, vcc_lo, s5, v16, vcc_lo
	s_clause 0x1
	global_store_b32 v[13:14], v5, off
	global_store_b32 v[3:4], v12, off
	v_add_nc_u32_e32 v3, s20, v11
	v_ashrrev_i32_e32 v2, 31, v1
	global_store_b32 v[15:16], v17, off
	v_mul_f32_e32 v16, v0, v10
	v_ashrrev_i32_e32 v12, 31, v11
	v_add_nc_u32_e32 v10, s20, v3
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	v_ashrrev_i32_e32 v4, 31, v3
	v_mul_f32_e32 v17, v0, v9
	v_lshlrev_b64 v[12:13], 2, v[11:12]
	v_ashrrev_i32_e32 v11, 31, v10
	v_add_nc_u32_e32 v9, s20, v10
	v_add_co_u32 v1, vcc_lo, s4, v1
	v_lshlrev_b64 v[3:4], 2, v[3:4]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[14:15], 2, v[10:11]
	v_ashrrev_i32_e32 v10, 31, v9
	v_add_co_ci_u32_e32 v2, vcc_lo, s5, v2, vcc_lo
	v_add_co_u32 v12, vcc_lo, s4, v12
	v_add_co_ci_u32_e32 v13, vcc_lo, s5, v13, vcc_lo
	v_mul_f32_e32 v18, v0, v8
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_mul_f32_e32 v11, v0, v7
	v_mul_f32_e32 v0, v0, v6
	v_lshlrev_b64 v[7:8], 2, v[9:10]
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	v_add_co_u32 v9, vcc_lo, s4, v14
	v_add_co_ci_u32_e32 v10, vcc_lo, s5, v15, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s4, v7
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v8, vcc_lo
	s_clause 0x4
	global_store_b32 v[1:2], v16, off
	global_store_b32 v[12:13], v17, off
	global_store_b32 v[3:4], v18, off
	global_store_b32 v[9:10], v11, off
	global_store_b32 v[5:6], v0, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 352
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 93
		.amdhsa_next_free_sgpr 38
		.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	_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_, .Lfunc_end0-_Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 3756
; NumSgprs: 40
; NumVgprs: 93
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 4
; VGPRBlocks: 11
; NumSGPRsForWavesPerEU: 40
; NumVGPRsForWavesPerEU: 93
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 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:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         44
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     by_value
      - .offset:         52
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .offset:         80
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         88
        .size:           8
        .value_kind:     global_buffer
      - .offset:         96
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         100
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         104
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         108
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         110
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         112
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         114
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         116
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         118
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         136
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         144
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         152
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         160
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 352
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_
    .private_segment_fixed_size: 0
    .sgpr_count:     40
    .sgpr_spill_count: 0
    .symbol:         _Z33normal_eqs_disparity_weighted_GPUPfPKfPK15HIP_vector_typeIfLj4EEPKifffffiS7_S7_S1_fS1_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     93
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
