/* The matrix addition example on CUDA
This program first reads two matrices from two text files namely matA.txt and matB.txt
Then it does matrix addition on the two matrices
Finally it saves the answer matrix in a text file named ans.txt
The program also measures the time taken for matrix operation on CUDA
*/

#include <stdio.h>

//define file names here
#define MATRIXA "matA.txt"
#define MATRIXB "matB.txt"
#define MATRIXANS "ans.txt"

//define the sizes of the matrices here
#define ROWS 16
#define COLS 16
#define SIZE ROWS*COLS

//kernel that does the matrix addition. Just add each element to the respective one
__global__ void addMatrix(int *ans_cuda,int *matA_cuda,int *matB_cuda){
	int row = threadIdx.y;
	int col = threadIdx.x;
	int position = row*COLS + col;
	ans_cuda[position]=matA_cuda[position]+matB_cuda[position];
}

int main(){
	
	//open the files 
	FILE *filematA = fopen(MATRIXA, "r");
	FILE *filematB = fopen(MATRIXB, "r");
	FILE *fileans = fopen(MATRIXANS, "w");
	
	//allocate matrices
	int matA[SIZE];
	int matB[SIZE];
	int ans[SIZE];
	
	//read the input matrices from file
	int row,col;
	for(row=0;row<ROWS;row++){
		for(col=0;col<COLS;col++){
			int position = row*COLS + col;
			fscanf(filematA, "%d", &matA[position]);
			fscanf(filematB, "%d", &matB[position]);
			ans[position]=0;
		}
	}
	

/*************************CUDA STUFF STARTS HERE************************/	
	
	//variables for time measurements
	cudaEvent_t start,stop;
	float elapsedtime;
	
	//pointers for cuda memory locations
	int *matA_cuda;
	int *matB_cuda;
	int *ans_cuda;	

	//the moment at which we start measuring the time
	cudaEventCreate(&start);
	cudaEventRecord(start,0);
	
	//allocate memory in cuda
	cudaMalloc((void **)&matA_cuda,sizeof(int)*SIZE);
	cudaMalloc((void **)&matB_cuda,sizeof(int)*SIZE);	
	cudaMalloc((void **)&ans_cuda,sizeof(int)*SIZE);
		
	//copy contents from ram to cuda
	cudaMemcpy(matA_cuda, matA, sizeof(int)*SIZE, cudaMemcpyHostToDevice);
	cudaMemcpy(matB_cuda, matB, sizeof(int)*SIZE, cudaMemcpyHostToDevice);	

	//thread configuration 
	dim3 numBlocks(1,1);
	dim3 threadsPerBlock(COLS,ROWS);
	
	//do the matrix addition on CUDA
	addMatrix<<<numBlocks,threadsPerBlock>>>(ans_cuda,matA_cuda,matB_cuda);

	//copy the answer back
	cudaMemcpy(ans, ans_cuda, sizeof(int)*SIZE, cudaMemcpyDeviceToHost);	

	//the moment at which we stop measuring time 
	cudaEventCreate(&stop);
	cudaEventRecord(stop,0);
	cudaEventSynchronize(stop);
	
	//free the memory we allocated on CUDA
	cudaFree(matA_cuda);
	cudaFree(matB_cuda);
	cudaFree(ans_cuda);
	
/*************************CUDA STUFF ENDS HERE************************/


	//write the answer to the text file
	for(row=0;row<ROWS;row++){
		for(col=0;col<COLS;col++){
			int position = row*COLS + col;
			fprintf(fileans, "%d ", ans[position]);
		}
		fprintf(fileans, "\n");
	}
	
	//Find and print the elapsed time
	cudaEventElapsedTime(&elapsedtime,start,stop);
	printf("Time spent for operation is %.10f seconds\n",elapsedtime/(float)1000);

	fclose(filematA);
	fclose(filematB);
	fclose(fileans);
		
	return 0;
}
	.file	"tmpxft_00374034_00000000-6_matrixadd.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	_Z32__device_stub__Z9addMatrixPiS_S_PiS_S_
	.type	_Z32__device_stub__Z9addMatrixPiS_S_PiS_S_, @function
