#include "includes.h"
//Training of the CNN is done using Keras. After training for 10 epochs, the obtained accuracy on the training data set is 99.70 and on the test data set is 99.14.
//This model implements the following layes in order- 2DConvolution---->Maxpooling---->2D Convolution---->Maxpooling---->Fully_connected layer---->Fully_connected layer.
//The image is a 28*28 greyscale image. The specifications of the layers are as follows:
//Layer_0: Convolution: 32 3*3 kernels with no padding and 1 stride.
//Layer_1: Maxpooling: 2*2 filters with with no padding and 1 stride.
//Layer_2: Convolution: 64 3*3 kernels with no padding and 1 stride.
//Layer_3: Maxpooling: 2*2 filters with with no padding and 1 stride.
//Layer_4: Flattening
//Layer_5: Fully connected / dense layer with 1024 output units.
//Layer_6: Dropout (done during training only).
//Layer_7: Fully connected / dense layer with 10 output units.

//All arrays and matrices are designed to be row ordered in this implementation.



//Kernel that does convolution. This convolution is done by each thread identifying that patch or portion of the image that it is responsible for its result and does the multiplication and addition of it's patche's values with the suitable kernel.
//The depth of the output image is the number of kernels.

//Kernel that does maxpooling.

//This kernel implements the fully connected layers.


__global__ void dense_kernel(int num_input, int num_output, double* gpu_in, double* weights, double* biases, double* gpu_out, int num_classes)
{
int tid = blockDim.x*blockIdx.x + threadIdx.x;
if (tid >= num_output) return;
double sum = 0.0l;
//The weights are extracted from Keras such that all the weights to one output node appears together, followed by weights to the next node and so on.
//Thus, each output node will be a multiply add of adjacent weight values with the input nodes.
for (int count = 0; count < num_input; count++) {
sum += gpu_in[count] * weights[tid*num_input + count];
}
sum += biases[tid];

//Activation: If the layer is the final layer, then don't do anything, otherwise relu activation max(0,value) is taken.
if ((num_output) != num_classes) {
if (sum < 0.0) {
sum = 0.0l;
}
}
gpu_out[tid] = sum;
}
	.file	"tmpxft_0027cf02_00000000-6_dense_kernel.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z41__device_stub__Z12dense_kerneliiPdS_S_S_iiiPdS_S_S_i
	.type	_Z41__device_stub__Z12dense_kerneliiPdS_S_S_iiiPdS_S_S_i, @function
