#include "includes.h"
__global__ void cudaSNormalizeROIs_kernel( unsigned int inputSizeX, unsigned int inputSizeY, unsigned int nbProposals, unsigned int batchSize, unsigned int scoreIdx, unsigned int nbCls, unsigned int maxParts, unsigned int maxTemplates, bool keepMax, bool generateParts, bool generateTemplates, const float normX, const float normY, const float* means, const float* std, const unsigned int* numPartsPerClass, const unsigned int* numTemplatesPerClass, const float* ROIRef, const float* ROIEst, const float* ValuesEst, const float* partsEst, const float* partsVisibilityEst, const float* templatesEst, float* outputs, int* argMax, float* partsPrediction, float* partsVisibilityPrediction, float* templatesPrediction, float scoreThreshold)
{
const int batchPos = blockIdx.z*nbProposals;
const int index = (threadIdx.x & 0x1f) + blockIdx.x*blockDim.x;

if(index < nbProposals)
{

unsigned int indexMin = scoreIdx;
unsigned int indexMax = nbCls;

if(keepMax)
{
unsigned int cls = scoreIdx;
float maxVal = 0.0;

for(unsigned int i = indexMin; i < indexMax; ++i)
{
unsigned int inputIdx = i + index*nbCls + batchPos*nbCls;

if (ValuesEst[inputIdx] >= maxVal)
{
maxVal = ValuesEst[inputIdx];
cls = i;
}

}
argMax[index + batchPos] = cls;

indexMin = cls;
indexMax = cls + 1;
}
else
argMax[index + batchPos] = -1;

for(unsigned int clsIdx = indexMin; clsIdx < indexMax; ++clsIdx)
{

unsigned int bboxRefIdx = index*4 + batchPos*4;
unsigned int bboxEstIdx = clsIdx*4 + index*4*nbCls + batchPos*4*nbCls;
unsigned int valEstIdx = clsIdx + index*nbCls + batchPos*nbCls;

unsigned int outputIdx = keepMax ? index*4*(nbCls - scoreIdx) + batchPos*4*(nbCls - scoreIdx)
: (clsIdx - scoreIdx)*4 + index*4*(nbCls - scoreIdx) + batchPos*4*(nbCls - scoreIdx);


const float xbbRef = ROIRef[0 + bboxRefIdx]*normX;
const float ybbRef = ROIRef[1 + bboxRefIdx]*normY;
const float wbbRef = ROIRef[2 + bboxRefIdx]*normX;
const float hbbRef = ROIRef[3 + bboxRefIdx]*normY;


const float xbbEst = ROIEst[0 + bboxEstIdx]*std[0] + means[0];

const float ybbEst = ROIEst[1 + bboxEstIdx]*std[1] + means[1];

const float wbbEst = ROIEst[2 + bboxEstIdx]*std[2] + means[2];

const float hbbEst = ROIEst[3 + bboxEstIdx]*std[3] + means[3];


float x = xbbEst*wbbRef + xbbRef + wbbRef/2.0 - (wbbRef/2.0)*exp(wbbEst);
float y = ybbEst*hbbRef + ybbRef + hbbRef/2.0 - (hbbRef/2.0)*exp(hbbEst);
float w = wbbRef*exp(wbbEst);
float h = hbbRef*exp(hbbEst);

/**Clip values**/
if(x < 0.0)
{
w += x;
x = 0.0;
}

if(y < 0.0)
{
h += y;
y = 0.0;
}

w = ((w + x) > 1.0) ? (1.0 - x) / normX : w / normX;
h = ((h + y) > 1.0) ? (1.0 - y) / normY : h / normY;

x /= normX;
y /= normY;

if(ValuesEst[valEstIdx] >= scoreThreshold)
{
outputs[0 + outputIdx] = x;
outputs[1 + outputIdx] = y;
outputs[2 + outputIdx] = w;
outputs[3 + outputIdx] = h;

if(generateParts)
{
unsigned int partsIdx = 0;
unsigned int proposalPartIdx = 0;
for(int idxCls = 0; idxCls < clsIdx; ++idxCls)
partsIdx += numPartsPerClass[idxCls];

proposalPartIdx = partsIdx;

for(int idxCls = clsIdx; idxCls < nbCls; ++idxCls)
proposalPartIdx += numPartsPerClass[idxCls];

for(unsigned int part = 0; part < numPartsPerClass[clsIdx];
++part)
{
//const unsigned int partIdx = (partsIdx + part)*2;
const unsigned int inPartIdx = batchPos*2*proposalPartIdx
+ index*2*proposalPartIdx
+ partsIdx*2
+ part*2;

const unsigned int outPartIdx = batchPos*2*maxParts*nbCls
+ index*2*maxParts*nbCls
+ clsIdx*2*maxParts
+ part*2;

const float partY = partsEst[0 + inPartIdx];
const float partX = partsEst[1 + inPartIdx];
partsPrediction[0 + outPartIdx] = ((partY + 0.5) * hbbRef + ybbRef) / normY;
partsPrediction[1 + outPartIdx] = ((partX + 0.5) * wbbRef + xbbRef) / normX;

/// PARTS VISIBILITY PROCESSING
const unsigned int inPartVisibilityIdx = batchPos*4*proposalPartIdx
+ index*4*proposalPartIdx
+ partsIdx
+ part;

const unsigned int outPartVisibilityIdx = batchPos*maxParts*nbCls
+ index*maxParts*nbCls
+ clsIdx*maxParts
+ part;

float idxMax = 0.0;
float valueMax = partsVisibilityEst[inPartVisibilityIdx];

for(unsigned int v = 1; v < 4; ++v)
{
if(partsVisibilityEst[v*proposalPartIdx + inPartVisibilityIdx] >
valueMax)
{
idxMax = (float) v;
valueMax = partsVisibilityEst[v*proposalPartIdx + inPartVisibilityIdx];
}
}
partsVisibilityPrediction[outPartVisibilityIdx] = idxMax;
}
}


if(generateTemplates)
{
unsigned int templatesIdx = 0;
unsigned int proposalTemplateIdx = 0;
for(int idxCls = 0; idxCls < clsIdx; ++idxCls)
templatesIdx += numTemplatesPerClass[idxCls];

proposalTemplateIdx = templatesIdx;

for(int idxCls = clsIdx; idxCls < nbCls; ++idxCls)
proposalTemplateIdx += numTemplatesPerClass[idxCls];

for(unsigned int tpl = 0; tpl < numTemplatesPerClass[clsIdx];
++tpl)
{
const unsigned int inTemplateIdx = batchPos*3*proposalTemplateIdx
+ index*3*proposalTemplateIdx
+ templatesIdx*3
+ tpl*3;

const unsigned int outTemplateIdx = batchPos*3*maxTemplates*nbCls
+ index*3*maxTemplates*nbCls
+ clsIdx*3*maxTemplates
+ tpl*3;

const float templateA = expf(templatesEst[0 + inTemplateIdx]);
const float templateB = expf(templatesEst[1 + inTemplateIdx]);
const float templateC = expf(templatesEst[2 + inTemplateIdx]);

templatesPrediction[0 + outTemplateIdx] = templateA;
templatesPrediction[1 + outTemplateIdx] = templateB;
templatesPrediction[2 + outTemplateIdx] = templateC;
}
}
}
else
{
outputs[0 + outputIdx] = 0.0;
outputs[1 + outputIdx] = 0.0;
outputs[2 + outputIdx] = 0.0;
outputs[3 + outputIdx] = 0.0;
}
}
}

}
	.file	"tmpxft_00259f18_00000000-6_cudaSNormalizeROIs_kernel.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z100__device_stub__Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_fjjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.type	_Z100__device_stub__Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_fjjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f, @function