_Z32__device_stub__Z9addMatrixPiS_S_PiS_S_:
.LFB2052:
	.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
	movq	%rdx, 8(%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	8(%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	.L2
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z9addMatrixPiS_S_(%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
.L2:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2052:
	.size	_Z32__device_stub__Z9addMatrixPiS_S_PiS_S_, .-_Z32__device_stub__Z9addMatrixPiS_S_PiS_S_
	.globl	_Z9addMatrixPiS_S_
	.type	_Z9addMatrixPiS_S_, @function
_Z9addMatrixPiS_S_:
.LFB2053:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z9addMatrixPiS_S_PiS_S_
	.cfi_endproc
.LFE2053:
	.size	_Z9addMatrixPiS_S_, .-_Z9addMatrixPiS_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"r"
.LC1:
	.string	"matA.txt"
.LC2:
	.string	"matB.txt"
.LC3:
	.string	"w"
.LC4:
	.string	"ans.txt"
.LC5:
	.string	"%d"
.LC6:
	.string	"%d "
.LC7:
	.string	"\n"
.LC9:
	.string	"Time spent for operation is %.10f seconds\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC1(%rip), %rdi
	xorl	%r15d, %r15d
	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
	leaq	.LC0(%rip), %rbx
	movq	%rbx, %rsi
	subq	$3192, %rsp
	.cfi_def_cfa_offset 3248
	movq	%fs:40, %rax
	movq	%rax, 3176(%rsp)
	xorl	%eax, %eax
	call	fopen@PLT
	movq	%rbx, %rsi
	leaq	.LC2(%rip), %rdi
	movq	%rax, 16(%rsp)
	call	fopen@PLT
	leaq	.LC3(%rip), %rsi
	leaq	.LC4(%rip), %rdi
	movq	%rax, 24(%rsp)
	call	fopen@PLT
	movq	%rax, 8(%rsp)
.L9:
	movq	%r15, %rbp
	movl	$16, %ebx
	salq	$6, %rbp
.L10:
	movq	16(%rsp), %rdi
	leaq	104(%rsp), %r14
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	leaq	(%r14,%rbp), %rdx
	leaq	1128(%rsp), %r13
	call	__isoc23_fscanf@PLT
	movq	24(%rsp), %rdi
	leaq	0(%r13,%rbp), %rdx
	xorl	%eax, %eax
	leaq	.LC5(%rip), %rsi
	leaq	2152(%rsp), %r12
	call	__isoc23_fscanf@PLT
	xorl	%eax, %eax
	movl	%eax, 2152(%rsp,%rbp)
	addq	$4, %rbp
	decl	%ebx
	jne	.L10
	incq	%r15
	cmpq	$16, %r15
	jne	.L9
	leaq	40(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	40(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	leaq	56(%rsp), %rdi
	movl	$1024, %esi
	call	cudaMalloc@PLT
	leaq	64(%rsp), %rdi
	movl	$1024, %esi
	call	cudaMalloc@PLT
	leaq	72(%rsp), %rdi
	movl	$1024, %esi
	call	cudaMalloc@PLT
	movq	56(%rsp), %rdi
	movl	$1, %ecx
	movq	%r14, %rsi
	movl	$1024, %edx
	call	cudaMemcpy@PLT
	movq	64(%rsp), %rdi
	movl	$1, %ecx
	movq	%r13, %rsi
	movl	$1024, %edx
	call	cudaMemcpy@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$4294967297, %rdi
	movl	$1, %esi
	movabsq	$68719476752, %rdx
	movl	$1, 88(%rsp)
	movq	%rdi, 80(%rsp)
	movq	%rdx, 92(%rsp)
	movl	$1, 100(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L12
	movq	64(%rsp), %rdx
	movq	56(%rsp), %rsi
	movq	72(%rsp), %rdi
	call	_Z32__device_stub__Z9addMatrixPiS_S_PiS_S_
.L12:
	movq	72(%rsp), %rsi
	movl	$2, %ecx
	movq	%r12, %rdi
	xorl	%r13d, %r13d
	movl	$1024, %edx
	leaq	.LC6(%rip), %r14
	call	cudaMemcpy@PLT
	leaq	48(%rsp), %rdi
	call	cudaEventCreate@PLT
	movq	48(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	48(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	64(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rdi
	call	cudaFree@PLT
.L13:
	movq	%r13, %r8
	xorl	%r15d, %r15d
	salq	$6, %r8
	leaq	(%r12,%r8), %rbx
.L14:
	movl	(%rbx,%r15,4), %ecx
	movq	8(%rsp), %rdi
	movq	%r14, %rdx
	xorl	%eax, %eax
	movl	$2, %esi
	incq	%r15
	call	__fprintf_chk@PLT
	cmpq	$16, %r15
	jne	.L14
	movq	8(%rsp), %rdi
	movl	$2, %esi
	xorl	%eax, %eax
	incq	%r13
	leaq	.LC7(%rip), %rdx
	call	__fprintf_chk@PLT
	cmpq	$16, %r13
	jne	.L13
	movq	48(%rsp), %rdx
	movq	40(%rsp), %rsi
	leaq	36(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	movss	36(%rsp), %xmm0
	divss	.LC8(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	16(%rsp), %rdi
	call	fclose@PLT
	movq	24(%rsp), %rdi
	call	fclose@PLT
	movq	8(%rsp), %rdi
	call	fclose@PLT
	movq	3176(%rsp), %rax
	subq	%fs:40, %rax
	je	.L16
	call	__stack_chk_fail@PLT
.L16:
	addq	$3192, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	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	main, .-main
	.section	.rodata.str1.1
.LC10:
	.string	"_Z9addMatrixPiS_S_"
	.section	.text.startup
	.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	.LC10(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z9addMatrixPiS_S_(%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
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC8:
	.long	1148846080
	.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 matrix addition example on CUDA
This program first reads two matrices from two text files namely matA.txt and matB.txt
Then it does matrix addition on the two matrices
Finally it saves the answer matrix in a text file named ans.txt
The program also measures the time taken for matrix operation on CUDA
*/


#include <hip/hip_runtime.h>
#include <stdio.h>

//define file names here
#define MATRIXA "matA.txt"
#define MATRIXB "matB.txt"
#define MATRIXANS "ans.txt"

//define the sizes of the matrices here
#define ROWS 16
#define COLS 16
#define SIZE ROWS*COLS

//kernel that does the matrix addition. Just add each element to the respective one
__global__ void addMatrix(int *ans_cuda,int *matA_cuda,int *matB_cuda){
	int row = threadIdx.y;
	int col = threadIdx.x;
	int position = row*COLS + col;
	ans_cuda[position]=matA_cuda[position]+matB_cuda[position];
}

int main(){
	
	//open the files 
	FILE *filematA = fopen(MATRIXA, "r");
	FILE *filematB = fopen(MATRIXB, "r");
	FILE *fileans = fopen(MATRIXANS, "w");
	
	//allocate matrices
	int matA[SIZE];
	int matB[SIZE];
	int ans[SIZE];
	
	//read the input matrices from file
	int row,col;
	for(row=0;row<ROWS;row++){
		for(col=0;col<COLS;col++){
			int position = row*COLS + col;
			fscanf(filematA, "%d", &matA[position]);
			fscanf(filematB, "%d", &matB[position]);
			ans[position]=0;
		}
	}
	

/*************************CUDA STUFF STARTS HERE************************/	
	
	//variables for time measurements
	hipEvent_t start,stop;
	float elapsedtime;
	
	//pointers for cuda memory locations
	int *matA_cuda;
	int *matB_cuda;
	int *ans_cuda;	

	//the moment at which we start measuring the time
	hipEventCreate(&start);
	hipEventRecord(start,0);
	
	//allocate memory in cuda
	hipMalloc((void **)&matA_cuda,sizeof(int)*SIZE);
	hipMalloc((void **)&matB_cuda,sizeof(int)*SIZE);	
	hipMalloc((void **)&ans_cuda,sizeof(int)*SIZE);
		
	//copy contents from ram to cuda
	hipMemcpy(matA_cuda, matA, sizeof(int)*SIZE, hipMemcpyHostToDevice);
	hipMemcpy(matB_cuda, matB, sizeof(int)*SIZE, hipMemcpyHostToDevice);	

	//thread configuration 
	dim3 numBlocks(1,1);
	dim3 threadsPerBlock(COLS,ROWS);
	
	//do the matrix addition on CUDA
	addMatrix<<<numBlocks,threadsPerBlock>>>(ans_cuda,matA_cuda,matB_cuda);

	//copy the answer back
	hipMemcpy(ans, ans_cuda, sizeof(int)*SIZE, hipMemcpyDeviceToHost);	

	//the moment at which we stop measuring time 
	hipEventCreate(&stop);
	hipEventRecord(stop,0);
	hipEventSynchronize(stop);
	
	//free the memory we allocated on CUDA
	hipFree(matA_cuda);
	hipFree(matB_cuda);
	hipFree(ans_cuda);
	
/*************************CUDA STUFF ENDS HERE************************/


	//write the answer to the text file
	for(row=0;row<ROWS;row++){
		for(col=0;col<COLS;col++){
			int position = row*COLS + col;
			fprintf(fileans, "%d ", ans[position]);
		}
		fprintf(fileans, "\n");
	}
	
	//Find and print the elapsed time
	hipEventElapsedTime(&elapsedtime,start,stop);
	printf("Time spent for operation is %.10f seconds\n",elapsedtime/(float)1000);

	fclose(filematA);
	fclose(filematB);
	fclose(fileans);
		
	return 0;
}
	.text
	.file	"matrixadd.hip"
	.globl	_Z24__device_stub__addMatrixPiS_S_ # -- Begin function _Z24__device_stub__addMatrixPiS_S_
	.type	_Z24__device_stub__addMatrixPiS_S_,@function
_Z24__device_stub__addMatrixPiS_S_:     # @_Z24__device_stub__addMatrixPiS_S_
	.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	24(%rsp), %rsi
	movq	%rdx, (%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	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	$_Z9addMatrixPiS_S_, %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_end0:
	.size	_Z24__device_stub__addMatrixPiS_S_, .Lfunc_end0-_Z24__device_stub__addMatrixPiS_S_
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI1_0:
	.long	0x447a0000                      # float 1000
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.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	$3160, %rsp                     # imm = 0xC58
	.cfi_def_cfa_offset 3216
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movl	$.L.str.2, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	movl	$.L.str.3, %edi
	movl	$.L.str.4, %esi
	callq	fopen
	movq	%rax, 40(%rsp)                  # 8-byte Spill
	leaq	80(%rsp), %r12
	leaq	1104(%rsp), %r13
	leaq	2128(%rsp), %rbp
	xorl	%eax, %eax
	movq	16(%rsp), %r14                  # 8-byte Reload
	movq	8(%rsp), %r15                   # 8-byte Reload
.LBB1_1:                                # %.preheader32
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_2 Depth 2
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	xorl	%ebx, %ebx
.LBB1_2:                                #   Parent Loop BB1_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	leaq	(%rbx,%rbp), %rdx
	movl	$.L.str.5, %esi
	movq	%r14, %rdi
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	leaq	(%rbx,%r13), %rdx
	movl	$.L.str.5, %esi
	movq	%r15, %rdi
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	movl	$0, (%r12,%rbx)
	addq	$4, %rbx
	cmpq	$64, %rbx
	jne	.LBB1_2
# %bb.3:                                #   in Loop: Header=BB1_1 Depth=1
	movq	48(%rsp), %rax                  # 8-byte Reload
	incq	%rax
	addq	$64, %r12
	addq	$64, %r13
	addq	$64, %rbp
	cmpq	$16, %rax
	jne	.LBB1_1
# %bb.4:
	leaq	64(%rsp), %r12
	movq	%r12, %rdi
	callq	hipEventCreate
	movq	(%r12), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	leaq	32(%rsp), %r12
	movl	$1024, %esi                     # imm = 0x400
	movq	%r12, %rdi
	callq	hipMalloc
	leaq	24(%rsp), %r13
	movl	$1024, %esi                     # imm = 0x400
	movq	%r13, %rdi
	callq	hipMalloc
	movq	%rsp, %rdi
	movl	$1024, %esi                     # imm = 0x400
	callq	hipMalloc
	movq	(%r12), %rdi
	leaq	2128(%rsp), %rsi
	movl	$1024, %edx                     # imm = 0x400
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r13), %rdi
	leaq	1104(%rsp), %rsi
	movl	$1024, %edx                     # imm = 0x400
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movabsq	$68719476752, %rdx              # imm = 0x1000000010
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_6
# %bb.5:
	movq	(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdx
	callq	_Z24__device_stub__addMatrixPiS_S_
.LBB1_6:
	movq	(%rsp), %rsi
	leaq	80(%rsp), %r12
	movl	$1024, %edx                     # imm = 0x400
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	leaq	56(%rsp), %r13
	movq	%r13, %rdi
	callq	hipEventCreate
	movq	(%r13), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%r13), %rdi
	callq	hipEventSynchronize
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	xorl	%r13d, %r13d
	movq	16(%rsp), %rbp                  # 8-byte Reload
	movq	8(%rsp), %r14                   # 8-byte Reload
	movq	40(%rsp), %r15                  # 8-byte Reload
.LBB1_7:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_8 Depth 2
	xorl	%ebx, %ebx
.LBB1_8:                                #   Parent Loop BB1_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%r12,%rbx,4), %edx
	movl	$.L.str.6, %esi
	movq	%r15, %rdi
	xorl	%eax, %eax
	callq	fprintf
	incq	%rbx
	cmpq	$16, %rbx
	jne	.LBB1_8
# %bb.9:                                #   in Loop: Header=BB1_7 Depth=1
	movl	$10, %edi
	movq	%r15, %rsi
	callq	fputc@PLT
	incq	%r13
	addq	$64, %r12
	cmpq	$16, %r13
	jne	.LBB1_7
# %bb.10:
	movq	64(%rsp), %rsi
	movq	56(%rsp), %rdx
	leaq	76(%rsp), %r12
	movq	%r12, %rdi
	callq	hipEventElapsedTime
	movss	(%r12), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	divss	.LCPI1_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.8, %edi
	movb	$1, %al
	callq	printf
	movq	%rbp, %rdi
	callq	fclose
	movq	%r14, %rdi
	callq	fclose
	movq	%r15, %rdi
	callq	fclose
	xorl	%eax, %eax
	addq	$3160, %rsp                     # imm = 0xC58
	.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_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:
	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	$_Z9addMatrixPiS_S_, %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	_Z9addMatrixPiS_S_,@object      # @_Z9addMatrixPiS_S_
	.section	.rodata,"a",@progbits
	.globl	_Z9addMatrixPiS_S_
	.p2align	3, 0x0
_Z9addMatrixPiS_S_:
	.quad	_Z24__device_stub__addMatrixPiS_S_
	.size	_Z9addMatrixPiS_S_, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"matA.txt"
	.size	.L.str, 9

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"r"
	.size	.L.str.1, 2

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"matB.txt"
	.size	.L.str.2, 9

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"ans.txt"
	.size	.L.str.3, 8

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"w"
	.size	.L.str.4, 2

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"%d"
	.size	.L.str.5, 3

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"%d "
	.size	.L.str.6, 4

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Time spent for operation is %.10f seconds\n"
	.size	.L.str.8, 43

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z9addMatrixPiS_S_"
	.size	.L__unnamed_1, 19

	.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 _Z24__device_stub__addMatrixPiS_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z9addMatrixPiS_S_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9addMatrixPiS_S_      ; -- Begin function _Z9addMatrixPiS_S_
	.globl	_Z9addMatrixPiS_S_
	.p2align	8
	.type	_Z9addMatrixPiS_S_,@function
_Z9addMatrixPiS_S_:                     ; @_Z9addMatrixPiS_S_
; %bb.0:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_lshrrev_b32_e32 v1, 6, v0
	v_and_b32_e32 v0, 0x3ff, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v1, 0x3ff0, v1
	v_add_lshl_u32 v0, v1, v0, 2
	s_waitcnt lgkmcnt(0)
	s_clause 0x1
	global_load_b32 v1, v0, s[6:7]
	global_load_b32 v2, v0, s[0:1]
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v1, v2, v1
	global_store_b32 v0, v1, s[4:5]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9addMatrixPiS_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.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 1
		.amdhsa_next_free_vgpr 3
		.amdhsa_next_free_sgpr 8
		.amdhsa_reserve_vcc 0
		.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	_Z9addMatrixPiS_S_, .Lfunc_end0-_Z9addMatrixPiS_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 104
; NumSgprs: 8
; NumVgprs: 3
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 0
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 8
; NumVGPRsForWavesPerEU: 3
; 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: 1
	.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
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9addMatrixPiS_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     8
    .sgpr_spill_count: 0
    .symbol:         _Z9addMatrixPiS_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     3
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
