//##############################################################################################################################################################################################################//
//Aquila - An Open-Source GPU-Accelerated Toolkit for Cognitive and Neuro-Robotics Research																														//
//																																																				//
//Copyright (c) <2012>, <Martin Peniak - www.martinpeniak.com>																																					//
//All rights reserved.																																															//
//																																																				//
//Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:																//
//																																																				//
// - Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.																				//
// - Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution.	//
//																																																				//
//THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR	//
//A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT	//
//LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR	//
//TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.																//
//                                                                                                                                                                                                              //
//The views and conclusions contained in the software and documentation are those of the authors and should not be interpreted                                                                                  //
//as representing official policies,either expressed or implied, of the FreeBSD Project.                                                                                                                        //
//##############################################################################################################################################################################################################//

#include <cuda.h>
#include <cuda_runtime.h>

/*!
 * \brief Calculates the euclidean distance between input vector and weight vector of a neuron and saves these distances in the output array.
 * \param[in] inputs - inputs
 * \param[in] weights - weights
 * \param[in] numInputs - number of inputs
 * \param[in] sequenceId - sequence id
 * \param[in] numOutputs - number of outputs
 * \param[out] outputs - outputs
 */
__global__ void propogateInputKernel(float *inputs, float *weights, float *outputs, int numInputs, int sequenceId, int numOutputs)
{
    //get the thread id
    int idx = blockIdx.x*blockDim.x + threadIdx.x;

	if(idx<numOutputs)
	{
		//euclidean distance
		float distance = 0.0;

		//calculate the euclidean distance between each node's weight vector and the current input vector
		for(int i=0; i<numInputs; i++)
		{
			distance += (inputs[(sequenceId * numInputs) + i] - weights[idx+(numOutputs*i)]) * (inputs[(sequenceId * numInputs) + i] - weights[idx+(numOutputs*i)]);
		}

		distance = sqrtf(distance);

		//wait until all threads are finished
		__syncthreads();

		//assign the euclidean distance as the output having a particular index
		outputs[idx] = distance;
	}
}

/*!
 * \brief Finds the best matching unit having the lowest euclidean distance - part 1/2.
 * \param[in] outputs - outputs
 * \param[in] numOutputs - number of outputs
 * \param[out] winner - winner
 */
__global__ void findBestMatchPass1Kernel(float *outputs,  int *winner, int numOutputs)
{
    //get local thread id
    int tid  = threadIdx.x;

    //get global thread id
    int idx = blockIdx.x*blockDim.x + threadIdx.x;

    if(idx<numOutputs)
    {
    	//initialise
        winner[idx] = idx;
    	
        //synchronise threads to make sure all data was loaded
        __syncthreads();

        for(int s=1; s<blockDim.x; s*=2)
    	{
        	int index = 2 * s * tid;

            if(index<blockDim.x && (blockIdx.x*blockDim.x)+index+s<numOutputs)
        	{
                if(outputs[winner[(blockIdx.x*blockDim.x)+index]] > outputs[winner[(blockIdx.x*blockDim.x)+index+s]])
                {
                    winner[(blockIdx.x*blockDim.x)+index] = winner[(blockIdx.x*blockDim.x)+index+s];
                }
        	}

        	//wait until all threads are finished with current pass
        	__syncthreads();
    	}
    }
}

/*!
 * \brief Finds the best matching unit having the lowest euclidean distance - part 2/2.
 * \param[in] outputs - outputs
 * \param[in] numOutputs - number of outputs
 * \param[out] winner - winner
 */
__global__ void findBestMatchPass2Kernel(float *outputs, int *winner, int numOutputs)
{
    for(int s=1; s<blockDim.x; s*=2)
    {
        int index = 2*s*threadIdx.x;

        if(index<blockDim.x && 2*(s*(threadIdx.x*blockDim.x))+(s*blockDim.x)<numOutputs)
        {
            if(outputs[winner[2*(s*(threadIdx.x*blockDim.x))]] > outputs[winner[2*(s*(threadIdx.x*blockDim.x))+(s*blockDim.x)]])
			{
				winner[2*(s*(threadIdx.x*blockDim.x))] = winner[2*(s*(threadIdx.x*blockDim.x))+(s*blockDim.x)];
			}
        }

        //wait until all threads are finished with current pass
        __syncthreads();
    }

    if(threadIdx.x==0)
    {
        //find out X and Y coordinates and save them in activity_index[1] and activity_index[2] respectively
        int Y = -1;
        int X = int(winner[0]);

        int dim = int(sqrtf((float)numOutputs));

        while(X>-1)
        {
            X -= dim;
            Y ++;
        }

        //assign X,Y possitions
        winner[1] = X + dim;
        winner[2] = Y;
    }
    __syncthreads();
}

/*!
 * \brief Updates weights of the self-organising map.
 * \param[in] inputs - inputs
 * \param[in] winner - winner
 * \param[in] sigma - sigma
 * \param[in] numInputs - number of inputs
 * \param[in] sequenceId - sequence id
 * \param[in] numOutputs - number of outputs
 * \param[in] neighbourhoodSize - neighbourhood size
 * \param[in] numIterations - number of iterations
 * \param[in] currentIteration - current iteration
 * \param[out] weights - weights
 */
__global__ void updateWeightsKernel(float *inputs, float *weights, int *winner, float sigma, int numInputs, int sequenceId, int numOutputs, int neighbourhoodSize, float initLearningRate,int numIterations, int currentIteration)
{
    int idx = blockIdx.x*blockDim.x + threadIdx.x;

	if(idx<numOutputs)
	{
		//dimension of the map
		int dim = int(sqrtf((float)numOutputs));

		//euclidean distance
		float distance = 0.0;

		//current learning rate
		float currentLearningRate = 0.0;

		//convert current index to x,y
		int curY = -1;
		int curX = idx;
		do{
			curX -= dim;
			curY ++;
        }while(curX>-1);
		curX += dim;

		//calculate the distance from winner to current in the matrix
        distance = (float)((curX-winner[1])*(curX-winner[1]))+((curY-winner[2])*(curY-winner[2]));

        if(distance>(neighbourhoodSize*neighbourhoodSize))
		{
			distance = 0.0;
		}
		else
		{
            if(sigma>0.0)
			{
                float s = sigma*expf(-(float)currentIteration/numIterations);
                distance = expf(-(distance)/(2*(s*s)));
			}
			else
			{
                distance = expf(-(distance)/(2*dim));
			}
		}

		//adjust the learning rate
        currentLearningRate = initLearningRate*__expf(-(float)currentIteration/numIterations);

		//update weights
		for(int i=0; i<numInputs; i++) 
		{
            weights[idx+(numOutputs*i)] += (inputs[(sequenceId*numInputs)+i]-weights[idx+(numOutputs*i)])*distance*currentLearningRate;
		}	
	}
}

/*!
 * \brief Wrapper for findBestMatchPass1Kernel and findBestMatchPass2Kernel.
 * \param[in] grid - CUDA grid size
 * \param[in] block - CUDA block size
 * \param[in] outputs - outputs
 * \param[in] numOutputs - number of outputs
 * \param[out] winner - winner
 */
void findBestMatchOnDevice(dim3 grid, dim3 block, float *outputs,  int *winner, int numOutputs)
{
    findBestMatchPass1Kernel<<<grid,block>>>(outputs, winner, numOutputs);
    findBestMatchPass2Kernel<<<1,block>>>(outputs, winner, numOutputs);
}

