/***
	Copyright 2012 Injung Kim

	Licensed under the Apache License, Version 2.0 (the "License");
	you may not use this file except in compliance with the License.
	You may obtain a copy of the License at

		http://www.apache.org/licenses/LICENSE-2.0

	Unless required by applicable law or agreed to in writing, software
	distributed under the License is distributed on an "AS IS" BASIS,
	WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
	See the License for the specific language governing permissions and
	limitations under the License.
***/


#include <stdio.h>

#include <cuda_runtime.h>

__global__ void Propagate_kernel(float *pDevInput, float *pDevWeight, int inDim, int outDim, float *pDevOutput);

void PropagateOnDevice(float *pInput, float *pWeight, int inDim, int outDim, float *pOutput)
{
	float *pDevInput = NULL, *pDevOutput = NULL, *pDevWeight = NULL;

	// allocate device memory	
	cudaMalloc((void**)&pDevInput, inDim * sizeof(float));
	cudaMalloc((void**)&pDevOutput, outDim * sizeof(float));
	cudaMalloc((void**)&pDevWeight, (inDim+1)*outDim * sizeof(float));

	// copy input and weight to device memory
	cudaMemcpy(pDevInput, pInput, inDim * sizeof(float), cudaMemcpyHostToDevice);
	cudaMemcpy(pDevWeight, pWeight, (inDim+1)*outDim * sizeof(float), cudaMemcpyHostToDevice);
	
	// compute output on CUDA device
	Propagate_kernel<<<1, outDim>>>(pDevInput, pDevWeight, inDim, outDim, pDevOutput);
	
	// copy output to host memory
	cudaMemcpy(pOutput, pDevOutput, outDim * sizeof(float), cudaMemcpyDeviceToHost);
	
	// deallocate device memory
	cudaFree(pDevInput);
	cudaFree(pDevOutput);
	cudaFree(pDevWeight);
}

__global__ void Propagate_kernel(float *pDevInput, float *pDevWeight, int inDim, int outDim, float *pDevOutput)
{
	// idx is thread index
	int idx = blockIdx.x * blockDim.x + threadIdx.x;

	while(idx < outDim){
		float *w = &pDevWeight[idx * (inDim + 1)];

		float net = 0.F;
		for(int i = 0; i < inDim; i++)
				net += pDevInput[i] * w[i];
		net += w[inDim];						// add bias

		pDevOutput[idx] = 1.F/(1.F + (float)exp(-net));

		idx += blockDim.x * gridDim.x; 	//	blockDim.x * gridDim.x is # of threads
	}
}

	.file	"tmpxft_00331b61_00000000-6_HGULayer_CUDA.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z42__device_stub__Z16Propagate_kernelPfS_iiS_PfS_iiS_
	.type	_Z42__device_stub__Z16Propagate_kernelPfS_iiS_PfS_iiS_, @function
_Z42__device_stub__Z16Propagate_kernelPfS_iiS_PfS_iiS_:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 8(%rsp)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%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)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z16Propagate_kernelPfS_iiS_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 176
	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 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z42__device_stub__Z16Propagate_kernelPfS_iiS_PfS_iiS_, .-_Z42__device_stub__Z16Propagate_kernelPfS_iiS_PfS_iiS_
	.globl	_Z16Propagate_kernelPfS_iiS_
	.type	_Z16Propagate_kernelPfS_iiS_, @function
_Z16Propagate_kernelPfS_iiS_:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z16Propagate_kernelPfS_iiS_PfS_iiS_
	.cfi_endproc
.LFE2053:
	.size	_Z16Propagate_kernelPfS_iiS_, .-_Z16Propagate_kernelPfS_iiS_
	.globl	_Z17PropagateOnDevicePfS_iiS_
	.type	_Z17PropagateOnDevicePfS_iiS_, @function
