#include "includes.h"
__global__ void cudaS_ssdToOutput_kernels(  unsigned int batchSize, unsigned int nbClass, unsigned int nbAnchors, unsigned int channelWidth, unsigned int channelHeight, unsigned int nbProposals, unsigned int* nbValidROIs, unsigned int cls, unsigned int totalParts, unsigned int totalTemplates, unsigned int maxParts, unsigned int maxTemplates, unsigned int cumulParts, unsigned int cumulTemplates, unsigned int nbParts, unsigned int nbTemplates, float xRatio, float yRatio, float xOutputRatio, float yOutputRatio, const float* roi_bbox, const float* roi_anchors, const float* anchors, const float* inputs_parts, const float* inputs_templates, float* outputs)
{
const int batchPos = blockIdx.z;
const int proposal = (threadIdx.x & 0x1f) + blockIdx.x*blockDim.x;
const int ptIdx = blockIdx.y;

const int nbDetectedObject  = (int) nbValidROIs[batchPos];
const int nbIdx = 6;
if(proposal < nbProposals)
{
const unsigned int n = proposal + cls*nbProposals + batchPos*nbProposals*nbClass;

if(proposal < nbDetectedObject)
{
if(ptIdx == 0)
{
outputs[0 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = roi_bbox[0 + 5*proposal + batchPos*nbProposals*5];
outputs[1 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = roi_bbox[1 + 5*proposal + batchPos*nbProposals*5];
outputs[2 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = roi_bbox[2 + 5*proposal + batchPos*nbProposals*5];
outputs[3 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = roi_bbox[3 + 5*proposal + batchPos*nbProposals*5];
outputs[4 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = roi_bbox[4 + 5*proposal + batchPos*nbProposals*5];
outputs[5 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = (float) cls;
}

if(ptIdx < nbParts && totalParts > 0)
{
const unsigned int xa   = roi_anchors[0 + 5*proposal + batchPos*nbProposals*5];
const unsigned int ya   = roi_anchors[1 + 5*proposal + batchPos*nbProposals*5];
const unsigned int k    = roi_anchors[2 + 5*proposal + batchPos*nbProposals*5];


const int yIdx = xa
+ ya*channelWidth
+ (k*nbParts*2 + cumulParts + ptIdx*2)*channelHeight*channelWidth
+ batchPos*channelHeight*channelWidth*nbAnchors*2*totalParts;
const int xIdx = xa
+ ya*channelWidth
+ (k*nbParts*2 + cumulParts + ptIdx*2 + 1)*channelHeight*channelWidth
+ batchPos*channelHeight*channelWidth*nbAnchors*2*totalParts;


const float partY = inputs_parts[yIdx];
const float partX = inputs_parts[xIdx];

const int xa0 = (int)(anchors[k*4] + xa * xRatio);
const int ya0 = (int)(anchors[k*4 + 1] + ya * yRatio);
const int xa1 = (int)(anchors[k*4 + 2] + xa * xRatio);
const int ya1 = (int)(anchors[k*4 + 3] + ya * yRatio);

// Anchors width and height
const int wa = xa1 - xa0;
const int ha = ya1 - ya0;

// Anchor center coordinates (xac, yac)
const float xac = xa0 + wa / 2.0;
const float yac = ya0 + ha / 2.0;
const float predPartY = ((partY) * ha + yac)*yOutputRatio ;
const float predPartX = ((partX) * wa + xac)*xOutputRatio ;

outputs[ptIdx*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = predPartY;
outputs[ptIdx*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = predPartX;

}
else if(ptIdx < maxParts && totalParts > 0)
{
outputs[ptIdx*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
}

///for(unsigned int t = 0; t < nbTemplates; ++t)
if(ptIdx < nbTemplates && totalTemplates > 0)
{
const unsigned int xa   = roi_anchors[0 + 5*proposal + batchPos*nbProposals*5];
const unsigned int ya   = roi_anchors[1 + 5*proposal + batchPos*nbProposals*5];
const unsigned int k    = roi_anchors[2 + 5*proposal + batchPos*nbProposals*5];

const int yIdx = xa
+ ya*channelWidth
+ (k*nbTemplates*3 + cumulTemplates + ptIdx*3)*channelHeight*channelWidth
+ batchPos*channelHeight*channelWidth*nbAnchors*3*totalTemplates;
const int xIdx = xa
+ ya*channelWidth
+ (k*nbTemplates*3 + cumulTemplates + ptIdx*3 + 1)*channelHeight*channelWidth
+ batchPos*channelHeight*channelWidth*nbAnchors*3*totalTemplates;
const int zIdx = xa
+ ya*channelWidth
+ (k*nbTemplates*3 + cumulTemplates + ptIdx*3 + 2)*channelHeight*channelWidth
+ batchPos*channelHeight*channelWidth*nbAnchors*3*totalTemplates;


const float templateY = expf(inputs_templates[yIdx]);
const float templateX = expf(inputs_templates[xIdx]);
const float templateZ = expf(inputs_templates[zIdx]);

outputs[ptIdx*3 + maxParts*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = templateY;
outputs[ptIdx*3 + maxParts*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = templateX;
outputs[ptIdx*3 + maxParts*2 + 2 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = templateZ;

}
else if(ptIdx < maxTemplates && totalTemplates > 0)
{
outputs[ptIdx*3 + maxParts*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*3 + maxParts*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*3 + maxParts*2 + 2 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
}

}
else
{
outputs[0 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = 0.0;
outputs[1 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = 0.0;
outputs[2 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = 0.0;
outputs[3 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = 0.0;
outputs[4 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = 0.0;

//for(unsigned int p = 0; p < nbParts; ++p)
if(ptIdx < maxParts && totalParts > 0)
{
outputs[ptIdx*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
}

//for(unsigned int t = 0;t < nbTemplates; ++t)
if(ptIdx < maxTemplates && totalTemplates > 0)
{
outputs[ptIdx*3 + maxParts*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*3 + maxParts*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*3 + maxParts*2 + 2 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
}

}
}
}
	.file	"tmpxft_002c69ec_00000000-6_cudaS_ssdToOutput_kernels.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	_Z81__device_stub__Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_PfjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.type	_Z81__device_stub__Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_PfjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf, @function
_Z81__device_stub__Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_PfjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf:
.LFB2032:
	.cfi_startproc
	endbr64
	subq	$392, %rsp
	.cfi_def_cfa_offset 400
	movq	400(%rsp), %rax
	movl	%edi, 92(%rsp)
	leaq	120(%rsp), %rdi
	movl	%esi, 88(%rsp)
	leaq	132(%rsp), %rsi
	movq	%rax, 64(%rsp)
	movq	480(%rsp), %rax
	movl	%edx, 84(%rsp)
	leaq	104(%rsp), %rdx
	movq	%rax, 40(%rsp)
	movq	488(%rsp), %rax
	movl	%ecx, 80(%rsp)
	leaq	112(%rsp), %rcx
	movq	%rax, 32(%rsp)
	movq	496(%rsp), %rax
	movl	%r8d, 76(%rsp)
	movq	%rax, 24(%rsp)
	movq	504(%rsp), %rax
	movl	%r9d, 72(%rsp)
	movq	%rax, 16(%rsp)
	movq	512(%rsp), %rax
	movss	%xmm0, 60(%rsp)
	movq	%rax, 8(%rsp)
	movq	520(%rsp), %rax
	movss	%xmm1, 56(%rsp)
	movq	%rax, (%rsp)
	movss	%xmm2, 52(%rsp)
	movss	%xmm3, 48(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 376(%rsp)
	xorl	%eax, %eax
	leaq	92(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	88(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	84(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	80(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	76(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	72(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	408(%rsp), %rax
	movq	%rax, 224(%rsp)
	leaq	416(%rsp), %rax
	movq	%rax, 232(%rsp)
	leaq	424(%rsp), %rax
	movq	%rax, 240(%rsp)
	leaq	432(%rsp), %rax
	movq	%rax, 248(%rsp)
	leaq	440(%rsp), %rax
	movq	%rax, 256(%rsp)
	leaq	448(%rsp), %rax
	movq	%rax, 264(%rsp)
	leaq	456(%rsp), %rax
	movq	%rax, 272(%rsp)
	leaq	464(%rsp), %rax
	movq	%rax, 280(%rsp)
	leaq	472(%rsp), %rax
	movq	%rax, 288(%rsp)
	leaq	60(%rsp), %rax
	movq	%rax, 296(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 304(%rsp)
	leaq	52(%rsp), %rax
	movq	%rax, 312(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 320(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 328(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 336(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 344(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 352(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 360(%rsp)
	movq	%rsp, %rax
	movq	%rax, 368(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 120(%rsp)
	movl	$1, 128(%rsp)
	movq	%rax, 132(%rsp)
	movl	$1, 140(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	112(%rsp)
	.cfi_def_cfa_offset 408
	leaq	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf(%rip), %rdi
	pushq	112(%rsp)
	.cfi_def_cfa_offset 416
	movq	148(%rsp), %rcx
	movl	156(%rsp), %r8d
	movq	136(%rsp), %rsi
	movl	144(%rsp), %edx
	leaq	184(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 408
	popq	%rdx
	.cfi_def_cfa_offset 400
.L2:
	movq	376(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$392, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2032:
	.size	_Z81__device_stub__Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_PfjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf, .-_Z81__device_stub__Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_PfjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.globl	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.type	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf, @function
_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf:
.LFB2033:
	.cfi_startproc
	endbr64
	jmp	_Z81__device_stub__Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_PfjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.cfi_endproc
.LFE2033:
	.size	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf, .-_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf"
	.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	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf(%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 cudaS_ssdToOutput_kernels(  unsigned int batchSize, unsigned int nbClass, unsigned int nbAnchors, unsigned int channelWidth, unsigned int channelHeight, unsigned int nbProposals, unsigned int* nbValidROIs, unsigned int cls, unsigned int totalParts, unsigned int totalTemplates, unsigned int maxParts, unsigned int maxTemplates, unsigned int cumulParts, unsigned int cumulTemplates, unsigned int nbParts, unsigned int nbTemplates, float xRatio, float yRatio, float xOutputRatio, float yOutputRatio, const float* roi_bbox, const float* roi_anchors, const float* anchors, const float* inputs_parts, const float* inputs_templates, float* outputs)
{
const int batchPos = blockIdx.z;
const int proposal = (threadIdx.x & 0x1f) + blockIdx.x*blockDim.x;
const int ptIdx = blockIdx.y;

const int nbDetectedObject  = (int) nbValidROIs[batchPos];
const int nbIdx = 6;
if(proposal < nbProposals)
{
const unsigned int n = proposal + cls*nbProposals + batchPos*nbProposals*nbClass;

if(proposal < nbDetectedObject)
{
if(ptIdx == 0)
{
outputs[0 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = roi_bbox[0 + 5*proposal + batchPos*nbProposals*5];
outputs[1 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = roi_bbox[1 + 5*proposal + batchPos*nbProposals*5];
outputs[2 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = roi_bbox[2 + 5*proposal + batchPos*nbProposals*5];
outputs[3 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = roi_bbox[3 + 5*proposal + batchPos*nbProposals*5];
outputs[4 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = roi_bbox[4 + 5*proposal + batchPos*nbProposals*5];
outputs[5 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = (float) cls;
}

if(ptIdx < nbParts && totalParts > 0)
{
const unsigned int xa   = roi_anchors[0 + 5*proposal + batchPos*nbProposals*5];
const unsigned int ya   = roi_anchors[1 + 5*proposal + batchPos*nbProposals*5];
const unsigned int k    = roi_anchors[2 + 5*proposal + batchPos*nbProposals*5];


const int yIdx = xa
+ ya*channelWidth
+ (k*nbParts*2 + cumulParts + ptIdx*2)*channelHeight*channelWidth
+ batchPos*channelHeight*channelWidth*nbAnchors*2*totalParts;
const int xIdx = xa
+ ya*channelWidth
+ (k*nbParts*2 + cumulParts + ptIdx*2 + 1)*channelHeight*channelWidth
+ batchPos*channelHeight*channelWidth*nbAnchors*2*totalParts;


const float partY = inputs_parts[yIdx];
const float partX = inputs_parts[xIdx];

const int xa0 = (int)(anchors[k*4] + xa * xRatio);
const int ya0 = (int)(anchors[k*4 + 1] + ya * yRatio);
const int xa1 = (int)(anchors[k*4 + 2] + xa * xRatio);
const int ya1 = (int)(anchors[k*4 + 3] + ya * yRatio);

// Anchors width and height
const int wa = xa1 - xa0;
const int ha = ya1 - ya0;

// Anchor center coordinates (xac, yac)
const float xac = xa0 + wa / 2.0;
const float yac = ya0 + ha / 2.0;
const float predPartY = ((partY) * ha + yac)*yOutputRatio ;
const float predPartX = ((partX) * wa + xac)*xOutputRatio ;

outputs[ptIdx*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = predPartY;
outputs[ptIdx*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = predPartX;

}
else if(ptIdx < maxParts && totalParts > 0)
{
outputs[ptIdx*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
}

///for(unsigned int t = 0; t < nbTemplates; ++t)
if(ptIdx < nbTemplates && totalTemplates > 0)
{
const unsigned int xa   = roi_anchors[0 + 5*proposal + batchPos*nbProposals*5];
const unsigned int ya   = roi_anchors[1 + 5*proposal + batchPos*nbProposals*5];
const unsigned int k    = roi_anchors[2 + 5*proposal + batchPos*nbProposals*5];

const int yIdx = xa
+ ya*channelWidth
+ (k*nbTemplates*3 + cumulTemplates + ptIdx*3)*channelHeight*channelWidth
+ batchPos*channelHeight*channelWidth*nbAnchors*3*totalTemplates;
const int xIdx = xa
+ ya*channelWidth
+ (k*nbTemplates*3 + cumulTemplates + ptIdx*3 + 1)*channelHeight*channelWidth
+ batchPos*channelHeight*channelWidth*nbAnchors*3*totalTemplates;
const int zIdx = xa
+ ya*channelWidth
+ (k*nbTemplates*3 + cumulTemplates + ptIdx*3 + 2)*channelHeight*channelWidth
+ batchPos*channelHeight*channelWidth*nbAnchors*3*totalTemplates;


const float templateY = expf(inputs_templates[yIdx]);
const float templateX = expf(inputs_templates[xIdx]);
const float templateZ = expf(inputs_templates[zIdx]);

outputs[ptIdx*3 + maxParts*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = templateY;
outputs[ptIdx*3 + maxParts*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = templateX;
outputs[ptIdx*3 + maxParts*2 + 2 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = templateZ;

}
else if(ptIdx < maxTemplates && totalTemplates > 0)
{
outputs[ptIdx*3 + maxParts*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*3 + maxParts*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*3 + maxParts*2 + 2 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
}

}
else
{
outputs[0 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = 0.0;
outputs[1 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = 0.0;
outputs[2 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = 0.0;
outputs[3 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = 0.0;
outputs[4 + n*(nbIdx + maxParts*2 + maxTemplates*3)] = 0.0;

//for(unsigned int p = 0; p < nbParts; ++p)
if(ptIdx < maxParts && totalParts > 0)
{
outputs[ptIdx*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
}

//for(unsigned int t = 0;t < nbTemplates; ++t)
if(ptIdx < maxTemplates && totalTemplates > 0)
{
outputs[ptIdx*3 + maxParts*2 + 0 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*3 + maxParts*2 + 1 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
outputs[ptIdx*3 + maxParts*2 + 2 + nbIdx + n*(nbIdx + maxParts*2 + maxTemplates*3) ] = 0.0;
}

}
}
}
	.text
	.file	"cudaS_ssdToOutput_kernels.hip"
	.globl	_Z40__device_stub__cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf # -- Begin function _Z40__device_stub__cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.type	_Z40__device_stub__cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf,@function
_Z40__device_stub__cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf: # @_Z40__device_stub__cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.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	$304, %rsp                      # imm = 0x130
	.cfi_def_cfa_offset 352
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	44(%rsp), %r10
	movl	%edi, (%r10)
	leaq	40(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	36(%rsp), %r11
	movl	%edx, (%r11)
	leaq	32(%rsp), %r14
	movl	%ecx, (%r14)
	leaq	28(%rsp), %r15
	movl	%r8d, (%r15)
	leaq	24(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	20(%rsp), %rax
	movss	%xmm0, (%rax)
	leaq	16(%rsp), %rcx
	movss	%xmm1, (%rcx)
	leaq	12(%rsp), %rdx
	movss	%xmm2, (%rdx)
	leaq	8(%rsp), %rsi
	movss	%xmm3, (%rsi)
	leaq	96(%rsp), %rbx
	movq	%r10, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%r11, 16(%rbx)
	movq	%r14, 24(%rbx)
	movq	%r15, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	352(%rsp), %rdi
	movq	%rdi, 48(%rbx)
	leaq	360(%rsp), %rdi
	movq	%rdi, 56(%rbx)
	leaq	368(%rsp), %rdi
	movq	%rdi, 64(%rbx)
	leaq	376(%rsp), %rdi
	movq	%rdi, 72(%rbx)
	leaq	384(%rsp), %rdi
	movq	%rdi, 80(%rbx)
	leaq	392(%rsp), %rdi
	movq	%rdi, 88(%rbx)
	leaq	400(%rsp), %rdi
	movq	%rdi, 96(%rbx)
	leaq	408(%rsp), %rdi
	movq	%rdi, 104(%rbx)
	leaq	416(%rsp), %rdi
	movq	%rdi, 112(%rbx)
	leaq	424(%rsp), %rdi
	movq	%rdi, 120(%rbx)
	movq	%rax, 128(%rbx)
	movq	%rcx, 136(%rbx)
	movq	%rdx, 144(%rbx)
	movq	%rsi, 152(%rbx)
	leaq	432(%rsp), %rax
	movq	%rax, 160(%rbx)
	leaq	440(%rsp), %rax
	movq	%rax, 168(%rbx)
	leaq	448(%rsp), %rax
	movq	%rax, 176(%rbx)
	leaq	456(%rsp), %rax
	movq	%rax, 184(%rbx)
	leaq	464(%rsp), %rax
	movq	%rax, 192(%rbx)
	leaq	472(%rsp), %rax
	movq	%rax, 200(%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	$_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$320, %rsp                      # imm = 0x140
	.cfi_adjust_cfa_offset -320
	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__cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf, .Lfunc_end0-_Z40__device_stub__cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.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	$_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf, %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	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf,@object # @_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.section	.rodata,"a",@progbits
	.globl	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.p2align	3, 0x0
_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf:
	.quad	_Z40__device_stub__cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.size	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf, 8

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

	.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__cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf ; -- Begin function _Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.globl	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
	.p2align	8
	.type	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf,@function
_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf: ; @_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
; %bb.0:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x94
	s_load_b32 s12, s[0:1], 0x14
	v_and_b32_e32 v0, 31, v0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s3, s3, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_mad_u64_u32 v[2:3], null, s13, s3, v[0:1]
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_u32_e64 s12, v2
	s_cbranch_execz .LBB0_32
; %bb.1:
	s_clause 0x1
	s_load_b64 s[16:17], s[0:1], 0x18
	s_load_b128 s[28:31], s[0:1], 0x4
	s_mov_b32 s2, s15
	s_ashr_i32 s3, s15, 31
	s_load_b256 s[4:11], s[0:1], 0x20
	s_lshl_b64 s[18:19], s[2:3], 2
	s_mul_i32 s15, s15, s12
                                        ; implicit-def: $vgpr1
	s_waitcnt lgkmcnt(0)
	s_add_u32 s16, s16, s18
	s_addc_u32 s17, s17, s19
	s_mul_i32 s13, s15, s28
	s_load_b32 s3, s[16:17], 0x0
	s_load_b128 s[24:27], s[0:1], 0x78
	s_mul_i32 s12, s4, s12
                                        ; implicit-def: $sgpr16
	s_delay_alu instid0(SALU_CYCLE_1)
	v_add3_u32 v4, s12, s13, v2
	s_mov_b32 s12, 0
	s_waitcnt lgkmcnt(0)
	v_cmp_le_i32_e32 vcc_lo, s3, v2
	s_mul_i32 s3, s8, 3
	s_and_saveexec_b32 s13, vcc_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s13, exec_lo, s13
	s_cbranch_execz .LBB0_8
; %bb.2:
	s_lshl_b32 s16, s7, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s12, s16, s3
	s_add_i32 s12, s12, 6
	s_cmp_ge_u32 s14, s7
	v_mul_lo_u32 v0, v4, s12
	v_mov_b32_e32 v1, 0
	s_cselect_b32 s17, -1, 0
	s_cmp_eq_u32 s5, 0
	s_mov_b32 s12, 0
	s_cselect_b32 s18, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	s_or_b32 s17, s18, s17
	v_dual_mov_b32 v3, v1 :: v_dual_add_nc_u32 v2, 1, v0
	v_add_nc_u32_e32 v4, 2, v0
	v_lshlrev_b64 v[6:7], 2, v[0:1]
	v_dual_mov_b32 v5, v1 :: v_dual_add_nc_u32 v8, 3, v0
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	v_dual_mov_b32 v9, v1 :: v_dual_add_nc_u32 v10, 4, v0
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	v_add_co_u32 v6, vcc_lo, s26, v6
	v_mov_b32_e32 v11, v1
	v_add_co_ci_u32_e32 v7, vcc_lo, s27, v7, vcc_lo
	v_add_co_u32 v2, vcc_lo, s26, v2
	v_lshlrev_b64 v[8:9], 2, v[8:9]
	v_add_co_ci_u32_e32 v3, vcc_lo, s27, v3, vcc_lo
	v_add_co_u32 v4, vcc_lo, s26, v4
	v_lshlrev_b64 v[10:11], 2, v[10:11]
	v_add_co_ci_u32_e32 v5, vcc_lo, s27, v5, vcc_lo
	v_add_co_u32 v8, vcc_lo, s26, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s27, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v10, vcc_lo, s26, v10
	v_add_co_ci_u32_e32 v11, vcc_lo, s27, v11, vcc_lo
	s_and_b32 vcc_lo, exec_lo, s17
	s_clause 0x4
	global_store_b32 v[6:7], v1, off
	global_store_b32 v[2:3], v1, off
	global_store_b32 v[4:5], v1, off
	global_store_b32 v[8:9], v1, off
	global_store_b32 v[10:11], v1, off
	s_cbranch_vccnz .LBB0_4
; %bb.3:
	s_lshl_b32 s17, s14, 1
	v_mov_b32_e32 v3, v1
	v_add3_u32 v2, s17, 6, v0
	v_mov_b32_e32 v5, v1
	v_add3_u32 v4, s17, 7, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v2, vcc_lo, s26, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s27, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v4, vcc_lo, s26, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s27, v5, vcc_lo
	s_clause 0x1
	global_store_b32 v[2:3], v1, off
	global_store_b32 v[4:5], v1, off
.LBB0_4:
	s_cmp_ge_u32 s14, s8
	s_cselect_b32 s17, -1, 0
	s_cmp_eq_u32 s6, 0
	s_cselect_b32 s18, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s17, s18, s17
	s_and_b32 vcc_lo, exec_lo, s17
	s_cbranch_vccnz .LBB0_6
; %bb.5:
	s_mul_i32 s12, s14, 3
	v_mov_b32_e32 v2, 0
	s_add_i32 s16, s16, s12
	s_mov_b32 s12, -1
	v_add3_u32 v1, s16, 6, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[3:4], 2, v[1:2]
	v_add3_u32 v1, s16, 7, v0
	v_lshlrev_b64 v[5:6], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v3, vcc_lo, s26, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s27, v4, vcc_lo
	v_add3_u32 v1, s16, 8, v0
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s26, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s27, v6, vcc_lo
	s_mov_b32 s16, 0
	s_clause 0x1
	global_store_b32 v[3:4], v2, off
	global_store_b32 v[5:6], v2, off
	s_branch .LBB0_7
.LBB0_6:
                                        ; implicit-def: $sgpr16
                                        ; implicit-def: $vgpr1
.LBB0_7:                                ; %Flow441
	s_and_b32 s12, s12, exec_lo
                                        ; implicit-def: $vgpr4
                                        ; implicit-def: $vgpr2_vgpr3
.LBB0_8:                                ; %Flow440
	s_or_saveexec_b32 s13, s13
	v_mov_b32_e32 v3, s16
	s_xor_b32 exec_lo, exec_lo, s13
	s_cbranch_execz .LBB0_30
; %bb.9:
	s_clause 0x2
	s_load_b256 s[16:23], s[0:1], 0x58
	s_load_b128 s[36:39], s[0:1], 0x40
	s_load_b32 s1, s[0:1], 0x50
	v_add_nc_u32_e32 v0, s15, v2
	s_cmp_lg_u32 s14, 0
	s_mov_b32 s15, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_lshl_add_u32 v0, v0, 2, v0
	s_cbranch_scc1 .LBB0_11
; %bb.10:
	v_mov_b32_e32 v1, 0
	s_lshl_b32 s0, s7, 1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v5, 1, v0
	s_add_i32 s0, s0, s3
	v_mov_b32_e32 v6, v1
	v_lshlrev_b64 v[2:3], 2, v[0:1]
	s_add_i32 s0, s0, 6
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[5:6], 2, v[5:6]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s16, v2
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_add_co_ci_u32_e32 v3, vcc_lo, s17, v3, vcc_lo
	global_load_b32 v9, v[2:3], off
	v_mul_lo_u32 v2, v4, s0
	v_mov_b32_e32 v3, v1
	v_lshlrev_b64 v[7:8], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v7, vcc_lo, s26, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s27, v8, vcc_lo
	v_add_co_u32 v5, vcc_lo, s16, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s17, v6, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[7:8], v9, off
	global_load_b32 v3, v[5:6], off
	v_dual_mov_b32 v6, v1 :: v_dual_add_nc_u32 v5, 1, v2
	v_dual_mov_b32 v8, v1 :: v_dual_add_nc_u32 v7, 2, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[5:6], 2, v[5:6]
	v_lshlrev_b64 v[7:8], 2, v[7:8]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v5, vcc_lo, s26, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s27, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v7, vcc_lo, s16, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s17, v8, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[5:6], v3, off
	global_load_b32 v3, v[7:8], off
	v_dual_mov_b32 v6, v1 :: v_dual_add_nc_u32 v5, 2, v2
	v_dual_mov_b32 v8, v1 :: v_dual_add_nc_u32 v7, 3, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[5:6], 2, v[5:6]
	v_lshlrev_b64 v[7:8], 2, v[7:8]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v5, vcc_lo, s26, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s27, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v7, vcc_lo, s16, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s17, v8, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[5:6], v3, off
	global_load_b32 v3, v[7:8], off
	v_dual_mov_b32 v6, v1 :: v_dual_add_nc_u32 v5, 3, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], 2, v[5:6]
	v_add_co_u32 v5, vcc_lo, s26, v5
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v6, vcc_lo, s27, v6, vcc_lo
	s_waitcnt vmcnt(0)
	global_store_b32 v[5:6], v3, off
	v_add_nc_u32_e32 v5, 4, v2
	v_add_nc_u32_e32 v2, 5, v2
	v_dual_mov_b32 v8, v1 :: v_dual_add_nc_u32 v7, 4, v0
	v_mov_b32_e32 v6, v1
	v_mov_b32_e32 v3, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[7:8], 2, v[7:8]
	v_lshlrev_b64 v[5:6], 2, v[5:6]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[1:2], 2, v[2:3]
	v_add_co_u32 v7, vcc_lo, s16, v7
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_ci_u32_e32 v8, vcc_lo, s17, v8, vcc_lo
	v_add_co_u32 v5, vcc_lo, s26, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s27, v6, vcc_lo
	global_load_b32 v7, v[7:8], off
	v_add_co_u32 v1, vcc_lo, s26, v1
	v_cvt_f32_u32_e32 v8, s4
	v_add_co_ci_u32_e32 v2, vcc_lo, s27, v2, vcc_lo
	s_waitcnt vmcnt(0)
	s_clause 0x1
	global_store_b32 v[5:6], v7, off
	global_store_b32 v[1:2], v8, off
.LBB0_11:
	s_cmp_ge_u32 s14, s11
	s_cselect_b32 s4, -1, 0
	s_cmp_lg_u32 s5, 0
	s_cselect_b32 s0, -1, 0
	s_cmp_eq_u32 s5, 0
	s_waitcnt lgkmcnt(0)
	s_cselect_b32 s16, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s4, s16, s4
	s_and_b32 vcc_lo, exec_lo, s4
	s_cbranch_vccz .LBB0_14
; %bb.12:
	s_cmp_lt_u32 s14, s7
	s_cselect_b32 s4, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 s0, s0, s4
	s_mov_b32 s4, 0
	s_and_b32 vcc_lo, exec_lo, s0
	s_cbranch_vccz .LBB0_15
; %bb.13:
	s_lshl_b32 s0, s14, 1
	s_mov_b32 s15, -1
	s_mov_b32 s16, 0
	s_branch .LBB0_16
.LBB0_14:
	s_mov_b32 s4, -1
.LBB0_15:
                                        ; implicit-def: $sgpr16
                                        ; implicit-def: $sgpr0
.LBB0_16:                               ; %Flow
	v_dual_mov_b32 v1, s16 :: v_dual_mov_b32 v2, s16
	s_and_b32 vcc_lo, exec_lo, s4
	s_mul_i32 s4, s31, s30
	s_delay_alu instid0(SALU_CYCLE_1)
	s_mul_i32 s2, s4, s2
	s_cbranch_vccz .LBB0_18
; %bb.17:
	v_dual_mov_b32 v1, 0 :: v_dual_add_nc_u32 v2, 2, v0
	s_lshl_b32 s0, s14, 1
	s_mov_b32 s15, -1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mov_b32_e32 v3, v1
	v_lshlrev_b64 v[5:6], 2, v[0:1]
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s18, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s19, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	v_add_co_u32 v5, vcc_lo, s18, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s19, v6, vcc_lo
	global_load_b32 v7, v[2:3], off
	v_dual_mov_b32 v3, v1 :: v_dual_add_nc_u32 v2, 1, v0
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s18, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s19, v3, vcc_lo
	s_clause 0x1
	global_load_b32 v9, v[5:6], off
	global_load_b32 v3, v[2:3], off
	s_waitcnt vmcnt(2)
	v_cvt_u32_f32_e32 v10, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_mov_b32 v2, v1 :: v_dual_lshlrev_b32 v1, 2, v10
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v1, vcc_lo, s20, v1
	v_add_co_ci_u32_e32 v2, vcc_lo, s21, v2, vcc_lo
	s_waitcnt vmcnt(0)
	v_cvt_u32_f32_e32 v11, v3
	global_load_b128 v[5:8], v[1:2], off
	v_mul_lo_u32 v1, s11, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b32_e32 v2, 1, v1
	v_cvt_u32_f32_e32 v1, v9
	v_add3_u32 v9, s0, s9, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_mad_u64_u32 v[2:3], null, v11, s30, v[1:2]
	s_mul_i32 s9, s2, s29
	v_cvt_f32_u32_e32 v1, v1
	v_mul_lo_u32 v3, s4, v9
	s_mul_i32 s9, s9, s5
	s_delay_alu instid0(SALU_CYCLE_1)
	s_lshl_b32 s5, s9, 1
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_add3_u32 v2, v2, s5, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v9, s4, v2
	v_ashrrev_i32_e32 v3, 31, v2
	v_ashrrev_i32_e32 v10, 31, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	v_lshlrev_b64 v[9:10], 2, v[9:10]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v2, vcc_lo, s22, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s23, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v9, vcc_lo, s22, v9
	v_add_co_ci_u32_e32 v10, vcc_lo, s23, v10, vcc_lo
	s_clause 0x1
	global_load_b32 v3, v[2:3], off
	global_load_b32 v12, v[9:10], off
	v_cvt_f32_u32_e32 v2, v11
	s_waitcnt vmcnt(2)
	v_fma_f32 v5, s37, v1, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_fma_f32 v6, s38, v2, v6
	v_fmac_f32_e32 v8, s38, v2
	v_fma_f32 v1, s37, v1, v7
	v_cvt_i32_f32_e32 v2, v5
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cvt_i32_f32_e32 v5, v6
	v_cvt_i32_f32_e32 v6, v8
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_i32_f32_e32 v1, v1
	v_sub_nc_u32_e32 v11, v6, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_4)
	v_sub_nc_u32_e32 v13, v1, v2
	v_cvt_f64_i32_e32 v[1:2], v2
	v_cvt_f64_i32_e32 v[5:6], v5
	v_cvt_f64_i32_e32 v[7:8], v11
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_f64_i32_e32 v[9:10], v13
	v_fma_f64 v[5:6], v[7:8], 0.5, v[5:6]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f64 v[1:2], v[9:10], 0.5, v[1:2]
	v_cvt_f32_f64_e32 v5, v[5:6]
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_cvt_f32_f64_e32 v1, v[1:2]
	v_cvt_f32_i32_e32 v2, v11
	v_cvt_f32_i32_e32 v6, v13
	s_waitcnt vmcnt(1)
	v_fmac_f32_e32 v5, v3, v2
	s_waitcnt vmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmac_f32 v1, v12, v6 :: v_dual_mul_f32 v2, s1, v5
	v_mul_f32_e32 v1, s39, v1
.LBB0_18:                               ; %Flow439
	s_and_not1_b32 vcc_lo, exec_lo, s15
	s_cbranch_vccnz .LBB0_20
; %bb.19:                               ; %.sink.split
	s_lshl_b32 s1, s7, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s1, s1, s3
	s_add_i32 s1, s1, 6
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[5:6], null, v4, s1, s[0:1]
	v_dual_mov_b32 v7, 0 :: v_dual_add_nc_u32 v6, 6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[8:9], 2, v[6:7]
	v_add_nc_u32_e32 v6, 7, v5
	v_lshlrev_b64 v[5:6], 2, v[6:7]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v7, vcc_lo, s26, v8
	v_add_co_ci_u32_e32 v8, vcc_lo, s27, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s26, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s27, v6, vcc_lo
	s_clause 0x1
	global_store_b32 v[7:8], v2, off
	global_store_b32 v[5:6], v1, off
.LBB0_20:
	s_cmp_ge_u32 s14, s36
	s_cselect_b32 s1, -1, 0
	s_cmp_lg_u32 s6, 0
	s_cselect_b32 s0, -1, 0
	s_cmp_eq_u32 s6, 0
	s_cselect_b32 s5, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s1, s5, s1
	s_and_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccz .LBB0_23
; %bb.21:
	s_cmp_lt_u32 s14, s8
	s_cselect_b32 s1, -1, 0
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, s0, s1
	s_and_b32 vcc_lo, exec_lo, s0
	s_cbranch_vccz .LBB0_24
; %bb.22:
	s_lshl_b32 s0, s7, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s1, s0, s3
	s_add_i32 s1, s1, 6
	s_delay_alu instid0(SALU_CYCLE_1)
	v_mul_lo_u32 v1, v4, s1
	s_mul_i32 s1, s14, 3
	s_delay_alu instid0(VALU_DEP_1) | instid1(SALU_CYCLE_1)
	v_add3_u32 v3, s0, s1, v1
	v_mov_b32_e32 v2, 0
	s_mov_b32 s1, 0
	s_or_b32 s0, s12, exec_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v1, 6, v3
	v_lshlrev_b64 v[5:6], 2, v[1:2]
	v_add_nc_u32_e32 v1, 7, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[7:8], 2, v[1:2]
	v_add_co_u32 v5, vcc_lo, s26, v5
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_4)
	v_add_co_ci_u32_e32 v6, vcc_lo, s27, v6, vcc_lo
	v_add_nc_u32_e32 v1, 8, v3
	v_add_co_u32 v7, vcc_lo, s26, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s27, v8, vcc_lo
	s_clause 0x1
	global_store_b32 v[5:6], v2, off
	global_store_b32 v[7:8], v2, off
	s_branch .LBB0_25
.LBB0_23:
	s_mov_b32 s5, -1
	s_mov_b32 s0, s12
                                        ; implicit-def: $sgpr1
                                        ; implicit-def: $vgpr1
	s_branch .LBB0_26
.LBB0_24:
	s_mov_b32 s0, s12
                                        ; implicit-def: $sgpr1
                                        ; implicit-def: $vgpr1
.LBB0_25:                               ; %Flow444
	s_mov_b32 s5, 0
.LBB0_26:                               ; %Flow443
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 vcc_lo, exec_lo, s5
	s_cbranch_vccz .LBB0_28
; %bb.27:
	v_dual_mov_b32 v1, 0 :: v_dual_add_nc_u32 v2, 1, v0
	s_mul_i32 s1, s36, 3
	s_mul_i32 s14, s14, 3
	s_or_b32 s0, s0, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[5:6], 2, v[0:1]
	v_dual_mov_b32 v3, v1 :: v_dual_add_nc_u32 v0, 2, 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[2:3], 2, v[2:3]
	v_add_co_u32 v7, vcc_lo, s18, v7
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v8, vcc_lo, s19, v8, vcc_lo
	v_add_co_u32 v5, vcc_lo, s18, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s19, v6, vcc_lo
	global_load_b32 v0, v[7:8], off
	v_add_co_u32 v2, vcc_lo, s18, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s19, v3, vcc_lo
	s_clause 0x1
	global_load_b32 v5, v[5:6], off
	global_load_b32 v2, v[2:3], off
	s_waitcnt vmcnt(2)
	v_cvt_u32_f32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_lo_u32 v3, s1, v0
	s_waitcnt vmcnt(1)
	v_cvt_u32_f32_e32 v0, v5
	s_waitcnt vmcnt(0)
	v_cvt_u32_f32_e32 v5, v2
	s_mul_i32 s1, s2, s29
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_mul_i32 s1, s1, s6
	s_mul_i32 s1, s1, 3
	v_add3_u32 v6, s14, s10, v3
	v_mad_u64_u32 v[2:3], null, v5, s30, v[0:1]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v0, s4, v6
	v_add3_u32 v2, v2, s1, v0
	s_lshl_b32 s1, s7, 1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_add_i32 s2, s1, s3
	v_ashrrev_i32_e32 v3, 31, v2
	v_add_nc_u32_e32 v5, s4, v2
	s_add_i32 s2, s2, 6
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_lo_u32 v0, v4, s2
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s24, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s25, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_add3_u32 v11, s1, s14, v0
	global_load_b32 v9, v[2:3], off
	s_waitcnt vmcnt(0)
	v_mul_f32_e32 v4, 0x3fb8aa3b, v9
	v_add_nc_u32_e32 v7, s4, v5
	v_ashrrev_i32_e32 v6, 31, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v8, 31, v7
	v_lshlrev_b64 v[5:6], 2, v[5:6]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[7:8], 2, v[7:8]
	v_add_co_u32 v2, vcc_lo, s24, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v3, vcc_lo, s25, v6, vcc_lo
	v_add_co_u32 v5, vcc_lo, s24, v7
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v6, vcc_lo, s25, v8, vcc_lo
	s_clause 0x1
	global_load_b32 v8, v[2:3], off
	global_load_b32 v10, v[5:6], off
	v_fma_f32 v5, 0x3fb8aa3b, v9, -v4
	v_add_nc_u32_e32 v0, 6, v11
	v_rndne_f32_e32 v7, v4
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_fmac_f32_e32 v5, 0x32a5705f, v9
	v_lshlrev_b64 v[2:3], 2, v[0:1]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_sub_f32_e32 v4, v4, v7
	v_cvt_i32_f32_e32 v7, v7
	v_add_f32_e32 v4, v4, v5
	s_waitcnt vmcnt(1)
	v_mul_f32_e32 v6, 0x3fb8aa3b, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fma_f32 v12, 0x3fb8aa3b, v8, -v6
	v_rndne_f32_e32 v13, v6
	v_fmac_f32_e32 v12, 0x32a5705f, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_sub_f32_e32 v6, v6, v13
	s_waitcnt vmcnt(0)
	v_dual_mul_f32 v0, 0x3fb8aa3b, v10 :: v_dual_add_f32 v5, v6, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_rndne_f32_e32 v15, v0
	v_fma_f32 v14, 0x3fb8aa3b, v10, -v0
	v_exp_f32_e32 v12, v4
	v_add_co_u32 v4, vcc_lo, s26, v2
	v_sub_f32_e32 v16, v0, v15
	v_add_nc_u32_e32 v0, 7, v11
	v_fmac_f32_e32 v14, 0x32a5705f, v10
	v_cvt_i32_f32_e32 v2, v13
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	v_add_f32_e32 v6, v16, v14
	v_exp_f32_e32 v14, v5
	v_add_co_ci_u32_e32 v5, vcc_lo, s27, v3, vcc_lo
	v_ldexp_f32 v12, v12, v7
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_4) | instid1(TRANS32_DEP_2)
	v_exp_f32_e32 v16, v6
	v_add_co_u32 v6, vcc_lo, s26, v0
	v_add_co_ci_u32_e32 v7, vcc_lo, s27, v1, vcc_lo
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v9
	v_cvt_i32_f32_e32 v3, v15
	v_ldexp_f32 v0, v14, v2
	v_cndmask_b32_e32 v2, 0, v12, vcc_lo
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v8
	s_delay_alu instid0(TRANS32_DEP_1) | instid1(VALU_DEP_4)
	v_ldexp_f32 v1, v16, v3
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v0, 0, v0, vcc_lo
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v10
	v_cndmask_b32_e32 v1, 0, v1, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v9
	v_cndmask_b32_e32 v2, 0x7f800000, v2, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v8
	v_cndmask_b32_e32 v0, 0x7f800000, v0, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v10
	s_clause 0x1
	global_store_b32 v[4:5], v2, off
	global_store_b32 v[6:7], v0, off
	v_cndmask_b32_e32 v3, 0x7f800000, v1, vcc_lo
	v_add_nc_u32_e32 v1, 8, v11
	s_branch .LBB0_29
.LBB0_28:
	v_mov_b32_e32 v3, s1
.LBB0_29:                               ; %Flow445
	s_and_not1_b32 s1, s12, exec_lo
	s_and_b32 s0, s0, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s12, s1, s0
.LBB0_30:                               ; %Flow442
	s_or_b32 exec_lo, exec_lo, s13
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_b32 exec_lo, exec_lo, s12
	s_cbranch_execz .LBB0_32
; %bb.31:                               ; %.sink.split355
	v_mov_b32_e32 v2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	v_add_co_u32 v0, vcc_lo, s26, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s27, v1, vcc_lo
	global_store_b32 v[0:1], v3, off
.LBB0_32:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 392
		.amdhsa_user_sgpr_count 13
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 17
		.amdhsa_next_free_sgpr 40
		.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	_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf, .Lfunc_end0-_Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 3120
; NumSgprs: 42
; NumVgprs: 17
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 5
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 42
; NumVGPRsForWavesPerEU: 17
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 13
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .offset:         8
        .size:           4
        .value_kind:     by_value
      - .offset:         12
        .size:           4
        .value_kind:     by_value
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .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
      - .offset:         56
        .size:           4
        .value_kind:     by_value
      - .offset:         60
        .size:           4
        .value_kind:     by_value
      - .offset:         64
        .size:           4
        .value_kind:     by_value
      - .offset:         68
        .size:           4
        .value_kind:     by_value
      - .offset:         72
        .size:           4
        .value_kind:     by_value
      - .offset:         76
        .size:           4
        .value_kind:     by_value
      - .offset:         80
        .size:           4
        .value_kind:     by_value
      - .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
      - .offset:         136
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         140
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         144
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         148
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         150
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         152
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         154
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         156
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         158
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         176
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         184
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         192
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         200
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 392
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf
    .private_segment_fixed_size: 0
    .sgpr_count:     42
    .sgpr_spill_count: 0
    .symbol:         _Z25cudaS_ssdToOutput_kernelsjjjjjjPjjjjjjjjjjffffPKfS1_S1_S1_S1_Pf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     17
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