/*!
 * \brief Wrapper for propogateInputKernel.
 * \param[in] grid - CUDA grid size
 * \param[in] block - CUDA block size
 * \param[in] inputs - inputs
 * \param[in] weights - weights
 * \param[in] numInputs - number of inputs
 * \param[in] sequenceId - sequence id
 * \param[in] numOutputs - number of outputs
 * \param[out] outputs - outputs
 */
void propogateInputOnDevice(dim3 grid, dim3 block, float *inputs, float *weights, float *outputs, int numInputs, int sequenceId, int numOutputs)
{
    propogateInputKernel<<<grid,block>>>(inputs, weights, outputs, numInputs, sequenceId, numOutputs);
}

/*!
 * \brief Wrapper for updateWeightsKernel.
 * \param[in] grid - CUDA grid size
 * \param[in] block - CUDA block size
 * \param[in] inputs - inputs
 * \param[in] winner - winner
 * \param[in] sigma - sigma
 * \param[in] numInputs - number of inputs
 * \param[in] sequenceId - sequence id
 * \param[in] numOutputs - number of outputs
 * \param[in] neighbourhoodSize - neighbourhood size
 * \param[in] numIterations - number of iterations
 * \param[in] currentIteration - current iteration
 * \param[out] weights - weights
 */
void updateWeightsOnDevice(dim3 grid, dim3 block, float *inputs, float *weights, int *winner, float sigma, int numInputs, int sequenceId, int numOutputs, int neighbourhoodSize, float initLearningRate,int numIterations, int currentIteration)
{
    updateWeightsKernel<<<grid,block>>>(inputs, weights, winner, sigma, numInputs, sequenceId, numOutputs, neighbourhoodSize, initLearningRate, numIterations, currentIteration);
}

	.file	"tmpxft_0037341a_00000000-6_somKernels.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2013:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2013:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z47__device_stub__Z20propogateInputKernelPfS_S_iiiPfS_S_iii
	.type	_Z47__device_stub__Z20propogateInputKernelPfS_S_iiiPfS_S_iii, @function
