/*The number of threads per block and the number of blocks per grid specified in the <<<...>>> syntax can be of type int or dim3.
 Two-dimensional blocks or grids can be specified as in the example above.

Each block within the grid can be identified by a one-dimensional, two-dimensional, or three-dimensional index accessible 
within the kernel through the built-in blockIdx variable. The dimension of the thread block is accessible within the kernel
through the built-in blockDim variable.

Extending the previous MatAdd() example to handle multiple blocks, the code becomes as follows.

*/

#include <stdio.h>
#define N 1024
__device__ int A[N][N];
__device__ int B[N][N];
__device__ int C[N][N];

__global__ void MatAdd()
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>();
        cudaDeviceSynchronize();
}

/*A thread block size of 16x16 (256 threads), although arbitrary in this case, is a common choice. The grid is created with enough 
blocks to have one thread per matrix element as before. For simplicity, this example assumes that the number of threads per 
grid in each dimension is evenly divisible by the number of threads per block in that dimension, although that need not be 
the case.

Thread blocks are required to execute independently: It must be possible to execute them in any order, in parallel or in series. 
This independence requirement allows thread blocks to be scheduled in any order across any number of cores as illustrated by 
Figure 5, enabling programmers to write code that scales with the number of cores.

Threads within a block can cooperate by sharing data through some shared memory and by synchronizing their execution to
 coordinate memory accesses. More precisely, one can specify synchronization points in the kernel by calling the __syncthreads() 
 intrinsic function; __syncthreads() acts as a barrier at which all threads in the block must wait before any is allowed to
  proceed. Shared Memory gives an example of using shared memory.


For efficient cooperation, the shared memory is expected to be a low-latency memory near each processor core 
(much like an L1 cache) and __syncthreads() is expected to be lightweight.
*/
	.file	"tmpxft_002ec423_00000000-6_3-matAdd.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	_Z24__device_stub__Z6MatAddvv
	.type	_Z24__device_stub__Z6MatAddvv, @function