_Z100__device_stub__Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_fjjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f:
.LFB2032:
	.cfi_startproc
	endbr64
	subq	$488, %rsp
	.cfi_def_cfa_offset 496
	movl	512(%rsp), %eax
	movl	%edi, 172(%rsp)
	leaq	192(%rsp), %rdi
	movl	%esi, 168(%rsp)
	leaq	204(%rsp), %rsi
	movb	%al, 148(%rsp)
	movl	520(%rsp), %eax
	movl	%edx, 164(%rsp)
	leaq	176(%rsp), %rdx
	movb	%al, 144(%rsp)
	movl	528(%rsp), %eax
	movl	%ecx, 160(%rsp)
	leaq	184(%rsp), %rcx
	movb	%al, 140(%rsp)
	movq	536(%rsp), %rax
	movl	%r8d, 156(%rsp)
	movq	%rax, 120(%rsp)
	movq	544(%rsp), %rax
	movl	%r9d, 152(%rsp)
	movq	%rax, 112(%rsp)
	movq	552(%rsp), %rax
	movss	%xmm0, 136(%rsp)
	movq	%rax, 104(%rsp)
	movq	560(%rsp), %rax
	movss	%xmm1, 132(%rsp)
	movq	%rax, 96(%rsp)
	movq	568(%rsp), %rax
	movss	%xmm2, 128(%rsp)
	movq	%rax, 88(%rsp)
	movq	576(%rsp), %rax
	movq	%rax, 80(%rsp)
	movq	584(%rsp), %rax
	movq	%rax, 72(%rsp)
	movq	592(%rsp), %rax
	movq	%rax, 64(%rsp)
	movq	600(%rsp), %rax
	movq	%rax, 56(%rsp)
	movq	608(%rsp), %rax
	movq	%rax, 48(%rsp)
	movq	616(%rsp), %rax
	movq	%rax, 40(%rsp)
	movq	624(%rsp), %rax
	movq	%rax, 32(%rsp)
	movq	632(%rsp), %rax
	movq	%rax, 24(%rsp)
	movq	640(%rsp), %rax
	movq	%rax, 16(%rsp)
	movq	648(%rsp), %rax
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 472(%rsp)
	xorl	%eax, %eax
	leaq	172(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	168(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	164(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	160(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	156(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	152(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	496(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	504(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	148(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	144(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	140(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	136(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	132(%rsp), %rax
	movq	%rax, 336(%rsp)
	leaq	120(%rsp), %rax
	movq	%rax, 344(%rsp)
	leaq	112(%rsp), %rax
	movq	%rax, 352(%rsp)
	leaq	104(%rsp), %rax
	movq	%rax, 360(%rsp)
	leaq	96(%rsp), %rax
	movq	%rax, 368(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 376(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 384(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 392(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 400(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 408(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 416(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 424(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 432(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 440(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 448(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 456(%rsp)
	leaq	128(%rsp), %rax
	movq	%rax, 464(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 192(%rsp)
	movl	$1, 200(%rsp)
	movq	%rax, 204(%rsp)
	movl	$1, 212(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	184(%rsp)
	.cfi_def_cfa_offset 504
	leaq	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f(%rip), %rdi
	pushq	184(%rsp)
	.cfi_def_cfa_offset 512
	movq	220(%rsp), %rcx
	movl	228(%rsp), %r8d
	movq	208(%rsp), %rsi
	movl	216(%rsp), %edx
	leaq	256(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 504
	popq	%rdx
	.cfi_def_cfa_offset 496
.L2:
	movq	472(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$488, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2032:
	.size	_Z100__device_stub__Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_fjjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f, .-_Z100__device_stub__Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_fjjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.globl	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.type	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f, @function
_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f:
.LFB2033:
	.cfi_startproc
	endbr64
	movzbl	40(%rsp), %eax
	movl	%eax, 40(%rsp)
	movzbl	32(%rsp), %eax
	movl	%eax, 32(%rsp)
	movzbl	24(%rsp), %eax
	movl	%eax, 24(%rsp)
	jmp	_Z100__device_stub__Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_fjjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.cfi_endproc
.LFE2033:
	.size	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f, .-_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2035:
	.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	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f(%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
.LFE2035:
	.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"
__global__ void cudaSNormalizeROIs_kernel( unsigned int inputSizeX, unsigned int inputSizeY, unsigned int nbProposals, unsigned int batchSize, unsigned int scoreIdx, unsigned int nbCls, unsigned int maxParts, unsigned int maxTemplates, bool keepMax, bool generateParts, bool generateTemplates, const float normX, const float normY, const float* means, const float* std, const unsigned int* numPartsPerClass, const unsigned int* numTemplatesPerClass, const float* ROIRef, const float* ROIEst, const float* ValuesEst, const float* partsEst, const float* partsVisibilityEst, const float* templatesEst, float* outputs, int* argMax, float* partsPrediction, float* partsVisibilityPrediction, float* templatesPrediction, float scoreThreshold)
{
const int batchPos = blockIdx.z*nbProposals;
const int index = (threadIdx.x & 0x1f) + blockIdx.x*blockDim.x;

if(index < nbProposals)
{

unsigned int indexMin = scoreIdx;
unsigned int indexMax = nbCls;

if(keepMax)
{
unsigned int cls = scoreIdx;
float maxVal = 0.0;

for(unsigned int i = indexMin; i < indexMax; ++i)
{
unsigned int inputIdx = i + index*nbCls + batchPos*nbCls;

if (ValuesEst[inputIdx] >= maxVal)
{
maxVal = ValuesEst[inputIdx];
cls = i;
}

}
argMax[index + batchPos] = cls;

indexMin = cls;
indexMax = cls + 1;
}
else
argMax[index + batchPos] = -1;

for(unsigned int clsIdx = indexMin; clsIdx < indexMax; ++clsIdx)
{

unsigned int bboxRefIdx = index*4 + batchPos*4;
unsigned int bboxEstIdx = clsIdx*4 + index*4*nbCls + batchPos*4*nbCls;
unsigned int valEstIdx = clsIdx + index*nbCls + batchPos*nbCls;

unsigned int outputIdx = keepMax ? index*4*(nbCls - scoreIdx) + batchPos*4*(nbCls - scoreIdx)
: (clsIdx - scoreIdx)*4 + index*4*(nbCls - scoreIdx) + batchPos*4*(nbCls - scoreIdx);


const float xbbRef = ROIRef[0 + bboxRefIdx]*normX;
const float ybbRef = ROIRef[1 + bboxRefIdx]*normY;
const float wbbRef = ROIRef[2 + bboxRefIdx]*normX;
const float hbbRef = ROIRef[3 + bboxRefIdx]*normY;


const float xbbEst = ROIEst[0 + bboxEstIdx]*std[0] + means[0];

const float ybbEst = ROIEst[1 + bboxEstIdx]*std[1] + means[1];

const float wbbEst = ROIEst[2 + bboxEstIdx]*std[2] + means[2];

const float hbbEst = ROIEst[3 + bboxEstIdx]*std[3] + means[3];


float x = xbbEst*wbbRef + xbbRef + wbbRef/2.0 - (wbbRef/2.0)*exp(wbbEst);
float y = ybbEst*hbbRef + ybbRef + hbbRef/2.0 - (hbbRef/2.0)*exp(hbbEst);
float w = wbbRef*exp(wbbEst);
float h = hbbRef*exp(hbbEst);

/**Clip values**/
if(x < 0.0)
{
w += x;
x = 0.0;
}

if(y < 0.0)
{
h += y;
y = 0.0;
}

w = ((w + x) > 1.0) ? (1.0 - x) / normX : w / normX;
h = ((h + y) > 1.0) ? (1.0 - y) / normY : h / normY;

x /= normX;
y /= normY;

if(ValuesEst[valEstIdx] >= scoreThreshold)
{
outputs[0 + outputIdx] = x;
outputs[1 + outputIdx] = y;
outputs[2 + outputIdx] = w;
outputs[3 + outputIdx] = h;

if(generateParts)
{
unsigned int partsIdx = 0;
unsigned int proposalPartIdx = 0;
for(int idxCls = 0; idxCls < clsIdx; ++idxCls)
partsIdx += numPartsPerClass[idxCls];

proposalPartIdx = partsIdx;

for(int idxCls = clsIdx; idxCls < nbCls; ++idxCls)
proposalPartIdx += numPartsPerClass[idxCls];

for(unsigned int part = 0; part < numPartsPerClass[clsIdx];
++part)
{
//const unsigned int partIdx = (partsIdx + part)*2;
const unsigned int inPartIdx = batchPos*2*proposalPartIdx
+ index*2*proposalPartIdx
+ partsIdx*2
+ part*2;

const unsigned int outPartIdx = batchPos*2*maxParts*nbCls
+ index*2*maxParts*nbCls
+ clsIdx*2*maxParts
+ part*2;

const float partY = partsEst[0 + inPartIdx];
const float partX = partsEst[1 + inPartIdx];
partsPrediction[0 + outPartIdx] = ((partY + 0.5) * hbbRef + ybbRef) / normY;
partsPrediction[1 + outPartIdx] = ((partX + 0.5) * wbbRef + xbbRef) / normX;

/// PARTS VISIBILITY PROCESSING
const unsigned int inPartVisibilityIdx = batchPos*4*proposalPartIdx
+ index*4*proposalPartIdx
+ partsIdx
+ part;

const unsigned int outPartVisibilityIdx = batchPos*maxParts*nbCls
+ index*maxParts*nbCls
+ clsIdx*maxParts
+ part;

float idxMax = 0.0;
float valueMax = partsVisibilityEst[inPartVisibilityIdx];

for(unsigned int v = 1; v < 4; ++v)
{
if(partsVisibilityEst[v*proposalPartIdx + inPartVisibilityIdx] >
valueMax)
{
idxMax = (float) v;
valueMax = partsVisibilityEst[v*proposalPartIdx + inPartVisibilityIdx];
}
}
partsVisibilityPrediction[outPartVisibilityIdx] = idxMax;
}
}


if(generateTemplates)
{
unsigned int templatesIdx = 0;
unsigned int proposalTemplateIdx = 0;
for(int idxCls = 0; idxCls < clsIdx; ++idxCls)
templatesIdx += numTemplatesPerClass[idxCls];

proposalTemplateIdx = templatesIdx;

for(int idxCls = clsIdx; idxCls < nbCls; ++idxCls)
proposalTemplateIdx += numTemplatesPerClass[idxCls];

for(unsigned int tpl = 0; tpl < numTemplatesPerClass[clsIdx];
++tpl)
{
const unsigned int inTemplateIdx = batchPos*3*proposalTemplateIdx
+ index*3*proposalTemplateIdx
+ templatesIdx*3
+ tpl*3;

const unsigned int outTemplateIdx = batchPos*3*maxTemplates*nbCls
+ index*3*maxTemplates*nbCls
+ clsIdx*3*maxTemplates
+ tpl*3;

const float templateA = expf(templatesEst[0 + inTemplateIdx]);
const float templateB = expf(templatesEst[1 + inTemplateIdx]);
const float templateC = expf(templatesEst[2 + inTemplateIdx]);

templatesPrediction[0 + outTemplateIdx] = templateA;
templatesPrediction[1 + outTemplateIdx] = templateB;
templatesPrediction[2 + outTemplateIdx] = templateC;
}
}
}
else
{
outputs[0 + outputIdx] = 0.0;
outputs[1 + outputIdx] = 0.0;
outputs[2 + outputIdx] = 0.0;
outputs[3 + outputIdx] = 0.0;
}
}
}

}
	.text
	.file	"cudaSNormalizeROIs_kernel.hip"
	.globl	_Z40__device_stub__cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f # -- Begin function _Z40__device_stub__cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.type	_Z40__device_stub__cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f,@function
_Z40__device_stub__cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f: # @_Z40__device_stub__cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.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	$336, %rsp                      # imm = 0x150
	.cfi_def_cfa_offset 384
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movb	416(%rsp), %al
	movb	408(%rsp), %r10b
	movb	400(%rsp), %r11b
	leaq	44(%rsp), %r14
	movl	%edi, (%r14)
	leaq	40(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	36(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	32(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	28(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	24(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	11(%rsp), %r9
	movb	%r11b, (%r9)
	leaq	10(%rsp), %r11
	movb	%r10b, (%r11)
	leaq	9(%rsp), %r10
	movb	%al, (%r10)
	leaq	20(%rsp), %r15
	movss	%xmm0, (%r15)
	leaq	16(%rsp), %r12
	movss	%xmm1, (%r12)
	leaq	12(%rsp), %rax
	movss	%xmm2, (%rax)
	leaq	96(%rsp), %rbx
	movq	%r14, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	384(%rsp), %rcx
	movq	%rcx, 48(%rbx)
	leaq	392(%rsp), %rcx
	movq	%rcx, 56(%rbx)
	movq	%r9, 64(%rbx)
	movq	%r11, 72(%rbx)
	movq	%r10, 80(%rbx)
	movq	%r15, 88(%rbx)
	movq	%r12, 96(%rbx)
	leaq	424(%rsp), %rcx
	movq	%rcx, 104(%rbx)
	leaq	432(%rsp), %rcx
	movq	%rcx, 112(%rbx)
	leaq	440(%rsp), %rcx
	movq	%rcx, 120(%rbx)
	leaq	448(%rsp), %rcx
	movq	%rcx, 128(%rbx)
	leaq	456(%rsp), %rcx
	movq	%rcx, 136(%rbx)
	leaq	464(%rsp), %rcx
	movq	%rcx, 144(%rbx)
	leaq	472(%rsp), %rcx
	movq	%rcx, 152(%rbx)
	leaq	480(%rsp), %rcx
	movq	%rcx, 160(%rbx)
	leaq	488(%rsp), %rcx
	movq	%rcx, 168(%rbx)
	leaq	496(%rsp), %rcx
	movq	%rcx, 176(%rbx)
	leaq	504(%rsp), %rcx
	movq	%rcx, 184(%rbx)
	leaq	512(%rsp), %rcx
	movq	%rcx, 192(%rbx)
	leaq	520(%rsp), %rcx
	movq	%rcx, 200(%rbx)
	leaq	528(%rsp), %rcx
	movq	%rcx, 208(%rbx)
	leaq	536(%rsp), %rcx
	movq	%rcx, 216(%rbx)
	movq	%rax, 224(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	56(%rsp), %r12
	leaq	48(%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	$_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$352, %rsp                      # imm = 0x160
	.cfi_adjust_cfa_offset -352
	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	_Z40__device_stub__cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f, .Lfunc_end0-_Z40__device_stub__cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.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	$_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f, %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	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f,@object # @_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.section	.rodata,"a",@progbits
	.globl	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.p2align	3, 0x0
_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f:
	.quad	_Z40__device_stub__cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.size	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f, 8

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

	.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 _Z40__device_stub__cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f ; -- Begin function _Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.globl	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
	.p2align	8
	.type	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f,@function
_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f: ; @_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0xbc
	s_load_b32 s3, s[0:1], 0x8
	v_and_b32_e32 v2, 31, v0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[0:1], null, s14, s2, v[2:3]
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_u32_e64 s3, v0
	s_cbranch_execz .LBB0_51
; %bb.1:
	s_clause 0x4
	s_load_b32 s2, s[0:1], 0x20
	s_load_b128 s[44:47], s[0:1], 0x10
	s_load_b512 s[16:31], s[0:1], 0x30
	s_load_b256 s[4:11], s[0:1], 0x90
	s_load_b256 s[36:43], s[0:1], 0x70
	s_mul_i32 s15, s15, s3
	s_waitcnt lgkmcnt(0)
	s_mov_b32 s10, 0
	v_add_nc_u32_e32 v0, s15, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_ashrrev_i32_e32 v1, 31, v0
	s_bitcmp1_b32 s2, 0
	s_cselect_b32 s3, -1, 0
	s_bitcmp0_b32 s2, 0
	s_cbranch_scc0 .LBB0_3
; %bb.2:
	v_lshlrev_b64 v[2:3], 2, v[0:1]
	v_mov_b32_e32 v4, -1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v2, vcc_lo, s42, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s43, v3, vcc_lo
	global_store_b32 v[2:3], v4, off
	s_branch .LBB0_4
.LBB0_3:
	s_mov_b32 s10, -1
.LBB0_4:                                ; %Flow511
	v_dual_mov_b32 v18, s44 :: v_dual_mov_b32 v31, s45
	s_and_not1_b32 vcc_lo, exec_lo, s10
	s_cbranch_vccnz .LBB0_9
; %bb.5:                                ; %.preheader345
	v_mov_b32_e32 v18, s44
	s_cmp_le_u32 s45, s44
	s_cbranch_scc1 .LBB0_8
; %bb.6:                                ; %.lr.ph
	v_mul_lo_u32 v4, v0, s45
	v_dual_mov_b32 v18, s44 :: v_dual_mov_b32 v3, 0
	v_mov_b32_e32 v5, 0
	s_mov_b32 s10, s44
.LBB0_7:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_3) | instid1(SALU_CYCLE_1)
	v_add_nc_u32_e32 v2, s10, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[6:7], 2, v[2:3]
	v_add_co_u32 v6, vcc_lo, s28, v6
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v7, vcc_lo, s29, v7, vcc_lo
	global_load_b32 v2, v[6:7], off
	s_waitcnt vmcnt(0)
	v_cmp_nge_f32_e32 vcc_lo, v2, v5
	v_cndmask_b32_e32 v5, v2, v5, vcc_lo
	v_cndmask_b32_e32 v18, s10, v18, vcc_lo
	s_add_i32 s10, s10, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_ge_u32 s10, s45
	s_cbranch_scc0 .LBB0_7
.LBB0_8:                                ; %Flow510
	v_lshlrev_b64 v[1:2], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v31, 1, v18
	v_add_co_u32 v1, vcc_lo, s42, v1
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v2, vcc_lo, s43, v2, vcc_lo
	global_store_b32 v[1:2], v18, off
.LBB0_9:
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_lt_u32_e32 vcc_lo, v18, v31
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB0_51
; %bb.10:                               ; %.lr.ph375
	s_clause 0x1
	s_load_b128 s[12:15], s[0:1], 0x20
	s_load_b32 s10, s[0:1], 0xa8
	v_dual_mov_b32 v1, 0 :: v_dual_lshlrev_b32 v24, 2, v0
	v_lshl_add_u32 v32, v0, 1, v0
	v_lshl_add_u32 v2, v18, 1, v18
	s_bitcmp1_b32 s2, 8
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_mov_b32 v25, v1 :: v_dual_lshlrev_b32 v34, 1, v0
	s_cselect_b32 s11, -1, 0
	v_mad_u64_u32 v[3:4], null, v32, s45, v[2:3]
	s_delay_alu instid0(VALU_DEP_2)
	v_lshlrev_b64 v[4:5], 2, v[24:25]
	v_mul_lo_u32 v33, v24, s45
	v_mul_lo_u32 v35, v0, s45
	v_mul_lo_u32 v37, v34, s45
	v_or_b32_e32 v39, 1, v24
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s15, s47, 3
	v_mul_lo_u32 v38, s47, v3
	v_add_co_u32 v25, vcc_lo, s24, v4
	v_cvt_f64_f32_e32 v[20:21], s13
	v_cvt_f64_f32_e32 v[22:23], s14
	s_bitcmp1_b32 s12, 16
	v_add_co_ci_u32_e32 v26, vcc_lo, s25, v5, vcc_lo
	s_cselect_b32 s12, -1, 0
	s_sub_i32 s0, s45, s44
	s_mov_b32 s24, 0
	v_mul_lo_u32 v36, v24, s0
.LBB0_11:                               ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_23 Depth 2
                                        ;     Child Loop BB0_27 Depth 2
                                        ;     Child Loop BB0_31 Depth 2
                                        ;       Child Loop BB0_32 Depth 3
                                        ;     Child Loop BB0_38 Depth 2
                                        ;     Child Loop BB0_42 Depth 2
                                        ;     Child Loop BB0_46 Depth 2
	v_lshl_add_u32 v0, v18, 2, v33
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[0:1]
	v_add_co_u32 v14, vcc_lo, s26, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v15, vcc_lo, s27, v3, vcc_lo
	global_load_b128 v[6:9], v[25:26], off
	s_clause 0x1
	global_load_b128 v[10:13], v1, s[18:19]
	global_load_b128 v[2:5], v1, s[16:17]
	global_load_b128 v[14:17], v[14:15], off
	s_waitcnt vmcnt(3)
	v_mul_f32_e32 v40, s13, v6
	s_waitcnt vmcnt(0)
	v_fma_f32 v0, v16, v12, v4
	v_mul_f32_e32 v4, s13, v8
	v_fma_f32 v2, v14, v10, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_f32_e32 v6, 0x3fb8aa3b, v0
	v_cvt_f64_f32_e32 v[27:28], v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_fma_f32 v2, v4, v2, v40
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v0
	v_fma_f32 v8, 0x3fb8aa3b, v0, -v6
	v_rndne_f32_e32 v10, v6
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_f64_f32_e32 v[29:30], v2
	v_fmac_f32_e32 v8, 0x32a5705f, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v2, v6, v10
	v_cvt_i32_f32_e32 v6, v10
	v_add_f32_e32 v2, v2, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_exp_f32_e32 v2, v2
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v2, v2, v6
	v_cndmask_b32_e32 v2, 0, v2, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v0
	v_mul_f64 v[41:42], v[27:28], 0.5
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v2, 0x7f800000, v2, vcc_lo
	v_fma_f64 v[29:30], v[27:28], 0.5, v[29:30]
	v_cvt_f64_f32_e32 v[43:44], v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[29:30], -v[41:42], v[43:44], v[29:30]
	v_cvt_f32_f64_e32 v6, v[29:30]
                                        ; implicit-def: $vgpr29_vgpr30
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_cmp_gt_f32_e32 vcc_lo, 0, v6
	v_cndmask_b32_e32 v0, 0x80000000, v6, vcc_lo
	v_cndmask_b32_e64 v10, v6, 0, vcc_lo
	v_fmac_f32_e32 v0, v4, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v2, v10, v0
	v_cmpx_nlt_f32_e32 1.0, v2
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB0_13
; %bb.12:                               ;   in Loop: Header=BB0_11 Depth=1
	v_div_scale_f32 v2, null, s13, s13, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v4, v2
	s_waitcnt_depctr 0xfff
	v_fma_f32 v6, -v2, v4, 1.0
	v_fmac_f32_e32 v4, v6, v4
	v_div_scale_f32 v6, vcc_lo, v0, s13, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v8, v6, v4
	v_fma_f32 v12, -v2, v8, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v8, v12, v4
	v_fma_f32 v2, -v2, v8, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v2, v2, v4, v8
	v_div_fixup_f32 v0, v2, s13, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_cvt_f64_f32_e32 v[29:30], v0
.LBB0_13:                               ; %Flow506
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_and_not1_saveexec_b32 s0, s0
	s_cbranch_execz .LBB0_15
; %bb.14:                               ;   in Loop: Header=BB0_11 Depth=1
	v_cvt_f64_f32_e32 v[29:30], v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[29:30], -v[29:30], 1.0
	v_div_scale_f64 v[41:42], null, v[20:21], v[20:21], v[29:30]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[43:44], v[41:42]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[45:46], -v[41:42], v[43:44], 1.0
	v_fma_f64 v[43:44], v[43:44], v[45:46], v[43:44]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[45:46], -v[41:42], v[43:44], 1.0
	v_fma_f64 v[43:44], v[43:44], v[45:46], v[43:44]
	v_div_scale_f64 v[45:46], vcc_lo, v[29:30], v[20:21], v[29:30]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[47:48], v[45:46], v[43:44]
	v_fma_f64 v[41:42], -v[41:42], v[47:48], v[45:46]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[41:42], v[41:42], v[43:44], v[47:48]
	v_div_fixup_f64 v[29:30], v[41:42], v[20:21], v[29:30]
.LBB0_15:                               ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_dual_mul_f32 v8, s14, v7 :: v_dual_fmac_f32 v5, v17, v13
	v_fma_f32 v0, v15, v11, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v4, 0x3fb8aa3b, v5
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v5
	v_rndne_f32_e32 v12, v4
	v_mul_f32_e32 v9, s14, v9
	v_fma_f32 v11, 0x3fb8aa3b, v5, -v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v0, v9, v0, v8
	v_fmac_f32_e32 v11, 0x32a5705f, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_cvt_f64_f32_e32 v[6:7], v0
	v_sub_f32_e32 v0, v4, v12
	v_cvt_f64_f32_e32 v[2:3], v9
	v_cvt_i32_f32_e32 v4, v12
	v_add_f32_e32 v0, v0, v11
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_exp_f32_e32 v0, v0
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v0, v0, v4
	v_cndmask_b32_e32 v0, 0, v0, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v13, 0x7f800000, v0, vcc_lo
	v_cvt_f64_f32_e32 v[11:12], v13
	v_mul_f64 v[4:5], v[2:3], 0.5
	v_fma_f64 v[6:7], v[2:3], 0.5, v[6:7]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[4:5], -v[4:5], v[11:12], v[6:7]
	v_cvt_f32_f64_e32 v4, v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cmp_gt_f32_e32 vcc_lo, 0, v4
	v_cndmask_b32_e32 v0, 0x80000000, v4, vcc_lo
	v_fmac_f32_e32 v0, v9, v13
	v_cndmask_b32_e64 v9, v4, 0, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v4, v9, v0
	v_cmp_nlt_f32_e32 vcc_lo, 1.0, v4
                                        ; implicit-def: $vgpr4_vgpr5
	s_and_saveexec_b32 s0, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB0_17
; %bb.16:                               ;   in Loop: Header=BB0_11 Depth=1
	v_div_scale_f32 v4, null, s14, s14, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v5, v4
	s_waitcnt_depctr 0xfff
	v_fma_f32 v6, -v4, v5, 1.0
	v_fmac_f32_e32 v5, v6, v5
	v_div_scale_f32 v6, vcc_lo, v0, s14, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v7, v6, v5
	v_fma_f32 v11, -v4, v7, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v7, v11, v5
	v_fma_f32 v4, -v4, v7, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v4, v4, v5, v7
	v_div_fixup_f32 v0, v4, s14, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_cvt_f64_f32_e32 v[4:5], v0
.LBB0_17:                               ; %Flow505
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_and_not1_saveexec_b32 s0, s0
	s_cbranch_execz .LBB0_19
; %bb.18:                               ;   in Loop: Header=BB0_11 Depth=1
	v_cvt_f64_f32_e32 v[4:5], v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f64 v[4:5], -v[4:5], 1.0
	v_div_scale_f64 v[6:7], null, v[22:23], v[22:23], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f64_e32 v[11:12], v[6:7]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[13:14], -v[6:7], v[11:12], 1.0
	v_fma_f64 v[11:12], v[11:12], v[13:14], v[11:12]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f64 v[13:14], -v[6:7], v[11:12], 1.0
	v_fma_f64 v[11:12], v[11:12], v[13:14], v[11:12]
	v_div_scale_f64 v[13:14], vcc_lo, v[4:5], v[22:23], v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[15:16], v[13:14], v[11:12]
	v_fma_f64 v[6:7], -v[6:7], v[15:16], v[13:14]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f64 v[6:7], v[6:7], v[11:12], v[15:16]
	v_div_fixup_f64 v[4:5], v[6:7], v[22:23], v[4:5]
.LBB0_19:                               ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v0, v18, v35
	s_mov_b32 s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[6:7], 2, v[0:1]
	v_add_co_u32 v6, vcc_lo, s28, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v7, vcc_lo, s29, v7, vcc_lo
	global_load_b32 v7, v[6:7], off
	v_subrev_nc_u32_e32 v6, s44, v18
	v_lshlrev_b32_e32 v6, 2, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v6, v6, 0, s3
	v_add_nc_u32_e32 v6, v36, v6
	s_waitcnt vmcnt(0)
	v_cmpx_le_f32_e32 s10, v7
	s_xor_b32 s25, exec_lo, s0
	s_cbranch_execz .LBB0_48
; %bb.20:                               ;   in Loop: Header=BB0_11 Depth=1
	v_div_scale_f32 v11, null, s14, s14, v9
	v_div_scale_f32 v14, null, s13, s13, v10
	v_div_scale_f32 v41, s0, v10, s13, v10
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_rcp_f32_e32 v15, v11
	v_rcp_f32_e32 v16, v14
	v_cvt_f32_f64_e32 v13, v[4:5]
	v_div_scale_f32 v17, vcc_lo, v9, s14, v9
	s_waitcnt_depctr 0xfff
	v_fma_f32 v7, -v11, v15, 1.0
	v_fma_f32 v12, -v14, v16, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_dual_fmac_f32 v15, v7, v15 :: v_dual_fmac_f32 v16, v12, v16
	v_cvt_f32_f64_e32 v12, v[29:30]
	v_mul_f32_e32 v5, v41, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v19, -v14, v5, v41
	v_dual_mul_f32 v4, v17, v15 :: v_dual_fmac_f32 v5, v19, v16
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_fma_f32 v7, -v11, v4, v17
	v_ashrrev_i32_e32 v19, 31, v18
	v_fma_f32 v14, -v14, v5, v41
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v4, v7, v15 :: v_dual_mov_b32 v7, v1
	v_fma_f32 v11, -v11, v4, v17
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[6:7], 2, v[6:7]
	v_div_fmas_f32 v11, v11, v15, v4
	s_mov_b32 vcc_lo, s0
	v_cmp_ne_u32_e64 s0, 0, v18
	v_div_fmas_f32 v14, v14, v16, v5
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_add_co_u32 v6, vcc_lo, s40, v6
	v_lshlrev_b64 v[4:5], 2, v[18:19]
	v_div_fixup_f32 v11, v11, s14, v9
	v_div_fixup_f32 v10, v14, s13, v10
	v_add_co_ci_u32_e32 v7, vcc_lo, s41, v7, vcc_lo
	s_and_not1_b32 vcc_lo, exec_lo, s11
	global_store_b128 v[6:7], v[10:13], off
	s_cbranch_vccnz .LBB0_35
; %bb.21:                               ; %.preheader343
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_mov_b32_e32 v13, 0
	s_mov_b32 s33, 0
	s_and_saveexec_b32 s2, s0
	s_cbranch_execz .LBB0_25
; %bb.22:                               ; %.lr.ph351.preheader
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_mov_b32_e32 v13, 0
	s_mov_b64 s[0:1], s[20:21]
	s_mov_b32 s34, 0
.LBB0_23:                               ; %.lr.ph351
                                        ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v6, v1, s[0:1]
	s_add_i32 s34, s34, 1
	s_add_u32 s0, s0, 4
	v_cmp_ge_u32_e32 vcc_lo, s34, v18
	s_addc_u32 s1, s1, 0
	s_or_b32 s33, vcc_lo, s33
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v13, v6, v13
	s_and_not1_b32 exec_lo, exec_lo, s33
	s_cbranch_execnz .LBB0_23
; %bb.24:                               ; %Flow501
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s33
.LBB0_25:                               ; %Flow502
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s2
	v_mov_b32_e32 v14, v13
	s_mov_b32 s1, exec_lo
	v_cmpx_gt_u32_e64 s45, v18
	s_cbranch_execz .LBB0_29
; %bb.26:                               ; %.lr.ph355.preheader
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_add_co_u32 v6, vcc_lo, s20, v4
	v_add_co_ci_u32_e32 v7, vcc_lo, s21, v5, vcc_lo
	v_dual_mov_b32 v9, v18 :: v_dual_mov_b32 v14, v13
	s_mov_b32 s2, 0
.LBB0_27:                               ; %.lr.ph355
                                        ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v10, v[6:7], off
	v_add_nc_u32_e32 v9, 1, v9
	v_add_co_u32 v6, vcc_lo, v6, 4
	v_add_co_ci_u32_e32 v7, vcc_lo, 0, v7, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_le_u32_e64 s0, s45, v9
	s_or_b32 s2, s0, s2
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v14, v10, v14
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_27
; %bb.28:                               ; %Flow499
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
.LBB0_29:                               ; %Flow500
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s1
	v_mov_b32_e32 v19, v1
	s_mov_b32 s1, exec_lo
	v_lshlrev_b64 v[6:7], 2, v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v6, vcc_lo, s20, v6
	v_add_co_ci_u32_e32 v7, vcc_lo, s21, v7, vcc_lo
	global_load_b32 v15, v[6:7], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u32_e32 0, v15
	s_cbranch_execz .LBB0_34
; %bb.30:                               ; %.lr.ph361
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_cvt_f64_f32_e32 v[6:7], v8
	v_cvt_f64_f32_e32 v[8:9], v40
	v_lshl_add_u32 v10, v18, 1, v37
	v_mul_lo_u32 v19, v14, v34
	v_mul_lo_u32 v17, v0, s46
	s_mov_b32 s2, 0
	s_mov_b32 s33, 0
	v_mul_lo_u32 v16, v10, s46
	v_mad_u64_u32 v[10:11], null, v14, v24, v[13:14]
	v_mad_u64_u32 v[11:12], null, v39, v14, v[13:14]
	v_lshl_add_u32 v12, v13, 1, v19
.LBB0_31:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_32 Depth 3
	s_lshl_b32 s34, s33, 1
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_dual_mov_b32 v13, 0 :: v_dual_add_nc_u32 v0, s34, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[29:30], 2, v[0:1]
	v_add_nc_u32_e32 v0, s33, v10
	v_add_co_u32 v29, vcc_lo, s30, v29
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v30, vcc_lo, s31, v30, vcc_lo
	global_load_b64 v[29:30], v[29:30], off
	s_waitcnt vmcnt(0)
	v_cvt_f64_f32_e32 v[40:41], v29
	v_cvt_f64_f32_e32 v[29:30], v30
	v_add_f64 v[40:41], v[40:41], 0.5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_f64 v[29:30], v[29:30], 0.5
	v_fma_f64 v[40:41], v[40:41], v[2:3], v[6:7]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[29:30], v[29:30], v[27:28], v[8:9]
	v_div_scale_f64 v[42:43], null, v[22:23], v[22:23], v[40:41]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_div_scale_f64 v[44:45], null, v[20:21], v[20:21], v[29:30]
	v_div_scale_f64 v[54:55], vcc_lo, v[40:41], v[22:23], v[40:41]
	v_rcp_f64_e32 v[46:47], v[42:43]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_rcp_f64_e32 v[48:49], v[44:45]
	s_waitcnt_depctr 0xfff
	v_fma_f64 v[50:51], -v[42:43], v[46:47], 1.0
	v_fma_f64 v[52:53], -v[44:45], v[48:49], 1.0
	v_fma_f64 v[46:47], v[46:47], v[50:51], v[46:47]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[48:49], v[48:49], v[52:53], v[48:49]
	v_fma_f64 v[50:51], -v[42:43], v[46:47], 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[52:53], -v[44:45], v[48:49], 1.0
	v_fma_f64 v[46:47], v[46:47], v[50:51], v[46:47]
	v_div_scale_f64 v[50:51], s0, v[29:30], v[20:21], v[29:30]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fma_f64 v[48:49], v[48:49], v[52:53], v[48:49]
	v_mul_f64 v[52:53], v[54:55], v[46:47]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_f64 v[56:57], v[50:51], v[48:49]
	v_fma_f64 v[42:43], -v[42:43], v[52:53], v[54:55]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[44:45], -v[44:45], v[56:57], v[50:51]
	v_div_fmas_f64 v[42:43], v[42:43], v[46:47], v[52:53]
	s_mov_b32 vcc_lo, s0
	s_mov_b32 s0, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_div_fmas_f64 v[44:45], v[44:45], v[48:49], v[56:57]
	v_div_fixup_f64 v[40:41], v[42:43], v[22:23], v[40:41]
	v_lshlrev_b64 v[42:43], 2, v[0:1]
	v_mov_b32_e32 v0, v11
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_div_fixup_f64 v[29:30], v[44:45], v[20:21], v[29:30]
	v_cvt_f32_f64_e32 v40, v[40:41]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cvt_f32_f64_e32 v41, v[29:30]
	v_dual_mov_b32 v30, v1 :: v_dual_add_nc_u32 v29, s34, v16
	v_lshlrev_b64 v[29:30], 2, v[29:30]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v29, vcc_lo, s4, v29
	v_add_co_ci_u32_e32 v30, vcc_lo, s5, v30, vcc_lo
	v_add_co_u32 v42, vcc_lo, s36, v42
	v_add_co_ci_u32_e32 v43, vcc_lo, s37, v43, vcc_lo
	global_store_b64 v[29:30], v[40:41], off
	global_load_b32 v19, v[42:43], off
.LBB0_32:                               ;   Parent Loop BB0_11 Depth=1
                                        ;     Parent Loop BB0_31 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[29:30], 2, v[0:1]
	v_add_co_u32 v29, vcc_lo, s36, v29
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(SALU_CYCLE_1)
	v_add_co_ci_u32_e32 v30, vcc_lo, s37, v30, vcc_lo
	global_load_b32 v29, v[29:30], off
	v_cvt_f32_u32_e32 v30, s0
	s_add_i32 s0, s0, 1
	s_cmp_eq_u32 s0, 4
	s_waitcnt vmcnt(0)
	v_cmp_gt_f32_e32 vcc_lo, v29, v19
	v_dual_cndmask_b32 v19, v19, v29 :: v_dual_add_nc_u32 v0, v0, v14
	v_cndmask_b32_e32 v13, v13, v30, vcc_lo
	s_cbranch_scc0 .LBB0_32
; %bb.33:                               ;   in Loop: Header=BB0_31 Depth=2
	v_add_nc_u32_e32 v0, s33, v17
	s_add_i32 s33, s33, 1
	v_add_nc_u32_e32 v11, 1, v11
	v_cmp_eq_u32_e32 vcc_lo, s33, v15
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[29:30], 2, v[0:1]
	s_or_b32 s2, vcc_lo, s2
	v_add_co_u32 v29, s0, s6, v29
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v30, s0, s7, v30, s0
	global_store_b32 v[29:30], v13, off
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_31
.LBB0_34:                               ; %Flow498
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
.LBB0_35:                               ; %.loopexit341
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s12
                                        ; implicit-def: $vgpr6
	s_cbranch_vccnz .LBB0_48
; %bb.36:                               ; %.preheader339
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_mov_b32_e32 v0, 0
	s_mov_b32 s33, 0
	s_mov_b32 s2, exec_lo
	v_cmpx_ne_u32_e32 0, v18
	s_cbranch_execz .LBB0_40
; %bb.37:                               ; %.lr.ph364.preheader
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_mov_b32_e32 v0, 0
	s_mov_b64 s[0:1], s[22:23]
	s_mov_b32 s34, 0
.LBB0_38:                               ; %.lr.ph364
                                        ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v2, v1, s[0:1]
	s_add_i32 s34, s34, 1
	s_add_u32 s0, s0, 4
	v_cmp_ge_u32_e32 vcc_lo, s34, v18
	s_addc_u32 s1, s1, 0
	s_or_b32 s33, vcc_lo, s33
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v0, v2, v0
	s_and_not1_b32 exec_lo, exec_lo, s33
	s_cbranch_execnz .LBB0_38
; %bb.39:                               ; %Flow494
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s33
.LBB0_40:                               ; %Flow495
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s2
	v_mov_b32_e32 v6, v0
	s_mov_b32 s1, exec_lo
	v_cmpx_gt_u32_e64 s45, v18
	s_cbranch_execz .LBB0_44
; %bb.41:                               ; %.lr.ph368.preheader
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_add_co_u32 v2, vcc_lo, s22, v4
	v_add_co_ci_u32_e32 v3, vcc_lo, s23, v5, vcc_lo
	v_mov_b32_e32 v4, v18
	v_mov_b32_e32 v6, v0
	s_mov_b32 s2, 0
.LBB0_42:                               ; %.lr.ph368
                                        ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v5, v[2:3], off
	v_add_nc_u32_e32 v4, 1, v4
	v_add_co_u32 v2, vcc_lo, v2, 4
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_le_u32_e64 s0, s45, v4
	s_or_b32 s2, s0, s2
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v6, v5, v6
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB0_42
; %bb.43:                               ; %Flow492
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
.LBB0_44:                               ; %Flow493
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_or_b32 exec_lo, exec_lo, s1
	v_mov_b32_e32 v19, v1
	s_mov_b32 s33, exec_lo
	v_lshlrev_b64 v[2:3], 2, v[18:19]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s22, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s23, v3, vcc_lo
	global_load_b32 v4, v[2:3], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u32_e32 0, v4
	s_cbranch_execz .LBB0_47
; %bb.45:                               ; %.lr.ph371
                                        ;   in Loop: Header=BB0_11 Depth=1
	v_lshl_add_u32 v0, v0, 1, v0
	s_mov_b32 s34, 0
	s_mov_b32 s35, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_mad_u64_u32 v[2:3], null, v6, v32, v[0:1]
.LBB0_46:                               ;   Parent Loop BB0_11 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, s35, v2
	v_dual_mov_b32 v6, v1 :: v_dual_add_nc_u32 v5, 1, v0
	v_lshlrev_b64 v[7:8], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[5:6], 2, v[5:6]
	v_add_co_u32 v7, vcc_lo, s38, v7
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v8, vcc_lo, s39, v8, vcc_lo
	v_add_co_u32 v5, vcc_lo, s38, v5
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v6, vcc_lo, s39, v6, vcc_lo
	global_load_b32 v3, v[7:8], off
	v_add_nc_u32_e32 v0, 2, v0
	v_lshlrev_b64 v[9:10], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v7, vcc_lo, s38, v9
	v_add_co_ci_u32_e32 v8, vcc_lo, s39, v10, vcc_lo
	s_clause 0x1
	global_load_b32 v12, v[5:6], off
	global_load_b32 v13, v[7:8], off
	v_dual_mov_b32 v6, v1 :: v_dual_add_nc_u32 v5, s35, v38
	s_add_i32 s35, s35, 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[6:7], 2, v[5:6]
	v_add_co_u32 v6, s0, s8, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_4) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v7, s0, s9, v7, s0
	s_waitcnt vmcnt(0)
	v_dual_mul_f32 v15, 0x3fb8aa3b, v12 :: v_dual_mul_f32 v16, 0x3fb8aa3b, v13
	v_cmp_ngt_f32_e64 s1, 0xc2ce8ed0, v13
	v_cmp_ngt_f32_e64 s2, 0xc2ce8ed0, v12
	v_fma_f32 v17, 0x3fb8aa3b, v12, -v15
	v_add_nc_u32_e32 v0, 1, v5
	v_fma_f32 v27, 0x3fb8aa3b, v13, -v16
	v_rndne_f32_e32 v28, v16
	v_rndne_f32_e32 v19, v15
	v_fmac_f32_e32 v17, 0x32a5705f, v12
	v_lshlrev_b64 v[8:9], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_dual_fmac_f32 v27, 0x32a5705f, v13 :: v_dual_sub_f32 v16, v16, v28
	v_add_nc_u32_e32 v0, 2, v5
	v_dual_sub_f32 v15, v15, v19 :: v_dual_add_nc_u32 v4, -1, v4
	v_add_co_u32 v8, s0, s8, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[10:11], 2, v[0:1]
	v_mul_f32_e32 v0, 0x3fb8aa3b, v3
	v_add_co_ci_u32_e64 v9, s0, s9, v9, s0
	v_cmp_eq_u32_e32 vcc_lo, 0, v4
	v_fma_f32 v5, 0x3fb8aa3b, v3, -v0
	v_rndne_f32_e32 v14, v0
	v_add_co_u32 v10, s0, s8, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v11, s0, s9, v11, s0
	v_dual_fmac_f32 v5, 0x32a5705f, v3 :: v_dual_sub_f32 v0, v0, v14
	v_cvt_i32_f32_e32 v14, v14
	v_cmp_ngt_f32_e64 s0, 0xc2ce8ed0, v3
	s_or_b32 s34, vcc_lo, s34
	s_delay_alu instid0(VALU_DEP_3)
	v_add_f32_e32 v0, v0, v5
	v_add_f32_e32 v5, v15, v17
	v_add_f32_e32 v15, v16, v27
	v_cvt_i32_f32_e32 v16, v19
	v_cvt_i32_f32_e32 v17, v28
	v_exp_f32_e32 v0, v0
	v_exp_f32_e32 v5, v5
	v_exp_f32_e32 v15, v15
	s_delay_alu instid0(TRANS32_DEP_3)
	v_ldexp_f32 v0, v0, v14
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v5, v5, v16
	v_ldexp_f32 v14, v15, v17
	v_cndmask_b32_e64 v0, 0, v0, s0
	v_cmp_nlt_f32_e64 s0, 0x42b17218, v3
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e64 v3, 0, v5, s2
	v_cndmask_b32_e64 v5, 0, v14, s1
	v_cmp_nlt_f32_e64 s1, 0x42b17218, v12
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v0, 0x7f800000, v0, s0
	v_cmp_nlt_f32_e64 s0, 0x42b17218, v13
	v_cndmask_b32_e64 v3, 0x7f800000, v3, s1
	s_delay_alu instid0(VALU_DEP_2)
	v_cndmask_b32_e64 v5, 0x7f800000, v5, s0
	s_clause 0x2
	global_store_b32 v[6:7], v0, off
	global_store_b32 v[8:9], v3, off
	global_store_b32 v[10:11], v5, off
	s_and_not1_b32 exec_lo, exec_lo, s34
	s_cbranch_execnz .LBB0_46
.LBB0_47:                               ; %Flow491
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s33
                                        ; implicit-def: $vgpr6
.LBB0_48:                               ; %Flow504
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_and_not1_saveexec_b32 s0, s25
	s_cbranch_execz .LBB0_50
; %bb.49:                               ;   in Loop: Header=BB0_11 Depth=1
	v_mov_b32_e32 v7, v1
	v_mov_b32_e32 v0, v1
	v_mov_b32_e32 v2, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[3:4], 2, v[6:7]
	v_add_co_u32 v5, vcc_lo, s40, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v6, vcc_lo, s41, v4, vcc_lo
	v_mov_b32_e32 v3, v1
	global_store_b128 v[5:6], v[0:3], off
.LBB0_50:                               ; %.loopexit
                                        ;   in Loop: Header=BB0_11 Depth=1
	s_or_b32 exec_lo, exec_lo, s0
	v_add_nc_u32_e32 v18, 1, v18
	v_add_nc_u32_e32 v38, s15, v38
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cmp_ge_u32_e32 vcc_lo, v18, v31
	s_or_b32 s24, vcc_lo, s24
	s_and_not1_b32 exec_lo, exec_lo, s24
	s_cbranch_execnz .LBB0_11
.LBB0_51:                               ; %.loopexit344
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 432
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 58
		.amdhsa_next_free_sgpr 48
		.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	_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f, .Lfunc_end0-_Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 3784
; NumSgprs: 50
; NumVgprs: 58
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 6
; VGPRBlocks: 7
; NumSGPRsForWavesPerEU: 50
; NumVGPRsForWavesPerEU: 58
; 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: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         12
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           1
        .value_kind:     by_value
      - .offset:         33
        .size:           1
        .value_kind:     by_value
      - .offset:         34
        .size:           1
        .value_kind:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         64
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         72
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         80
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         88
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         96
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         104
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         112
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         120
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         128
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         136
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         144
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         152
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         160
        .size:           8
        .value_kind:     global_buffer
      - .offset:         168
        .size:           4
        .value_kind:     by_value
      - .offset:         176
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         180
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         184
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         188
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         190
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         192
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         194
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         196
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         198
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         216
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         224
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         232
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         240
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 432
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f
    .private_segment_fixed_size: 0
    .sgpr_count:     50
    .sgpr_spill_count: 0
    .symbol:         _Z25cudaSNormalizeROIs_kerneljjjjjjjjbbbffPKfS0_PKjS2_S0_S0_S0_S0_S0_S0_PfPiS3_S3_S3_f.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     58
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