_Z47__device_stub__Z20propogateInputKernelPfS_S_iiiPfS_S_iii:
.LFB2035:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 16(%rsp)
	movl	%r9d, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z20propogateInputKernelPfS_S_iii(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L2:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2035:
	.size	_Z47__device_stub__Z20propogateInputKernelPfS_S_iiiPfS_S_iii, .-_Z47__device_stub__Z20propogateInputKernelPfS_S_iiiPfS_S_iii
	.globl	_Z20propogateInputKernelPfS_S_iii
	.type	_Z20propogateInputKernelPfS_S_iii, @function
_Z20propogateInputKernelPfS_S_iii:
.LFB2036:
	.cfi_startproc
	endbr64
	jmp	_Z47__device_stub__Z20propogateInputKernelPfS_S_iiiPfS_S_iii
	.cfi_endproc
.LFE2036:
	.size	_Z20propogateInputKernelPfS_S_iii, .-_Z20propogateInputKernelPfS_S_iii
	.globl	_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii
	.type	_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii, @function
_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii:
.LFB2009:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	%ecx, %ecx
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%r9, %rbp
	xorl	%r9d, %r9d
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%r8, %rbx
	xorl	%r8d, %r8d
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	96(%rsp), %r12
	movl	104(%rsp), %r13d
	movl	112(%rsp), %r14d
	movl	120(%rsp), %r15d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L8
	addq	$40, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 56
	movl	%r15d, %r9d
	movl	%r14d, %r8d
	movl	%r13d, %ecx
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	jmp	_Z47__device_stub__Z20propogateInputKernelPfS_S_iiiPfS_S_iii
.L8:
	.cfi_restore_state
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	ret
	.cfi_endproc
.LFE2009:
	.size	_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii, .-_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii
	.globl	_Z47__device_stub__Z24findBestMatchPass1KernelPfPiiPfPii
	.type	_Z47__device_stub__Z24findBestMatchPass1KernelPfPiiPfPii, @function
_Z47__device_stub__Z24findBestMatchPass1KernelPfPiiPfPii:
.LFB2037:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L11
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z24findBestMatchPass1KernelPfPii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L11:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2037:
	.size	_Z47__device_stub__Z24findBestMatchPass1KernelPfPiiPfPii, .-_Z47__device_stub__Z24findBestMatchPass1KernelPfPiiPfPii
	.globl	_Z24findBestMatchPass1KernelPfPii
	.type	_Z24findBestMatchPass1KernelPfPii, @function
_Z24findBestMatchPass1KernelPfPii:
.LFB2038:
	.cfi_startproc
	endbr64
	jmp	_Z47__device_stub__Z24findBestMatchPass1KernelPfPiiPfPii
	.cfi_endproc
.LFE2038:
	.size	_Z24findBestMatchPass1KernelPfPii, .-_Z24findBestMatchPass1KernelPfPii
	.globl	_Z47__device_stub__Z24findBestMatchPass2KernelPfPiiPfPii
	.type	_Z47__device_stub__Z24findBestMatchPass2KernelPfPiiPfPii, @function
_Z47__device_stub__Z24findBestMatchPass2KernelPfPiiPfPii:
.LFB2039:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L16
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z24findBestMatchPass2KernelPfPii(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L16:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L18
	call	__stack_chk_fail@PLT
.L18:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2039:
	.size	_Z47__device_stub__Z24findBestMatchPass2KernelPfPiiPfPii, .-_Z47__device_stub__Z24findBestMatchPass2KernelPfPiiPfPii
	.globl	_Z24findBestMatchPass2KernelPfPii
	.type	_Z24findBestMatchPass2KernelPfPii, @function
_Z24findBestMatchPass2KernelPfPii:
.LFB2040:
	.cfi_startproc
	endbr64
	jmp	_Z47__device_stub__Z24findBestMatchPass2KernelPfPiiPfPii
	.cfi_endproc
.LFE2040:
	.size	_Z24findBestMatchPass2KernelPfPii, .-_Z24findBestMatchPass2KernelPfPii
	.globl	_Z21findBestMatchOnDevice4dim3S_PfPii
	.type	_Z21findBestMatchOnDevice4dim3S_PfPii, @function
_Z21findBestMatchOnDevice4dim3S_PfPii:
.LFB2008:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%r9, %rbp
	xorl	%r9d, %r9d
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%r8, %rbx
	xorl	%r8d, %r8d
	subq	$48, %rsp
	.cfi_def_cfa_offset 80
	movl	%ecx, 8(%rsp)
	movl	%ecx, %ecx
	movl	80(%rsp), %r12d
	movq	%rdx, (%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L22
	movl	%r12d, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_Z47__device_stub__Z24findBestMatchPass1KernelPfPiiPfPii
.L22:
	movl	8(%rsp), %ecx
	movq	(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	addq	$48, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 32
	movl	%r12d, %edx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	jmp	_Z47__device_stub__Z24findBestMatchPass2KernelPfPiiPfPii
.L21:
	.cfi_restore_state
	addq	$48, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2008:
	.size	_Z21findBestMatchOnDevice4dim3S_PfPii, .-_Z21findBestMatchOnDevice4dim3S_PfPii
	.globl	_Z51__device_stub__Z19updateWeightsKernelPfS_PifiiiifiiPfS_Pifiiiifii
	.type	_Z51__device_stub__Z19updateWeightsKernelPfS_PifiiiifiiPfS_Pifiiiifii, @function
_Z51__device_stub__Z19updateWeightsKernelPfS_PifiiiifiiPfS_Pifiiiifii:
.LFB2041:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 12(%rsp)
	movl	%r9d, 8(%rsp)
	movss	%xmm0, 20(%rsp)
	movss	%xmm1, 4(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	232(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	240(%rsp), %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L25
	pushq	56(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z19updateWeightsKernelPfS_Pifiiiifii(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 240
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L25:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L27
	call	__stack_chk_fail@PLT
.L27:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2041:
	.size	_Z51__device_stub__Z19updateWeightsKernelPfS_PifiiiifiiPfS_Pifiiiifii, .-_Z51__device_stub__Z19updateWeightsKernelPfS_PifiiiifiiPfS_Pifiiiifii
	.globl	_Z19updateWeightsKernelPfS_Pifiiiifii
	.type	_Z19updateWeightsKernelPfS_Pifiiiifii, @function
_Z19updateWeightsKernelPfS_Pifiiiifii:
.LFB2042:
	.cfi_startproc
	endbr64
	jmp	_Z51__device_stub__Z19updateWeightsKernelPfS_PifiiiifiiPfS_Pifiiiifii
	.cfi_endproc
.LFE2042:
	.size	_Z19updateWeightsKernelPfS_Pifiiiifii, .-_Z19updateWeightsKernelPfS_Pifiiiifii
	.globl	_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii
	.type	_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii, @function
_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii:
.LFB2010:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	%ecx, %ecx
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	128(%rsp), %rax
	movq	%r8, (%rsp)
	xorl	%r8d, %r8d
	movq	%r9, 8(%rsp)
	xorl	%r9d, %r9d
	movl	136(%rsp), %r12d
	movss	%xmm0, 28(%rsp)
	movl	144(%rsp), %r13d
	movl	152(%rsp), %r14d
	movss	%xmm1, 24(%rsp)
	movl	160(%rsp), %r15d
	movl	168(%rsp), %ebx
	movq	%rax, 16(%rsp)
	movl	176(%rsp), %ebp
	call	__cudaPushCallConfiguration@PLT
	movss	24(%rsp), %xmm1
	movss	28(%rsp), %xmm0
	testl	%eax, %eax
	jne	.L30
	movl	%ebp, 144(%rsp)
	movl	%r14d, %r9d
	movl	%r13d, %r8d
	movl	%r12d, %ecx
	movq	16(%rsp), %rdx
	movq	(%rsp), %rdi
	movl	%ebx, 136(%rsp)
	movl	%r15d, 128(%rsp)
	movq	8(%rsp), %rsi
	addq	$72, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	jmp	_Z51__device_stub__Z19updateWeightsKernelPfS_PifiiiifiiPfS_Pifiiiifii
.L30:
	.cfi_restore_state
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.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
	ret
	.cfi_endproc
.LFE2010:
	.size	_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii, .-_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z19updateWeightsKernelPfS_Pifiiiifii"
.LC1:
	.string	"_Z24findBestMatchPass2KernelPfPii"
.LC2:
	.string	"_Z24findBestMatchPass1KernelPfPii"
.LC3:
	.string	"_Z20propogateInputKernelPfS_S_iii"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2044:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -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
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z19updateWeightsKernelPfS_Pifiiiifii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z24findBestMatchPass2KernelPfPii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z24findBestMatchPass1KernelPfPii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z20propogateInputKernelPfS_S_iii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	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	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2044:
	.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:

//##############################################################################################################################################################################################################//
//Aquila - An Open-Source GPU-Accelerated Toolkit for Cognitive and Neuro-Robotics Research																														//
//																																																				//
//Copyright (c) <2012>, <Martin Peniak - www.martinpeniak.com>																																					//
//All rights reserved.																																															//
//																																																				//
//Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:																//
//																																																				//
// - Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.																				//
// - Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution.	//
//																																																				//
//THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR	//
//A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT	//
//LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR	//
//TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.																//
//                                                                                                                                                                                                              //
//The views and conclusions contained in the software and documentation are those of the authors and should not be interpreted                                                                                  //
//as representing official policies,either expressed or implied, of the FreeBSD Project.                                                                                                                        //
//##############################################################################################################################################################################################################//

#include <hip/hip_runtime.h>


/*!
 * \brief Calculates the euclidean distance between input vector and weight vector of a neuron and saves these distances in the output array.
 * \param[in] inputs - inputs
 * \param[in] weights - weights
 * \param[in] numInputs - number of inputs
 * \param[in] sequenceId - sequence id
 * \param[in] numOutputs - number of outputs
 * \param[out] outputs - outputs
 */
__global__ void propogateInputKernel(float *inputs, float *weights, float *outputs, int numInputs, int sequenceId, int numOutputs)
{
    //get the thread id
    int idx = blockIdx.x*blockDim.x + threadIdx.x;

	if(idx<numOutputs)
	{
		//euclidean distance
		float distance = 0.0;

		//calculate the euclidean distance between each node's weight vector and the current input vector
		for(int i=0; i<numInputs; i++)
		{
			distance += (inputs[(sequenceId * numInputs) + i] - weights[idx+(numOutputs*i)]) * (inputs[(sequenceId * numInputs) + i] - weights[idx+(numOutputs*i)]);
		}

		distance = sqrtf(distance);

		//wait until all threads are finished
		__syncthreads();

		//assign the euclidean distance as the output having a particular index
		outputs[idx] = distance;
	}
}

/*!
 * \brief Finds the best matching unit having the lowest euclidean distance - part 1/2.
 * \param[in] outputs - outputs
 * \param[in] numOutputs - number of outputs
 * \param[out] winner - winner
 */
__global__ void findBestMatchPass1Kernel(float *outputs,  int *winner, int numOutputs)
{
    //get local thread id
    int tid  = threadIdx.x;

    //get global thread id
    int idx = blockIdx.x*blockDim.x + threadIdx.x;

    if(idx<numOutputs)
    {
    	//initialise
        winner[idx] = idx;
    	
        //synchronise threads to make sure all data was loaded
        __syncthreads();

        for(int s=1; s<blockDim.x; s*=2)
    	{
        	int index = 2 * s * tid;

            if(index<blockDim.x && (blockIdx.x*blockDim.x)+index+s<numOutputs)
        	{
                if(outputs[winner[(blockIdx.x*blockDim.x)+index]] > outputs[winner[(blockIdx.x*blockDim.x)+index+s]])
                {
                    winner[(blockIdx.x*blockDim.x)+index] = winner[(blockIdx.x*blockDim.x)+index+s];
                }
        	}

        	//wait until all threads are finished with current pass
        	__syncthreads();
    	}
    }
}

/*!
 * \brief Finds the best matching unit having the lowest euclidean distance - part 2/2.
 * \param[in] outputs - outputs
 * \param[in] numOutputs - number of outputs
 * \param[out] winner - winner
 */
__global__ void findBestMatchPass2Kernel(float *outputs, int *winner, int numOutputs)
{
    for(int s=1; s<blockDim.x; s*=2)
    {
        int index = 2*s*threadIdx.x;

        if(index<blockDim.x && 2*(s*(threadIdx.x*blockDim.x))+(s*blockDim.x)<numOutputs)
        {
            if(outputs[winner[2*(s*(threadIdx.x*blockDim.x))]] > outputs[winner[2*(s*(threadIdx.x*blockDim.x))+(s*blockDim.x)]])
			{
				winner[2*(s*(threadIdx.x*blockDim.x))] = winner[2*(s*(threadIdx.x*blockDim.x))+(s*blockDim.x)];
			}
        }

        //wait until all threads are finished with current pass
        __syncthreads();
    }

    if(threadIdx.x==0)
    {
        //find out X and Y coordinates and save them in activity_index[1] and activity_index[2] respectively
        int Y = -1;
        int X = int(winner[0]);

        int dim = int(sqrtf((float)numOutputs));

        while(X>-1)
        {
            X -= dim;
            Y ++;
        }

        //assign X,Y possitions
        winner[1] = X + dim;
        winner[2] = Y;
    }
    __syncthreads();
}

/*!
 * \brief Updates weights of the self-organising map.
 * \param[in] inputs - inputs
 * \param[in] winner - winner
 * \param[in] sigma - sigma
 * \param[in] numInputs - number of inputs
 * \param[in] sequenceId - sequence id
 * \param[in] numOutputs - number of outputs
 * \param[in] neighbourhoodSize - neighbourhood size
 * \param[in] numIterations - number of iterations
 * \param[in] currentIteration - current iteration
 * \param[out] weights - weights
 */
__global__ void updateWeightsKernel(float *inputs, float *weights, int *winner, float sigma, int numInputs, int sequenceId, int numOutputs, int neighbourhoodSize, float initLearningRate,int numIterations, int currentIteration)
{
    int idx = blockIdx.x*blockDim.x + threadIdx.x;

	if(idx<numOutputs)
	{
		//dimension of the map
		int dim = int(sqrtf((float)numOutputs));

		//euclidean distance
		float distance = 0.0;

		//current learning rate
		float currentLearningRate = 0.0;

		//convert current index to x,y
		int curY = -1;
		int curX = idx;
		do{
			curX -= dim;
			curY ++;
        }while(curX>-1);
		curX += dim;

		//calculate the distance from winner to current in the matrix
        distance = (float)((curX-winner[1])*(curX-winner[1]))+((curY-winner[2])*(curY-winner[2]));

        if(distance>(neighbourhoodSize*neighbourhoodSize))
		{
			distance = 0.0;
		}
		else
		{
            if(sigma>0.0)
			{
                float s = sigma*expf(-(float)currentIteration/numIterations);
                distance = expf(-(distance)/(2*(s*s)));
			}
			else
			{
                distance = expf(-(distance)/(2*dim));
			}
		}

		//adjust the learning rate
        currentLearningRate = initLearningRate*__expf(-(float)currentIteration/numIterations);

		//update weights
		for(int i=0; i<numInputs; i++) 
		{
            weights[idx+(numOutputs*i)] += (inputs[(sequenceId*numInputs)+i]-weights[idx+(numOutputs*i)])*distance*currentLearningRate;
		}	
	}
}

/*!
 * \brief Wrapper for findBestMatchPass1Kernel and findBestMatchPass2Kernel.
 * \param[in] grid - CUDA grid size
 * \param[in] block - CUDA block size
 * \param[in] outputs - outputs
 * \param[in] numOutputs - number of outputs
 * \param[out] winner - winner
 */
void findBestMatchOnDevice(dim3 grid, dim3 block, float *outputs,  int *winner, int numOutputs)
{
    findBestMatchPass1Kernel<<<grid,block>>>(outputs, winner, numOutputs);
    findBestMatchPass2Kernel<<<1,block>>>(outputs, winner, numOutputs);
}

/*!
 * \brief Wrapper for propogateInputKernel.
 * \param[in] grid - CUDA grid size
 * \param[in] block - CUDA block size
 * \param[in] inputs - inputs
 * \param[in] weights - weights
 * \param[in] numInputs - number of inputs
 * \param[in] sequenceId - sequence id
 * \param[in] numOutputs - number of outputs
 * \param[out] outputs - outputs
 */
void propogateInputOnDevice(dim3 grid, dim3 block, float *inputs, float *weights, float *outputs, int numInputs, int sequenceId, int numOutputs)
{
    propogateInputKernel<<<grid,block>>>(inputs, weights, outputs, numInputs, sequenceId, numOutputs);
}

/*!
 * \brief Wrapper for updateWeightsKernel.
 * \param[in] grid - CUDA grid size
 * \param[in] block - CUDA block size
 * \param[in] inputs - inputs
 * \param[in] winner - winner
 * \param[in] sigma - sigma
 * \param[in] numInputs - number of inputs
 * \param[in] sequenceId - sequence id
 * \param[in] numOutputs - number of outputs
 * \param[in] neighbourhoodSize - neighbourhood size
 * \param[in] numIterations - number of iterations
 * \param[in] currentIteration - current iteration
 * \param[out] weights - weights
 */
void updateWeightsOnDevice(dim3 grid, dim3 block, float *inputs, float *weights, int *winner, float sigma, int numInputs, int sequenceId, int numOutputs, int neighbourhoodSize, float initLearningRate,int numIterations, int currentIteration)
{
    updateWeightsKernel<<<grid,block>>>(inputs, weights, winner, sigma, numInputs, sequenceId, numOutputs, neighbourhoodSize, initLearningRate, numIterations, currentIteration);
}

	.text
	.file	"somKernels.hip"
	.globl	_Z35__device_stub__propogateInputKernelPfS_S_iii # -- Begin function _Z35__device_stub__propogateInputKernelPfS_S_iii
	.type	_Z35__device_stub__propogateInputKernelPfS_S_iii,@function
_Z35__device_stub__propogateInputKernelPfS_S_iii: # @_Z35__device_stub__propogateInputKernelPfS_S_iii
	.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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	20(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	16(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	12(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%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	$_Z20propogateInputKernelPfS_S_iii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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	_Z35__device_stub__propogateInputKernelPfS_S_iii, .Lfunc_end0-_Z35__device_stub__propogateInputKernelPfS_S_iii
	.cfi_endproc
                                        # -- End function
	.globl	_Z39__device_stub__findBestMatchPass1KernelPfPii # -- Begin function _Z39__device_stub__findBestMatchPass1KernelPfPii
	.type	_Z39__device_stub__findBestMatchPass1KernelPfPii,@function
_Z39__device_stub__findBestMatchPass1KernelPfPii: # @_Z39__device_stub__findBestMatchPass1KernelPfPii
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z24findBestMatchPass1KernelPfPii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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_end1:
	.size	_Z39__device_stub__findBestMatchPass1KernelPfPii, .Lfunc_end1-_Z39__device_stub__findBestMatchPass1KernelPfPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z39__device_stub__findBestMatchPass2KernelPfPii # -- Begin function _Z39__device_stub__findBestMatchPass2KernelPfPii
	.type	_Z39__device_stub__findBestMatchPass2KernelPfPii,@function
_Z39__device_stub__findBestMatchPass2KernelPfPii: # @_Z39__device_stub__findBestMatchPass2KernelPfPii
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	12(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z24findBestMatchPass2KernelPfPii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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_end2:
	.size	_Z39__device_stub__findBestMatchPass2KernelPfPii, .Lfunc_end2-_Z39__device_stub__findBestMatchPass2KernelPfPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z34__device_stub__updateWeightsKernelPfS_Pifiiiifii # -- Begin function _Z34__device_stub__updateWeightsKernelPfS_Pifiiiifii
	.type	_Z34__device_stub__updateWeightsKernelPfS_Pifiiiifii,@function
_Z34__device_stub__updateWeightsKernelPfS_Pifiiiifii: # @_Z34__device_stub__updateWeightsKernelPfS_Pifiiiifii
	.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	$192, %rsp
	.cfi_def_cfa_offset 240
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	20(%rsp), %rdx
	movss	%xmm0, (%rdx)
	leaq	16(%rsp), %r10
	movl	%ecx, (%r10)
	leaq	12(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	8(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	4(%rsp), %r9
	movss	%xmm1, (%r9)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%r10, 32(%rbx)
	movq	%rcx, 40(%rbx)
	movq	%r8, 48(%rbx)
	leaq	240(%rsp), %rax
	movq	%rax, 56(%rbx)
	movq	%r9, 64(%rbx)
	leaq	248(%rsp), %rax
	movq	%rax, 72(%rbx)
	leaq	256(%rsp), %rax
	movq	%rax, 80(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%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	$_Z19updateWeightsKernelPfS_Pifiiiifii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$208, %rsp
	.cfi_adjust_cfa_offset -208
	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_end3:
	.size	_Z34__device_stub__updateWeightsKernelPfS_Pifiiiifii, .Lfunc_end3-_Z34__device_stub__updateWeightsKernelPfS_Pifiiiifii
	.cfi_endproc
                                        # -- End function
	.globl	_Z21findBestMatchOnDevice4dim3S_PfPii # -- Begin function _Z21findBestMatchOnDevice4dim3S_PfPii
	.type	_Z21findBestMatchOnDevice4dim3S_PfPii,@function
_Z21findBestMatchOnDevice4dim3S_PfPii:  # @_Z21findBestMatchOnDevice4dim3S_PfPii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %rbx
	movq	%r8, %r14
	movl	%ecx, %r15d
	movq	%rdx, %r12
	movl	48(%rsp), %ebp
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_2
# %bb.1:
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movl	%ebp, %edx
	callq	_Z39__device_stub__findBestMatchPass1KernelPfPii
.LBB4_2:
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movl	$1, %esi
	movq	%r12, %rdx
	movl	%r15d, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB4_4
# %bb.3:
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB4_4:
	.cfi_def_cfa_offset 48
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movl	%ebp, %edx
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	_Z39__device_stub__findBestMatchPass2KernelPfPii # TAILCALL
.Lfunc_end4:
	.size	_Z21findBestMatchOnDevice4dim3S_PfPii, .Lfunc_end4-_Z21findBestMatchOnDevice4dim3S_PfPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii # -- Begin function _Z22propogateInputOnDevice4dim3S_PfS0_S0_iii
	.type	_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii,@function
_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii: # @_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %rbx
	movq	%r8, %r14
	movl	88(%rsp), %ebp
	movl	80(%rsp), %r15d
	movl	72(%rsp), %r12d
	movq	64(%rsp), %r13
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB5_2
# %bb.1:
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB5_2:
	.cfi_def_cfa_offset 64
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movq	%r13, %rdx
	movl	%r12d, %ecx
	movl	%r15d, %r8d
	movl	%ebp, %r9d
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	_Z35__device_stub__propogateInputKernelPfS_S_iii # TAILCALL
.Lfunc_end5:
	.size	_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii, .Lfunc_end5-_Z22propogateInputOnDevice4dim3S_PfS0_S0_iii
	.cfi_endproc
                                        # -- End function
	.globl	_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii # -- Begin function _Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii
	.type	_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii,@function
_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii: # @_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movss	%xmm1, 4(%rsp)                  # 4-byte Spill
	movss	%xmm0, (%rsp)                   # 4-byte Spill
	movq	%r9, %rbx
	movq	%r8, %r14
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	je	.LBB6_1
# %bb.2:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	jmp	.LBB6_3
.LBB6_1:
	.cfi_def_cfa_offset 32
	movl	80(%rsp), %eax
	movl	72(%rsp), %r10d
	movl	64(%rsp), %r11d
	movl	56(%rsp), %r9d
	movl	48(%rsp), %r8d
	movl	40(%rsp), %ecx
	movq	32(%rsp), %rdx
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movss	8(%rsp), %xmm0                  # 4-byte Reload
                                        # xmm0 = mem[0],zero,zero,zero
	movss	12(%rsp), %xmm1                 # 4-byte Reload
                                        # xmm1 = mem[0],zero,zero,zero
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	callq	_Z34__device_stub__updateWeightsKernelPfS_Pifiiiifii
	addq	$40, %rsp
	.cfi_adjust_cfa_offset -40
.LBB6_3:
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii, .Lfunc_end6-_Z21updateWeightsOnDevice4dim3S_PfS0_Pifiiiifii
	.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:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB7_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB7_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z20propogateInputKernelPfS_S_iii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z24findBestMatchPass1KernelPfPii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z24findBestMatchPass2KernelPfPii, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z19updateWeightsKernelPfS_Pifiiiifii, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end7:
	.size	__hip_module_ctor, .Lfunc_end7-__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	.LBB8_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
.LBB8_2:
	retq
.Lfunc_end8:
	.size	__hip_module_dtor, .Lfunc_end8-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z20propogateInputKernelPfS_S_iii,@object # @_Z20propogateInputKernelPfS_S_iii
	.section	.rodata,"a",@progbits
	.globl	_Z20propogateInputKernelPfS_S_iii
	.p2align	3, 0x0
_Z20propogateInputKernelPfS_S_iii:
	.quad	_Z35__device_stub__propogateInputKernelPfS_S_iii
	.size	_Z20propogateInputKernelPfS_S_iii, 8

	.type	_Z24findBestMatchPass1KernelPfPii,@object # @_Z24findBestMatchPass1KernelPfPii
	.globl	_Z24findBestMatchPass1KernelPfPii
	.p2align	3, 0x0
_Z24findBestMatchPass1KernelPfPii:
	.quad	_Z39__device_stub__findBestMatchPass1KernelPfPii
	.size	_Z24findBestMatchPass1KernelPfPii, 8

	.type	_Z24findBestMatchPass2KernelPfPii,@object # @_Z24findBestMatchPass2KernelPfPii
	.globl	_Z24findBestMatchPass2KernelPfPii
	.p2align	3, 0x0
_Z24findBestMatchPass2KernelPfPii:
	.quad	_Z39__device_stub__findBestMatchPass2KernelPfPii
	.size	_Z24findBestMatchPass2KernelPfPii, 8

	.type	_Z19updateWeightsKernelPfS_Pifiiiifii,@object # @_Z19updateWeightsKernelPfS_Pifiiiifii
	.globl	_Z19updateWeightsKernelPfS_Pifiiiifii
	.p2align	3, 0x0
_Z19updateWeightsKernelPfS_Pifiiiifii:
	.quad	_Z34__device_stub__updateWeightsKernelPfS_Pifiiiifii
	.size	_Z19updateWeightsKernelPfS_Pifiiiifii, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z24findBestMatchPass1KernelPfPii"
	.size	.L__unnamed_2, 34

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z24findBestMatchPass2KernelPfPii"
	.size	.L__unnamed_3, 34

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z19updateWeightsKernelPfS_Pifiiiifii"
	.size	.L__unnamed_4, 38

	.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 _Z35__device_stub__propogateInputKernelPfS_S_iii
	.addrsig_sym _Z39__device_stub__findBestMatchPass1KernelPfPii
	.addrsig_sym _Z39__device_stub__findBestMatchPass2KernelPfPii
	.addrsig_sym _Z34__device_stub__updateWeightsKernelPfS_Pifiiiifii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z20propogateInputKernelPfS_S_iii
	.addrsig_sym _Z24findBestMatchPass1KernelPfPii
	.addrsig_sym _Z24findBestMatchPass2KernelPfPii
	.addrsig_sym _Z19updateWeightsKernelPfS_Pifiiiifii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z20propogateInputKernelPfS_S_iii ; -- Begin function _Z20propogateInputKernelPfS_S_iii
	.globl	_Z20propogateInputKernelPfS_S_iii
	.p2align	8
	.type	_Z20propogateInputKernelPfS_S_iii,@function
_Z20propogateInputKernelPfS_S_iii:      ; @_Z20propogateInputKernelPfS_S_iii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x34
	s_load_b128 s[4:7], s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s6, v1
	s_cbranch_execz .LBB0_6
; %bb.1:                                ; %.preheader
	s_clause 0x1
	s_load_b128 s[8:11], s[0:1], 0x0
	s_load_b64 s[2:3], s[0:1], 0x10
	s_cmp_lt_i32 s4, 1
	s_cbranch_scc1 .LBB0_4
; %bb.2:                                ; %.lr.ph
	s_mul_i32 s0, s5, s4
	v_mov_b32_e32 v0, 0
	s_ashr_i32 s1, s0, 31
	v_mov_b32_e32 v2, v1
	s_lshl_b64 s[0:1], s[0:1], 2
	s_waitcnt lgkmcnt(0)
	s_add_u32 s0, s8, s0
	s_addc_u32 s1, s9, s1
.LBB0_3:                                ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v3, 31, v2
	s_load_b32 s5, s[0:1], 0x0
	s_add_i32 s4, s4, -1
	s_add_u32 s0, s0, 4
	s_addc_u32 s1, s1, 0
	v_lshlrev_b64 v[3:4], 2, v[2:3]
	s_cmp_eq_u32 s4, 0
	v_add_nc_u32_e32 v2, s6, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, s10, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s11, v4, vcc_lo
	global_load_b32 v3, v[3:4], off
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_sub_f32_e32 v3, s5, v3
	s_delay_alu instid0(VALU_DEP_1)
	v_fmac_f32_e32 v0, v3, v3
	s_cbranch_scc0 .LBB0_3
	s_branch .LBB0_5
.LBB0_4:
	v_mov_b32_e32 v0, 0
.LBB0_5:                                ; %Flow55
	s_delay_alu instid0(VALU_DEP_1)
	v_mul_f32_e32 v2, 0x4f800000, v0
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_cndmask_b32_e32 v3, v0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v0, v3
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v2, -1, v0
	v_add_nc_u32_e32 v4, 1, v0
	v_fma_f32 v5, -v2, v0, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v6, -v4, v0, v3
	v_cmp_ge_f32_e64 s0, 0, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v0, v0, v2, s0
	v_cmp_lt_f32_e64 s0, 0, v6
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v0, v0, v4, s0
	v_mul_f32_e32 v4, 0x37800000, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cndmask_b32_e32 v4, v0, v4, vcc_lo
	v_lshlrev_b64 v[0:1], 2, v[1:2]
	v_cmp_class_f32_e64 vcc_lo, v3, 0x260
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v2, v4, v3, vcc_lo
	v_add_co_u32 v0, vcc_lo, s2, v0
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_store_b32 v[0:1], v2, off
.LBB0_6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z20propogateInputKernelPfS_S_iii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 7
		.amdhsa_next_free_sgpr 16
		.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	_Z20propogateInputKernelPfS_S_iii, .Lfunc_end0-_Z20propogateInputKernelPfS_S_iii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 416
; NumSgprs: 18
; NumVgprs: 7
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 7
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z24findBestMatchPass1KernelPfPii ; -- Begin function _Z24findBestMatchPass1KernelPfPii
	.globl	_Z24findBestMatchPass1KernelPfPii
	.p2align	8
	.type	_Z24findBestMatchPass1KernelPfPii,@function
_Z24findBestMatchPass1KernelPfPii:      ; @_Z24findBestMatchPass1KernelPfPii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x24
	s_load_b32 s4, s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s5, s2, 0xffff
	s_mov_b32 s2, exec_lo
	s_mul_i32 s15, s15, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v1, s15, v0
	v_cmpx_gt_i32_e64 s4, v1
	s_cbranch_execz .LBB1_8
; %bb.1:
	s_load_b128 s[0:3], s[0:1], 0x0
	v_ashrrev_i32_e32 v2, 31, v1
	s_cmp_lt_u32 s5, 2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s2, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s3, v3, vcc_lo
	global_store_b32 v[2:3], v1, off
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB1_8
; %bb.2:                                ; %.lr.ph.preheader
	v_mov_b32_e32 v2, 0
	s_mov_b32 s6, 1
.LBB1_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	s_mov_b32 s8, s6
	s_lshl_b32 s6, s6, 1
	s_mov_b32 s7, exec_lo
	v_mul_lo_u32 v1, s6, v0
	v_cmpx_gt_u32_e64 s5, v1
	s_cbranch_execz .LBB1_7
; %bb.4:                                ;   in Loop: Header=BB1_3 Depth=1
	v_add_nc_u32_e32 v1, s15, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, s8, v1
	v_cmp_gt_u32_e32 vcc_lo, s4, v3
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB1_7
; %bb.5:                                ;   in Loop: Header=BB1_3 Depth=1
	v_mov_b32_e32 v4, v2
	v_lshlrev_b64 v[5:6], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[7:8], 2, v[3:4]
	v_add_co_u32 v3, vcc_lo, s2, v5
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v4, vcc_lo, s3, v6, vcc_lo
	v_add_co_u32 v5, vcc_lo, s2, v7
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_ci_u32_e32 v6, vcc_lo, s3, v8, vcc_lo
	s_clause 0x1
	global_load_b32 v7, v[3:4], off
	global_load_b32 v5, v[5:6], off
	s_waitcnt vmcnt(1)
	v_ashrrev_i32_e32 v8, 31, v7
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v6, 31, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[7:8], 2, v[7:8]
	v_lshlrev_b64 v[9:10], 2, v[5:6]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v6, vcc_lo, s0, v7
	v_add_co_ci_u32_e32 v7, vcc_lo, s1, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v8, vcc_lo, s0, v9
	v_add_co_ci_u32_e32 v9, vcc_lo, s1, v10, vcc_lo
	s_clause 0x1
	global_load_b32 v1, v[6:7], off
	global_load_b32 v6, v[8:9], off
	s_waitcnt vmcnt(0)
	v_cmp_gt_f32_e32 vcc_lo, v1, v6
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB1_7
; %bb.6:                                ;   in Loop: Header=BB1_3 Depth=1
	global_store_b32 v[3:4], v5, off
.LBB1_7:                                ;   in Loop: Header=BB1_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s7
	s_cmp_lt_u32 s6, s5
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB1_3
.LBB1_8:                                ; %.loopexit
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z24findBestMatchPass1KernelPfPii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 11
		.amdhsa_next_free_sgpr 16
		.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_end1:
	.size	_Z24findBestMatchPass1KernelPfPii, .Lfunc_end1-_Z24findBestMatchPass1KernelPfPii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 424
; NumSgprs: 18
; NumVgprs: 11
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 11
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z24findBestMatchPass2KernelPfPii ; -- Begin function _Z24findBestMatchPass2KernelPfPii
	.globl	_Z24findBestMatchPass2KernelPfPii
	.p2align	8
	.type	_Z24findBestMatchPass2KernelPfPii,@function
_Z24findBestMatchPass2KernelPfPii:      ; @_Z24findBestMatchPass2KernelPfPii
; %bb.0:
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x24
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b32 s0, s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s1, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lt_u32 s1, 2
	s_cbranch_scc1 .LBB2_7
; %bb.1:                                ; %.lr.ph
	v_dual_mov_b32 v2, 0 :: v_dual_lshlrev_b32 v7, 1, v0
	s_mov_b32 s2, 1
.LBB2_2:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_3) | instid1(VALU_DEP_1)
	s_mov_b32 s8, s2
	s_lshl_b32 s2, s2, 1
	s_mov_b32 s3, exec_lo
	v_mul_lo_u32 v1, s2, v0
	v_cmpx_gt_u32_e64 s1, v1
	s_cbranch_execz .LBB2_6
; %bb.3:                                ;   in Loop: Header=BB2_2 Depth=1
	v_mul_lo_u32 v1, v7, s8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, s8, v1
	v_mul_lo_u32 v3, v3, s1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_gt_u32_e32 vcc_lo, s0, v3
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB2_6
; %bb.4:                                ;   in Loop: Header=BB2_2 Depth=1
	v_mul_lo_u32 v1, v1, s1
	v_mov_b32_e32 v4, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[8:9], 2, v[3:4]
	v_lshlrev_b64 v[5:6], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, s6, v5
	v_add_co_ci_u32_e32 v4, vcc_lo, s7, v6, vcc_lo
	s_delay_alu instid0(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s6, v8
	v_add_co_ci_u32_e32 v6, vcc_lo, s7, v9, vcc_lo
	s_clause 0x1
	global_load_b32 v8, v[3:4], off
	global_load_b32 v5, v[5:6], off
	s_waitcnt vmcnt(1)
	v_ashrrev_i32_e32 v9, 31, v8
	s_waitcnt vmcnt(0)
	v_ashrrev_i32_e32 v6, 31, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[8:9], 2, v[8:9]
	v_lshlrev_b64 v[10:11], 2, v[5:6]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v8, vcc_lo, s4, v8
	v_add_co_ci_u32_e32 v9, vcc_lo, s5, v9, vcc_lo
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v10, vcc_lo, s4, v10
	v_add_co_ci_u32_e32 v11, vcc_lo, s5, v11, vcc_lo
	s_clause 0x1
	global_load_b32 v1, v[8:9], off
	global_load_b32 v6, v[10:11], off
	s_waitcnt vmcnt(0)
	v_cmp_gt_f32_e32 vcc_lo, v1, v6
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB2_6
; %bb.5:                                ;   in Loop: Header=BB2_2 Depth=1
	global_store_b32 v[3:4], v5, off
.LBB2_6:                                ;   in Loop: Header=BB2_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	s_cmp_ge_u32 s2, s1
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB2_2
.LBB2_7:                                ; %._crit_edge
	s_mov_b32 s1, exec_lo
	v_cmpx_eq_u32_e32 0, v0
	s_cbranch_execz .LBB2_11
; %bb.8:
	v_cvt_f32_i32_e32 v1, s0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v2, 0x4f800000, v1
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v1
	v_dual_mov_b32 v0, 0 :: v_dual_cndmask_b32 v1, v1, v2
	global_load_b32 v0, v0, s[6:7]
	v_sqrt_f32_e32 v2, v1
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v3, -1, v2
	v_add_nc_u32_e32 v4, 1, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v5, -v3, v2, v1
	v_fma_f32 v6, -v4, v2, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_ge_f32_e64 s0, 0, v5
	v_cndmask_b32_e64 v2, v2, v3, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_lt_f32_e64 s0, 0, v6
	v_cndmask_b32_e64 v2, v2, v4, s0
	s_mov_b32 s0, -1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v3, 0x37800000, v2
	v_cndmask_b32_e32 v2, v2, v3, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v1, 0x260
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v2, v1, vcc_lo
	v_cvt_i32_f32_e32 v1, v1
	s_waitcnt vmcnt(0)
	v_cmp_gt_i32_e32 vcc_lo, 0, v0
	s_cbranch_vccnz .LBB2_10
.LBB2_9:                                ; %.lr.ph37
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v0, v0, v1
	s_add_i32 s0, s0, 1
	v_cmp_lt_i32_e32 vcc_lo, -1, v0
	s_cbranch_vccnz .LBB2_9
.LBB2_10:                               ; %._crit_edge38
	s_delay_alu instid0(VALU_DEP_2)
	v_dual_mov_b32 v1, s0 :: v_dual_add_nc_u32 v0, v0, v1
	v_mov_b32_e32 v2, 0
	global_store_b64 v2, v[0:1], s[6:7] offset:4
.LBB2_11:                               ; %Flow54
	s_or_b32 exec_lo, exec_lo, s1
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z24findBestMatchPass2KernelPfPii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 12
		.amdhsa_next_free_sgpr 9
		.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_end2:
	.size	_Z24findBestMatchPass2KernelPfPii, .Lfunc_end2-_Z24findBestMatchPass2KernelPfPii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 616
; NumSgprs: 11
; NumVgprs: 12
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 11
; NumVGPRsForWavesPerEU: 12
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z19updateWeightsKernelPfS_Pifiiiifii ; -- Begin function _Z19updateWeightsKernelPfS_Pifiiiifii
	.globl	_Z19updateWeightsKernelPfS_Pifiiiifii
	.p2align	8
	.type	_Z19updateWeightsKernelPfS_Pifiiiifii,@function
_Z19updateWeightsKernelPfS_Pifiiiifii:  ; @_Z19updateWeightsKernelPfS_Pifiiiifii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x44
	s_load_b256 s[4:11], s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_cmpx_gt_i32_e64 s7, v1
	s_cbranch_execz .LBB3_12
; %bb.1:
	v_cvt_f32_i32_e32 v0, s7
	s_mov_b32 s3, -1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v2, 0x4f800000, v0
	v_cmp_gt_f32_e32 vcc_lo, 0xf800000, v0
	v_cndmask_b32_e32 v0, v0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_2)
	v_sqrt_f32_e32 v2, v0
	s_waitcnt_depctr 0xfff
	v_add_nc_u32_e32 v3, -1, v2
	v_add_nc_u32_e32 v4, 1, v2
	v_fma_f32 v5, -v3, v2, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fma_f32 v6, -v4, v2, v0
	v_cmp_ge_f32_e64 s2, 0, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cndmask_b32_e64 v2, v2, v3, s2
	v_cmp_lt_f32_e64 s2, 0, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e64 v2, v2, v4, s2
	s_mov_b32 s2, 0
	v_mul_f32_e32 v3, 0x37800000, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v2, v2, v3, vcc_lo
	v_cmp_class_f32_e64 vcc_lo, v0, 0x260
	v_dual_mov_b32 v3, v1 :: v_dual_cndmask_b32 v0, v2, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_cvt_i32_f32_e32 v2, v0
.LBB3_2:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_mov_b32_e32 v0, v3
	s_add_i32 s3, s3, 1
	v_mov_b32_e32 v4, s3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v3, v0, v2
	v_cmp_gt_i32_e32 vcc_lo, 0, v3
	s_or_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s2
	s_cbranch_execnz .LBB3_2
; %bb.3:
	s_or_b32 exec_lo, exec_lo, s2
	s_load_b64 s[2:3], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_load_b64 s[2:3], s[2:3], 0x4
	s_waitcnt lgkmcnt(0)
	v_subrev_nc_u32_e32 v0, s2, v0
	v_subrev_nc_u32_e32 v3, s3, v4
	s_mul_i32 s2, s8, s8
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_f32_i32_e32 v4, s2
	v_mul_lo_u32 v0, v0, v0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v3, v3, v3
	s_mov_b32 s2, exec_lo
	v_cvt_f32_i32_e32 v0, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_f32_i32_e32 v3, v3
	v_dual_add_f32 v0, v0, v3 :: v_dual_mov_b32 v3, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ngt_f32_e32 v0, v4
	s_cbranch_execz .LBB3_9
; %bb.4:
	v_cmp_ngt_f32_e64 s3, s4, 0
	s_delay_alu instid0(VALU_DEP_1)
	s_and_b32 vcc_lo, exec_lo, s3
	s_mov_b32 s3, 0
	s_cbranch_vccz .LBB3_6
; %bb.5:
	v_lshlrev_b32_e32 v2, 1, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_f32_i32_e32 v2, v2
	v_div_scale_f32 v3, null, v2, v2, -v0
	v_div_scale_f32 v6, vcc_lo, -v0, v2, -v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v4, v3
	s_waitcnt_depctr 0xfff
	v_fma_f32 v5, -v3, v4, 1.0
	v_fmac_f32_e32 v4, v5, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v5, v6, v4
	v_fma_f32 v7, -v3, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v5, v7, v4
	v_fma_f32 v3, -v3, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v3, v3, v4, v5
	v_div_fixup_f32 v2, v3, v2, -v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v3, 0x3fb8aa3b, v2
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v2
	v_fma_f32 v4, 0x3fb8aa3b, v2, -v3
	v_rndne_f32_e32 v5, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_dual_fmamk_f32 v4, v2, 0x32a5705f, v4 :: v_dual_sub_f32 v3, v3, v5
	v_add_f32_e32 v3, v3, v4
	v_cvt_i32_f32_e32 v4, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_exp_f32_e32 v3, v3
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v3, v3, v4
	v_cndmask_b32_e32 v3, 0, v3, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v2
	s_delay_alu instid0(VALU_DEP_2)
	v_cndmask_b32_e32 v3, 0x7f800000, v3, vcc_lo
	s_branch .LBB3_7
.LBB3_6:
	s_mov_b32 s3, -1
                                        ; implicit-def: $vgpr3
.LBB3_7:                                ; %Flow97
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 vcc_lo, exec_lo, s3
	s_cbranch_vccnz .LBB3_9
; %bb.8:
	v_cvt_f32_i32_e32 v2, s11
	v_cvt_f32_i32_e32 v3, s10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_scale_f32 v4, null, v3, v3, -v2
	v_div_scale_f32 v7, vcc_lo, -v2, v3, -v2
	v_rcp_f32_e32 v5, v4
	s_waitcnt_depctr 0xfff
	v_fma_f32 v6, -v4, v5, 1.0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v5, v6, v5
	v_mul_f32_e32 v6, v7, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v8, -v4, v6, v7
	v_fmac_f32_e32 v6, v8, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v4, -v4, v6, v7
	v_div_fmas_f32 v4, v4, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fixup_f32 v2, v4, v3, -v2
	v_mul_f32_e32 v3, 0x3fb8aa3b, v2
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_fma_f32 v4, 0x3fb8aa3b, v2, -v3
	v_rndne_f32_e32 v5, v3
	v_dual_fmac_f32 v4, 0x32a5705f, v2 :: v_dual_sub_f32 v3, v3, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f32_e32 v3, v3, v4
	v_cvt_i32_f32_e32 v4, v5
	v_exp_f32_e32 v3, v3
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v3, v3, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v3, 0, v3, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v2
	v_cndmask_b32_e32 v2, 0x7f800000, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v2, s4, v2
	v_mul_f32_e32 v2, v2, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v2, v2, v2
	v_div_scale_f32 v3, null, v2, v2, -v0
	v_div_scale_f32 v6, vcc_lo, -v0, v2, -v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v4, v3
	s_waitcnt_depctr 0xfff
	v_fma_f32 v5, -v3, v4, 1.0
	v_fmac_f32_e32 v4, v5, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v5, v6, v4
	v_fma_f32 v7, -v3, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v5, v7, v4
	v_fma_f32 v3, -v3, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_div_fmas_f32 v3, v3, v4, v5
	v_div_fixup_f32 v0, v3, v2, -v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v2, 0x3fb8aa3b, v0
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2ce8ed0, v0
	v_fma_f32 v3, 0x3fb8aa3b, v0, -v2
	v_rndne_f32_e32 v4, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v3, 0x32a5705f, v0
	v_sub_f32_e32 v2, v2, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f32_e32 v2, v2, v3
	v_cvt_i32_f32_e32 v3, v4
	v_exp_f32_e32 v2, v2
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v2, v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v2, 0, v2, vcc_lo
	v_cmp_nlt_f32_e32 vcc_lo, 0x42b17218, v0
	v_cndmask_b32_e32 v3, 0x7f800000, v2, vcc_lo
.LBB3_9:
	s_or_b32 exec_lo, exec_lo, s2
	s_cmp_lt_i32 s5, 1
	s_cbranch_scc1 .LBB3_12
; %bb.10:                               ; %.lr.ph
	v_cvt_f32_i32_e32 v0, s11
	v_cvt_f32_i32_e32 v2, s10
	s_load_b128 s[0:3], s[0:1], 0x0
	s_mul_i32 s10, s6, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_ashr_i32 s11, s10, 31
	v_div_scale_f32 v4, null, v2, v2, v0
	v_div_scale_f32 v7, vcc_lo, v0, v2, v0
	s_lshl_b64 s[10:11], s[10:11], 2
	s_delay_alu instid0(VALU_DEP_2)
	v_rcp_f32_e32 v5, v4
	s_waitcnt_depctr 0xfff
	v_fma_f32 v6, -v4, v5, 1.0
	s_waitcnt lgkmcnt(0)
	s_add_u32 s0, s0, s10
	s_addc_u32 s1, s1, s11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fmac_f32_e32 v5, v6, v5
	v_mul_f32_e32 v6, v7, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v8, -v4, v6, v7
	v_fmac_f32_e32 v6, v8, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_fma_f32 v4, -v4, v6, v7
	v_div_fmas_f32 v4, v4, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_div_fixup_f32 v0, v4, v2, v0
	v_mov_b32_e32 v4, 0
	v_mul_f32_e32 v0, 0xbfb8aa3b, v0
	s_delay_alu instid0(VALU_DEP_1)
	v_exp_f32_e32 v0, v0
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v0, s9, v0
.LBB3_11:                               ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v2, 31, v1
	s_add_i32 s5, s5, -1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_lshlrev_b64 v[5:6], 2, v[1:2]
	v_add_nc_u32_e32 v1, s7, v1
	v_add_co_u32 v5, vcc_lo, s2, v5
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v6, vcc_lo, s3, v6, vcc_lo
	global_load_b32 v2, v4, s[0:1]
	global_load_b32 v7, v[5:6], off
	s_add_u32 s0, s0, 4
	s_addc_u32 s1, s1, 0
	s_cmp_lg_u32 s5, 0
	s_waitcnt vmcnt(0)
	v_sub_f32_e32 v2, v2, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v2, v3, v2
	v_fmac_f32_e32 v7, v0, v2
	global_store_b32 v[5:6], v7, off
	s_cbranch_scc1 .LBB3_11
.LBB3_12:                               ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z19updateWeightsKernelPfS_Pifiiiifii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 312
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 9
		.amdhsa_next_free_sgpr 16
		.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_end3:
	.size	_Z19updateWeightsKernelPfS_Pifiiiifii, .Lfunc_end3-_Z19updateWeightsKernelPfS_Pifiiiifii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 1332
; NumSgprs: 18
; NumVgprs: 9
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 9
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z20propogateInputKernelPfS_S_iii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z20propogateInputKernelPfS_S_iii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z24findBestMatchPass1KernelPfPii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z24findBestMatchPass1KernelPfPii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     11
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z24findBestMatchPass2KernelPfPii
    .private_segment_fixed_size: 0
    .sgpr_count:     11
    .sgpr_spill_count: 0
    .symbol:         _Z24findBestMatchPass2KernelPfPii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     12
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         28
        .size:           4
        .value_kind:     by_value
      - .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:     hidden_block_count_x
      - .offset:         60
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         68
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         70
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         72
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         74
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         76
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         78
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         120
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 312
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z19updateWeightsKernelPfS_Pifiiiifii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z19updateWeightsKernelPfS_Pifiiiifii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     9
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