_Z24__device_stub__Z6MatAddvv:
.LFB2052:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	movabsq	$4294967297, %rax
	leaq	8(%rsp), %rcx
	movq	%rsp, %rdx
	leaq	28(%rsp), %rsi
	leaq	16(%rsp), %rdi
	movl	$1, 24(%rsp)
	movl	$1, 36(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 28(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	8(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z6MatAddv(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 112
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L2:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z24__device_stub__Z6MatAddvv, .-_Z24__device_stub__Z6MatAddvv
	.globl	_Z6MatAddv
	.type	_Z6MatAddv, @function
_Z6MatAddv:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z24__device_stub__Z6MatAddvv
	.cfi_endproc
.LFE2053:
	.size	_Z6MatAddv, .-_Z6MatAddv
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$68719476752, %rdx
	movl	$1, %esi
	movabsq	$274877907008, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	call	_Z24__device_stub__Z6MatAddvv
.L9:
	call	cudaDeviceSynchronize@PLT
	xorl	%eax, %eax
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z6MatAddv"
.LC1:
	.string	"A"
.LC2:
	.string	"B"
.LC3:
	.string	"C"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2055:
	.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	_Z6MatAddv(%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	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	movl	$4194304, %r9d
	leaq	_ZL1A(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC2(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$4194304, %r9d
	leaq	_ZL1B(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC3(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$4194304, %r9d
	movq	%rdx, %rcx
	leaq	_ZL1C(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r8
	.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
.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
	.local	_ZL1C
	.comm	_ZL1C,4194304,32
	.local	_ZL1B
	.comm	_ZL1B,4194304,32
	.local	_ZL1A
	.comm	_ZL1A,4194304,32
	.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:

/*The number of threads per block and the number of blocks per grid specified in the <<<...>>> syntax can be of type int or dim3.
 Two-dimensional blocks or grids can be specified as in the example above.

Each block within the grid can be identified by a one-dimensional, two-dimensional, or three-dimensional index accessible 
within the kernel through the built-in blockIdx variable. The dimension of the thread block is accessible within the kernel
through the built-in blockDim variable.

Extending the previous MatAdd() example to handle multiple blocks, the code becomes as follows.

*/


#include <hip/hip_runtime.h>
#include <stdio.h>
#define N 1024
__device__ int A[N][N];
__device__ int B[N][N];
__device__ int C[N][N];

__global__ void MatAdd()
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>();
        hipDeviceSynchronize();
}

/*A thread block size of 16x16 (256 threads), although arbitrary in this case, is a common choice. The grid is created with enough 
blocks to have one thread per matrix element as before. For simplicity, this example assumes that the number of threads per 
grid in each dimension is evenly divisible by the number of threads per block in that dimension, although that need not be 
the case.

Thread blocks are required to execute independently: It must be possible to execute them in any order, in parallel or in series. 
This independence requirement allows thread blocks to be scheduled in any order across any number of cores as illustrated by 
Figure 5, enabling programmers to write code that scales with the number of cores.

Threads within a block can cooperate by sharing data through some shared memory and by synchronizing their execution to
 coordinate memory accesses. More precisely, one can specify synchronization points in the kernel by calling the __syncthreads() 
 intrinsic function; __syncthreads() acts as a barrier at which all threads in the block must wait before any is allowed to
  proceed. Shared Memory gives an example of using shared memory.


For efficient cooperation, the shared memory is expected to be a low-latency memory near each processor core 
(much like an L1 cache) and __syncthreads() is expected to be lightweight.
*/
	.text
	.file	"3-matAdd.hip"
	.globl	_Z21__device_stub__MatAddv      # -- Begin function _Z21__device_stub__MatAddv
	.type	_Z21__device_stub__MatAddv,@function
_Z21__device_stub__MatAddv:             # @_Z21__device_stub__MatAddv
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rbx
	leaq	24(%rsp), %r14
	leaq	16(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movq	%r12, %rcx
	callq	__hipPopCallConfiguration
	movq	(%rbx), %rsi
	movl	8(%rbx), %edx
	movq	(%r14), %rcx
	movl	8(%r14), %r8d
	movq	%rsp, %r9
	movl	$_Z6MatAddv, %edi
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	pushq	(%r15)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$72, %rsp
	.cfi_adjust_cfa_offset -72
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z21__device_stub__MatAddv, .Lfunc_end0-_Z21__device_stub__MatAddv
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movabsq	$274877907008, %rdi             # imm = 0x4000000040
	movabsq	$68719476752, %rdx              # imm = 0x1000000010
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	callq	_Z21__device_stub__MatAddv
.LBB1_2:
	callq	hipDeviceSynchronize
	xorl	%eax, %eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.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
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	subq	$32, %rsp
	.cfi_adjust_cfa_offset 32
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z6MatAddv, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	movl	$A, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$4194304, %r9d                  # imm = 0x400000
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$B, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movl	$4194304, %r9d                  # imm = 0x400000
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$C, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movl	$4194304, %r9d                  # imm = 0x400000
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$__hip_module_dtor, %edi
	popq	%rbx
	.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	A,@object                       # @A
	.local	A
	.comm	A,4194304,16
	.type	B,@object                       # @B
	.local	B
	.comm	B,4194304,16
	.type	C,@object                       # @C
	.local	C
	.comm	C,4194304,16
	.type	_Z6MatAddv,@object              # @_Z6MatAddv
	.section	.rodata,"a",@progbits
	.globl	_Z6MatAddv
	.p2align	3, 0x0
_Z6MatAddv:
	.quad	_Z21__device_stub__MatAddv
	.size	_Z6MatAddv, 8

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

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"A"
	.size	.L__unnamed_2, 2

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

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"C"
	.size	.L__unnamed_4, 2

	.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 _Z21__device_stub__MatAddv
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym A
	.addrsig_sym B
	.addrsig_sym C
	.addrsig_sym _Z6MatAddv
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z6MatAddv              ; -- Begin function _Z6MatAddv
	.globl	_Z6MatAddv
	.p2align	8
	.type	_Z6MatAddv,@function
_Z6MatAddv:                             ; @_Z6MatAddv
; %bb.0:
	s_load_b32 s0, s[0:1], 0xc
	v_and_b32_e32 v2, 0x3ff, v0
	v_bfe_u32 v4, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s1, s0, 16
	s_and_b32 s0, s0, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[0:1], null, s14, s0, v[2:3]
	v_mad_u64_u32 v[2:3], null, s15, s1, v[4:5]
	s_mov_b32 s0, exec_lo
	v_max_i32_e32 v1, v0, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e32 0x400, v1
	s_cbranch_execz .LBB0_2
; %bb.1:
	v_ashrrev_i32_e32 v3, 31, v2
	v_ashrrev_i32_e32 v1, 31, v0
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, A@rel32@lo+4
	s_addc_u32 s1, s1, A@rel32@hi+12
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, B@rel32@lo+4
	s_addc_u32 s3, s3, B@rel32@hi+12
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	v_lshlrev_b64 v[0:1], 12, v[0:1]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, v0, v2
	v_add_co_ci_u32_e32 v5, vcc_lo, v1, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, v4, s0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v5, vcc_lo
	v_add_co_u32 v2, vcc_lo, v4, s2
	v_add_co_ci_u32_e32 v3, vcc_lo, s3, v5, vcc_lo
	global_load_b32 v6, v[0:1], off
	global_load_b32 v2, v[2:3], off
	s_getpc_b64 s[0:1]
	s_add_u32 s0, s0, C@rel32@lo+4
	s_addc_u32 s1, s1, C@rel32@hi+12
	v_add_co_u32 v0, vcc_lo, v4, s0
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v5, vcc_lo
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v2, v2, v6
	global_store_b32 v[0:1], v2, off
.LBB0_2:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z6MatAddv
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 256
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.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	_Z6MatAddv, .Lfunc_end0-_Z6MatAddv
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 268
; 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: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	A                       ; @A
	.type	A,@object
	.section	.bss,"aw",@nobits
	.globl	A
	.p2align	4, 0x0
A:
	.zero	4194304
	.size	A, 4194304

	.protected	B                       ; @B
	.type	B,@object
	.globl	B
	.p2align	4, 0x0
B:
	.zero	4194304
	.size	B, 4194304

	.protected	C                       ; @C
	.type	C,@object
	.globl	C
	.p2align	4, 0x0
C:
	.zero	4194304
	.size	C, 4194304

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.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 A
	.addrsig_sym B
	.addrsig_sym C
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         4
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         8
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         12
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         14
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         16
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         18
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         20
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         22
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         40
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         48
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         56
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         64
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 256
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z6MatAddv
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z6MatAddv.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     7
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