_Z17PropagateOnDevicePfS_iiS_:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movslq	%edx, %rdx
	movq	%rdi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rsi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%r8, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	salq	$2, %rdx
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%ecx, %ebp
	movq	%rdx, %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leal	1(%r12), %ebx
	imull	%ebp, %ebx
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rdi
	xorl	%eax, %eax
	movq	%rdx, 8(%rsp)
	movq	%rax, 24(%rsp)
	movslq	%ebx, %rbx
	movq	%rax, 32(%rsp)
	salq	$2, %rbx
	movq	%rax, 40(%rsp)
	call	cudaMalloc@PLT
	movslq	%ebp, %rax
	leaq	32(%rsp), %rdi
	salq	$2, %rax
	movq	%rax, %rsi
	movq	%rax, (%rsp)
	call	cudaMalloc@PLT
	leaq	40(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rdi
	movq	%r15, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rdi
	movq	%rbx, %rdx
	movq	%r14, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	%ebp, 60(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movq	%rdi, 64(%rsp)
	movl	68(%rsp), %ecx
	movq	60(%rsp), %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	movq	32(%rsp), %r8
	movq	40(%rsp), %rsi
	movl	%ebp, %ecx
	movl	%r12d, %edx
	movq	24(%rsp), %rdi
	call	_Z42__device_stub__Z16Propagate_kernelPfS_iiS_PfS_iiS_
.L9:
	movq	(%rsp), %rdx
	movq	32(%rsp), %rsi
	movl	$2, %ecx
	movq	%r13, %rdi
	call	cudaMemcpy@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$88, %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
.LFE2027:
	.size	_Z17PropagateOnDevicePfS_iiS_, .-_Z17PropagateOnDevicePfS_iiS_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z16Propagate_kernelPfS_iiS_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2055:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z16Propagate_kernelPfS_iiS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2055:
	.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:

/***
	Copyright 2012 Injung Kim

	Licensed under the Apache License, Version 2.0 (the "License");
	you may not use this file except in compliance with the License.
	You may obtain a copy of the License at

		http://www.apache.org/licenses/LICENSE-2.0

	Unless required by applicable law or agreed to in writing, software
	distributed under the License is distributed on an "AS IS" BASIS,
	WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
	See the License for the specific language governing permissions and
	limitations under the License.
***/


#include <stdio.h>

#include <hip/hip_runtime.h>

__global__ void Propagate_kernel(float *pDevInput, float *pDevWeight, int inDim, int outDim, float *pDevOutput);

void PropagateOnDevice(float *pInput, float *pWeight, int inDim, int outDim, float *pOutput)
{
	float *pDevInput = NULL, *pDevOutput = NULL, *pDevWeight = NULL;

	// allocate device memory	
	hipMalloc((void**)&pDevInput, inDim * sizeof(float));
	hipMalloc((void**)&pDevOutput, outDim * sizeof(float));
	hipMalloc((void**)&pDevWeight, (inDim+1)*outDim * sizeof(float));

	// copy input and weight to device memory
	hipMemcpy(pDevInput, pInput, inDim * sizeof(float), hipMemcpyHostToDevice);
	hipMemcpy(pDevWeight, pWeight, (inDim+1)*outDim * sizeof(float), hipMemcpyHostToDevice);
	
	// compute output on CUDA device
	Propagate_kernel<<<1, outDim>>>(pDevInput, pDevWeight, inDim, outDim, pDevOutput);
	
	// copy output to host memory
	hipMemcpy(pOutput, pDevOutput, outDim * sizeof(float), hipMemcpyDeviceToHost);
	
	// deallocate device memory
	hipFree(pDevInput);
	hipFree(pDevOutput);
	hipFree(pDevWeight);
}

__global__ void Propagate_kernel(float *pDevInput, float *pDevWeight, int inDim, int outDim, float *pDevOutput)
{
	// idx is thread index
	int idx = blockIdx.x * blockDim.x + threadIdx.x;

	while(idx < outDim){
		float *w = &pDevWeight[idx * (inDim + 1)];

		float net = 0.F;
		for(int i = 0; i < inDim; i++)
				net += pDevInput[i] * w[i];
		net += w[inDim];						// add bias

		pDevOutput[idx] = 1.F/(1.F + (float)exp(-net));

		idx += blockDim.x * gridDim.x; 	//	blockDim.x * gridDim.x is # of threads
	}
}

	.text
	.file	"HGULayer_CUDA.hip"
	.globl	_Z17PropagateOnDevicePfS_iiS_   # -- Begin function _Z17PropagateOnDevicePfS_iiS_
	.type	_Z17PropagateOnDevicePfS_iiS_,@function
_Z17PropagateOnDevicePfS_iiS_:          # @_Z17PropagateOnDevicePfS_iiS_
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	.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	%r8, 64(%rsp)                   # 8-byte Spill
	movl	%ecx, %ebx
	movq	%rsi, 48(%rsp)                  # 8-byte Spill
	movq	%rdi, 40(%rsp)                  # 8-byte Spill
	xorl	%eax, %eax
	leaq	32(%rsp), %r14
	movq	%rax, (%r14)
	leaq	16(%rsp), %r15
	movq	%rax, (%r15)
	leaq	8(%rsp), %rcx
	movq	%rax, (%rcx)
	movl	%edx, 24(%rsp)                  # 4-byte Spill
	movslq	%edx, %r13
	leaq	(,%r13,4), %r12
	movq	%r14, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	movslq	%ebx, %rbp
	leaq	(,%rbp,4), %rsi
	movq	%r15, %rdi
	movq	%rsi, 56(%rsp)                  # 8-byte Spill
	callq	hipMalloc
	incl	%r13d
	movl	%ebx, 28(%rsp)                  # 4-byte Spill
	imull	%ebx, %r13d
	movslq	%r13d, %r15
	shlq	$2, %r15
	leaq	8(%rsp), %rbx
	movq	%rbx, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movq	(%r14), %rdi
	movq	40(%rsp), %rsi                  # 8-byte Reload
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbx), %rdi
	movq	48(%rsp), %rsi                  # 8-byte Reload
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	%ebp, %edx
	btsq	$32, %rdx
	movabsq	$4294967296, %rdi               # imm = 0x100000000
	orq	$1, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB0_2
# %bb.1:
	movq	32(%rsp), %rdi
	movq	8(%rsp), %rsi
	movq	16(%rsp), %r8
	movl	24(%rsp), %edx                  # 4-byte Reload
	movl	28(%rsp), %ecx                  # 4-byte Reload
	callq	_Z31__device_stub__Propagate_kernelPfS_iiS_
.LBB0_2:
	movq	16(%rsp), %rsi
	movq	64(%rsp), %rdi                  # 8-byte Reload
	movq	56(%rsp), %rdx                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	addq	$72, %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
.Lfunc_end0:
	.size	_Z17PropagateOnDevicePfS_iiS_, .Lfunc_end0-_Z17PropagateOnDevicePfS_iiS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z31__device_stub__Propagate_kernelPfS_iiS_ # -- Begin function _Z31__device_stub__Propagate_kernelPfS_iiS_
	.type	_Z31__device_stub__Propagate_kernelPfS_iiS_,@function
_Z31__device_stub__Propagate_kernelPfS_iiS_: # @_Z31__device_stub__Propagate_kernelPfS_iiS_
	.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	$128, %rsp
	.cfi_def_cfa_offset 176
	.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), %rdi
	movq	%rsi, (%rdi)
	leaq	4(%rsp), %rsi
	movl	%edx, (%rsi)
	movq	%rsp, %rdx
	movl	%ecx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%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	$_Z16Propagate_kernelPfS_iiS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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	_Z31__device_stub__Propagate_kernelPfS_iiS_, .Lfunc_end1-_Z31__device_stub__Propagate_kernelPfS_iiS_
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z16Propagate_kernelPfS_iiS_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__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	.LBB3_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
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z16Propagate_kernelPfS_iiS_,@object # @_Z16Propagate_kernelPfS_iiS_
	.section	.rodata,"a",@progbits
	.globl	_Z16Propagate_kernelPfS_iiS_
	.p2align	3, 0x0