_Z41__device_stub__Z12dense_kerneliiPdS_S_S_iiiPdS_S_S_i:
.LFB2032:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 32(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 24(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 16(%rsp)
	movq	%r9, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z12dense_kerneliiPdS_S_S_i(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	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 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
.LFE2032:
	.size	_Z41__device_stub__Z12dense_kerneliiPdS_S_S_iiiPdS_S_S_i, .-_Z41__device_stub__Z12dense_kerneliiPdS_S_S_iiiPdS_S_S_i
	.globl	_Z12dense_kerneliiPdS_S_S_i
	.type	_Z12dense_kerneliiPdS_S_S_i, @function
_Z12dense_kerneliiPdS_S_S_i:
.LFB2033:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z12dense_kerneliiPdS_S_S_iiiPdS_S_S_i
	.cfi_endproc
.LFE2033:
	.size	_Z12dense_kerneliiPdS_S_S_i, .-_Z12dense_kerneliiPdS_S_S_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z12dense_kerneliiPdS_S_S_i"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2035:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z12dense_kerneliiPdS_S_S_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2035:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
#include "includes.h"
//Training of the CNN is done using Keras. After training for 10 epochs, the obtained accuracy on the training data set is 99.70 and on the test data set is 99.14.
//This model implements the following layes in order- 2DConvolution---->Maxpooling---->2D Convolution---->Maxpooling---->Fully_connected layer---->Fully_connected layer.
//The image is a 28*28 greyscale image. The specifications of the layers are as follows:
//Layer_0: Convolution: 32 3*3 kernels with no padding and 1 stride.
//Layer_1: Maxpooling: 2*2 filters with with no padding and 1 stride.
//Layer_2: Convolution: 64 3*3 kernels with no padding and 1 stride.
//Layer_3: Maxpooling: 2*2 filters with with no padding and 1 stride.
//Layer_4: Flattening
//Layer_5: Fully connected / dense layer with 1024 output units.
//Layer_6: Dropout (done during training only).
//Layer_7: Fully connected / dense layer with 10 output units.

//All arrays and matrices are designed to be row ordered in this implementation.



//Kernel that does convolution. This convolution is done by each thread identifying that patch or portion of the image that it is responsible for its result and does the multiplication and addition of it's patche's values with the suitable kernel.
//The depth of the output image is the number of kernels.

//Kernel that does maxpooling.

//This kernel implements the fully connected layers.


__global__ void dense_kernel(int num_input, int num_output, double* gpu_in, double* weights, double* biases, double* gpu_out, int num_classes)
{
int tid = blockDim.x*blockIdx.x + threadIdx.x;
if (tid >= num_output) return;
double sum = 0.0l;
//The weights are extracted from Keras such that all the weights to one output node appears together, followed by weights to the next node and so on.
//Thus, each output node will be a multiply add of adjacent weight values with the input nodes.
for (int count = 0; count < num_input; count++) {
sum += gpu_in[count] * weights[tid*num_input + count];
}
sum += biases[tid];

//Activation: If the layer is the final layer, then don't do anything, otherwise relu activation max(0,value) is taken.
if ((num_output) != num_classes) {
if (sum < 0.0) {
sum = 0.0l;
}
}
gpu_out[tid] = sum;
}
	.text
	.file	"dense_kernel.hip"
	.globl	_Z27__device_stub__dense_kerneliiPdS_S_S_i # -- Begin function _Z27__device_stub__dense_kerneliiPdS_S_S_i
	.type	_Z27__device_stub__dense_kerneliiPdS_S_S_i,@function
_Z27__device_stub__dense_kerneliiPdS_S_S_i: # @_Z27__device_stub__dense_kerneliiPdS_S_S_i
	.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	$160, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	56(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	48(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	32(%rsp), %r8
	movq	%r9, (%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	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%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	$_Z12dense_kerneliiPdS_S_S_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	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	_Z27__device_stub__dense_kerneliiPdS_S_S_i, .Lfunc_end0-_Z27__device_stub__dense_kerneliiPdS_S_S_i
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB1_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB1_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z12dense_kerneliiPdS_S_S_i, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end1:
	.size	__hip_module_ctor, .Lfunc_end1-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB2_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB2_2:
	retq
.Lfunc_end2:
	.size	__hip_module_dtor, .Lfunc_end2-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z12dense_kerneliiPdS_S_S_i,@object # @_Z12dense_kerneliiPdS_S_S_i
	.section	.rodata,"a",@progbits
	.globl	_Z12dense_kerneliiPdS_S_S_i
	.p2align	3, 0x0
_Z12dense_kerneliiPdS_S_S_i:
	.quad	_Z27__device_stub__dense_kerneliiPdS_S_S_i
	.size	_Z12dense_kerneliiPdS_S_S_i, 8

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

	.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 _Z27__device_stub__dense_kerneliiPdS_S_S_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z12dense_kerneliiPdS_S_S_i
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z12dense_kerneliiPdS_S_S_i ; -- Begin function _Z12dense_kerneliiPdS_S_S_i
	.globl	_Z12dense_kerneliiPdS_S_S_i
	.p2align	8
	.type	_Z12dense_kerneliiPdS_S_S_i,@function
_Z12dense_kerneliiPdS_S_S_i:            ; @_Z12dense_kerneliiPdS_S_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s4, s[0:1], 0x3c
	s_load_b64 s[2:3], s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s4, v[0:1]
	s_mov_b32 s4, exec_lo
	v_cmpx_gt_i32_e64 s3, v1
	s_cbranch_execz .LBB0_6
; %bb.1:                                ; %.preheader
	s_load_b256 s[4:11], s[0:1], 0x8
	s_cmp_lt_i32 s2, 1
	s_cbranch_scc1 .LBB0_4
; %bb.2:                                ; %.lr.ph
	v_mul_lo_u32 v2, v1, s2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[5:6], 3, v[2:3]
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s6, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s7, v6, vcc_lo
.LBB0_3:                                ; =>This Inner Loop Header: Depth=1
	global_load_b64 v[7:8], v[5:6], off
	s_load_b64 s[6:7], s[4:5], 0x0
	v_add_co_u32 v5, vcc_lo, v5, 8
	s_add_i32 s2, s2, -1
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v6, vcc_lo
	s_add_u32 s4, s4, 8
	s_addc_u32 s5, s5, 0
	s_cmp_eq_u32 s2, 0
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_fma_f64 v[3:4], s[6:7], v[7:8], v[3:4]
	s_cbranch_scc0 .LBB0_3
	s_branch .LBB0_5
.LBB0_4:
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
.LBB0_5:                                ; %Flow54
	v_ashrrev_i32_e32 v2, 31, v1
	s_load_b32 s0, s[0:1], 0x28
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v5, vcc_lo, s8, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v6, vcc_lo, s9, v1, vcc_lo
	global_load_b64 v[5:6], v[5:6], off
	s_cmp_lg_u32 s3, s0
	s_cselect_b32 s0, -1, 0
	s_waitcnt vmcnt(0)
	v_add_f64 v[2:3], v[3:4], v[5:6]
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_gt_f64_e32 vcc_lo, 0, v[2:3]
	s_and_b32 s0, s0, vcc_lo
	v_add_co_u32 v0, vcc_lo, s10, v0
	v_cndmask_b32_e64 v3, v3, 0, s0
	v_cndmask_b32_e64 v2, v2, 0, s0
	v_add_co_ci_u32_e32 v1, vcc_lo, s11, v1, vcc_lo
	global_store_b64 v[0:1], v[2:3], off
.LBB0_6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z12dense_kerneliiPdS_S_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.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_end0:
	.size	_Z12dense_kerneliiPdS_S_S_i, .Lfunc_end0-_Z12dense_kerneliiPdS_S_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 332
; 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:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .offset:         4
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z12dense_kerneliiPdS_S_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z12dense_kerneliiPdS_S_S_i.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