_Z16Propagate_kernelPfS_iiS_:
	.quad	_Z31__device_stub__Propagate_kernelPfS_iiS_
	.size	_Z16Propagate_kernelPfS_iiS_, 8

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

	.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 _Z31__device_stub__Propagate_kernelPfS_iiS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z16Propagate_kernelPfS_iiS_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z16Propagate_kernelPfS_iiS_ ; -- Begin function _Z16Propagate_kernelPfS_iiS_
	.globl	_Z16Propagate_kernelPfS_iiS_
	.p2align	8
	.type	_Z16Propagate_kernelPfS_iiS_,@function
_Z16Propagate_kernelPfS_iiS_:           ; @_Z16Propagate_kernelPfS_iiS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s6, s[0:1], 0x2c
	s_load_b64 s[2:3], s[0:1], 0x10
	s_add_u32 s4, s0, 32
	s_addc_u32 s5, s1, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s10, s6, 0xffff
	s_mov_b32 s6, exec_lo
	v_mad_u64_u32 v[1:2], null, s15, s10, v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e64 s3, v1
	s_cbranch_execz .LBB0_6
; %bb.1:                                ; %.lr.ph29
	s_load_b32 s15, s[4:5], 0x0
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[8:9], s[0:1], 0x18
	s_add_i32 s12, s2, 1
	s_cmp_gt_i32 s2, 0
	v_mul_lo_u32 v3, v1, s12
	v_mov_b32_e32 v0, 0
	s_mov_b32 s0, s2
	s_cselect_b32 s14, -1, 0
	s_ashr_i32 s1, s2, 31
	s_mov_b32 s13, 0
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s15, s15, s10
	s_lshl_b64 s[10:11], s[0:1], 2
	s_mul_i32 s16, s15, s12
.LBB0_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
	v_mov_b32_e32 v2, 0
	s_and_not1_b32 vcc_lo, exec_lo, s14
	s_cbranch_vccnz .LBB0_5
; %bb.3:                                ; %.lr.ph.preheader
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_ashrrev_i32_e32 v4, 31, v3
	v_mov_b32_e32 v2, 0
	s_mov_b64 s[0:1], s[4:5]
	s_mov_b32 s17, s2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[4:5], 2, v[3:4]
	v_add_co_u32 v4, vcc_lo, s6, v4
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v5, vcc_lo
.LBB0_4:                                ; %.lr.ph
                                        ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v6, v0, s[0:1]
	global_load_b32 v7, v[4:5], off
	v_add_co_u32 v4, vcc_lo, v4, 4
	s_add_i32 s17, s17, -1
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_add_u32 s0, s0, 4
	s_addc_u32 s1, s1, 0
	s_cmp_eq_u32 s17, 0
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v2, v6, v7
	s_cbranch_scc0 .LBB0_4
.LBB0_5:                                ; %Flow50
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_mul_lo_u32 v4, v1, s12
	s_add_u32 s0, s10, s6
	s_addc_u32 s1, s11, s7
	v_add_nc_u32_e32 v3, s16, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v5, 31, v4
	v_lshlrev_b64 v[4:5], 2, v[4:5]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, s0, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s1, v5, vcc_lo
	global_load_b32 v4, v[4:5], off
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v2, v2, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f32_e32 v4, 0xbfb8aa3b, v2
	v_cmp_nlt_f32_e32 vcc_lo, 0x42ce8ed0, v2
	v_fma_f32 v5, 0xbfb8aa3b, v2, -v4
	v_rndne_f32_e32 v6, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v5, 0xb2a5705f, v2
	v_sub_f32_e32 v4, v4, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add_f32_e32 v4, v4, v5
	v_cvt_i32_f32_e32 v5, v6
	v_exp_f32_e32 v4, v4
	s_waitcnt_depctr 0xfff
	v_ldexp_f32 v4, v4, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v4, 0, v4, vcc_lo
	v_cmp_ngt_f32_e32 vcc_lo, 0xc2b17218, v2
	v_cndmask_b32_e32 v2, 0x7f800000, v4, vcc_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_f32_e32 v6, 1.0, v2
	v_div_scale_f32 v4, null, v6, v6, 1.0
	v_div_scale_f32 v5, vcc_lo, 1.0, v6, 1.0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_f32_e32 v7, v4
	s_waitcnt_depctr 0xfff
	v_fma_f32 v2, -v4, v7, 1.0
	v_fmac_f32_e32 v7, v2, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f32_e32 v8, v5, v7
	v_fma_f32 v2, -v4, v8, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_fmac_f32_e32 v8, v2, v7
	v_ashrrev_i32_e32 v2, 31, v1
	v_fma_f32 v9, -v4, v8, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[4:5], 2, v[1:2]
	v_add_nc_u32_e32 v1, s15, v1
	v_div_fmas_f32 v2, v9, v7, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_le_i32_e32 vcc_lo, s3, v1
	v_add_co_u32 v4, s0, s8, v4
	s_delay_alu instid0(VALU_DEP_3)
	v_div_fixup_f32 v2, v2, v6, 1.0
	v_add_co_ci_u32_e64 v5, s0, s9, v5, s0
	s_or_b32 s13, vcc_lo, s13
	global_store_b32 v[4:5], v2, off
	s_and_not1_b32 exec_lo, exec_lo, s13
	s_cbranch_execnz .LBB0_2
.LBB0_6:                                ; %._crit_edge30
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z16Propagate_kernelPfS_iiS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 10
		.amdhsa_next_free_sgpr 18
		.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	_Z16Propagate_kernelPfS_iiS_, .Lfunc_end0-_Z16Propagate_kernelPfS_iiS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 608
; NumSgprs: 20
; NumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 20
; NumVGPRsForWavesPerEU: 10
; 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
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z16Propagate_kernelPfS_iiS_
    .private_segment_fixed_size: 0
    .sgpr_count:     20
    .sgpr_spill_count: 0
    .symbol:         _Z16Propagate_kernelPfS_iiS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     10
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
