//CS-4370 Parallel Programming for many core GPUs
//Name: Gesu Bal
/*
this is a simple cuda program calculating vector add for 2 dimensions on GPU device
I added two two-dimensional matrices A, B on the device GPU. 
After the device matrix addition kernel function is invoked, and the addition result is transferred back to the CPU. 
The program will also compute the  sum matrix of matrices A and B using the CPU.  
Then the program compares the device-computed result with the CPU-computed result. 
If it matches, it prints out Test PASSED to the screen before exiting.
*/
#include<stdio.h>
#include<cuda.h>
int N,blocksize;

//gpu function for addition
__global__ void add_gpu(int *d_a, int *d_b, int *d_c, int N)
{
     
     int row=blockIdx.y*blockDim.y+threadIdx.y;
     int col=blockIdx.x*blockDim.x+threadIdx.x;
     //int index =i+(j*N); 
     if((row <N) && (col <N))
     {
       
           d_c[row*N+col]=d_a[row*N+col]+d_b[row*N+col];
	   
      }
      
      
}

//cpu function for addition
void add_matrix_cpu(int *a, int *b, int *cpu_c, int N)
{ 
int i, j; 
for (i=0;i<N;i++) { 
       for (j=0;j<N;j++) {
             
             cpu_c[i*N+j]=a[i*N+j]+b[i*N+j];
     } 
  } 
} 

//match cpu and gpu results
int verify(int * a, int * b, int N)
{   
    int i,j;
    int error=0;
	for(i=0;i<N;i++)
	{
		for(j=0;j<N;j++)
		{
		    if(a[i*N+j]!=b[i*N+j])
		    {
		     error++; 
		    }
		}
	}
	
	if(error==0)
	{
	  printf("CPU and GPU results matched: Test Passed \n");
	}
	else
	{
	  printf("CPU and GPU results did not match");
	}
    return 1;
    
}

//print matrix fucntion
int printMatrix(int *a,int N)
{
  int i,j;
  for (i=0;i<N;i++)
    {
        for (j=0;j<N;j++)
        {
          printf("%d\t",a[i*N+j]);
        }
        printf("\n");
    }
return 1;
  
}

int main()
{
    //user input    
    int r, col;
	printf("Select one of the following options for vector addition: \n");
	printf("Press a for matrix size 8 * 8 \n");
	printf("Press b for matrix size 64 * 64 \n");
	printf("Press c for matrix size 128 * 128 \n");
	printf("Press d for matrix size 500 * 500 \n");
	printf("Press e for matrix size 1000 * 1000 \n");
    printf("Press any other key for exiting \n");
	char ch;
	scanf("%c",&ch);
	switch(ch)
        {
            case 'a':
                r=8;
		col=8;
		N=8;
		blocksize=4;
		printf("Array size is 8 * 8 \n");
		
                break;
            case 'b':
                r=64;
		col=64;
		N=64;
		blocksize=16;
		printf("Array size is 64 * 64 \n");
		
                break;
            case 'c':
                r=128;
		col=128;
		N=128;
		blocksize=16;
		printf("Array size is 128 * 128 \n");
		
                break;
	    case 'd':
                r=500;
		col=500;
		N=500;
		blocksize=16;
		printf("Array size is 500 * 500 \n");
		
                break;
            case 'e':
                r=1000;
		col=1000;
		N=1000;
		blocksize=16;
		printf("Array size is 1000 * 1000 \n");
		
                break;
	    default:
		exit(1);
                break;            
	}
  
    //vector initialization
	int *a, *b, *c, *cpu_c, *d_a, *d_b, *d_c;
	
	int a_size=r*col;
	int b_size=r*col;
	int c_size=r*col;
	int cpu_c_size=r*col;
	
    
    //memory allocation for vectors on host	
	a=(int*)malloc(sizeof(int)*a_size);
	b=(int*)malloc(sizeof(int)*b_size);
	c=(int*)malloc(sizeof(int)*c_size);
	cpu_c=(int*)malloc(sizeof(int)*cpu_c_size);
		
	
	//matrix initialization
    int i,j;
	int init=1325;
        for (i=0;i<N;i++)
	{
		for (j=0;j<N;j++)
		{
		    init=3125*init%65536;
		    a[i*col+j]=((init-32768)/16384);
		    b[i*col+j]=(init%1000);
		}
	}

	int cudaret=cudaMalloc((void **)(&d_a),(N*N)*sizeof(int));
	if(cudaret!=cudaSuccess)
	{printf("memory was not allocated on device \n");}
	
	cudaMalloc((void **)(&d_b),(N*N)*sizeof(int));
	cudaMalloc((void **)(&d_c),(N*N)*sizeof(int));
	

	//copying contents of a and b to device arrays
	cudaMemcpy(d_a,a,(N*N)*sizeof(int),cudaMemcpyHostToDevice);
    cudaMemcpy(d_b,b,(N*N)*sizeof(int),cudaMemcpyHostToDevice);
	
	//Initializing block count and block size
	dim3 dimBlock(blocksize,blocksize,1); 
	int blockCount_x = (N - 1)/(double(blocksize))+1;//Get number of blocks needed per direction.
	int blockCount_y = (N - 1)/(double(blocksize))+1;
	printf("the number of the thread blocks in x direction will be %d\n", blockCount_x);
	printf("the number of the thread blocks in y direction will be %d\n", blockCount_y);
	dim3 dimGrid(blockCount_x,blockCount_y,1);
        
	//calling CPU program
	printf("calculating results for CPU vector addition \n");
	printf("---------\n");
	add_matrix_cpu(a,b,cpu_c,N);
    
    //printMatrix(a,N);
    //pritnMatrix(b,N);
    //printMatrix(cpu_c,N);
	
	//call kernel for gpu functioning
	printf("calling kernel for gpu computations for vector addition \n");
	printf("---------\n");
	add_gpu<<<dimGrid,dimBlock>>>(d_a,d_b,d_c,N);
	printf("calculating results for gpu \n");
	printf("---------\n");
	
    //copying resulting back to cpu from gpu
    cudaMemcpy(c,d_c,(N*N)*sizeof(int),cudaMemcpyDeviceToHost);
        
	//matching cpu and gpu results
	printf("comparing results for CPU and GPU computations \n");
	printf("---------\n");
	verify(c,cpu_c,N);
	//printMatrix(c,N);
	
    //Deallocating memory
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
	
    return 0;
}

	.file	"tmpxft_00331690_00000000-6_GesuVecAdd.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2033:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z14add_matrix_cpuPiS_S_i
	.type	_Z14add_matrix_cpuPiS_S_i, @function
_Z14add_matrix_cpuPiS_S_i:
.LFB2027:
	.cfi_startproc
	endbr64
	movq	%rdi, %r9
	movq	%rsi, %r10
	movslq	%ecx, %rdi
	xorl	%esi, %esi
	salq	$2, %rdi
	xorl	%r8d, %r8d
	cmpl	%ecx, %esi
	jl	.L9
	ret
.L9:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
.L5:
	movq	%r8, %rax
	xorl	%r11d, %r11d
.L4:
	movl	(%r10,%rax), %ebx
	incl	%r11d
	addl	(%r9,%rax), %ebx
	movl	%ebx, (%rdx,%rax)
	addq	$4, %rax
	cmpl	%r11d, %ecx
	jne	.L4
	incl	%esi
	addq	%rdi, %r8
	cmpl	%ecx, %esi
	jl	.L5
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z14add_matrix_cpuPiS_S_i, .-_Z14add_matrix_cpuPiS_S_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"CPU and GPU results matched: Test Passed \n"
.LC1:
	.string	"CPU and GPU results did not match"
	.text
	.globl	_Z6verifyPiS_i
	.type	_Z6verifyPiS_i, @function
_Z6verifyPiS_i:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movslq	%edx, %r10
	xorl	%r9d, %r9d
	xorl	%ecx, %ecx
	salq	$2, %r10
	xorl	%r8d, %r8d
.L12:
	cmpl	%edx, %r8d
	jge	.L21
	movq	%r9, %rax
	xorl	%r11d, %r11d
.L14:
	movl	(%rsi,%rax), %ebx
	cmpl	%ebx, (%rdi,%rax)
	je	.L13
	incl	%ecx
.L13:
	incl	%r11d
	addq	$4, %rax
	cmpl	%r11d, %edx
	jne	.L14
	incl	%r8d
	addq	%r10, %r9
	jmp	.L12
.L21:
	leaq	.LC0(%rip), %rsi
	testl	%ecx, %ecx
	je	.L20
	leaq	.LC1(%rip), %rsi
.L20:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z6verifyPiS_i, .-_Z6verifyPiS_i
	.section	.rodata.str1.1
.LC2:
	.string	"%d\t"
.LC3:
	.string	"\n"
	.text
	.globl	_Z11printMatrixPii
	.type	_Z11printMatrixPii, @function
_Z11printMatrixPii:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	leaq	.LC2(%rip), %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movslq	%esi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%r13, %rbp
	salq	$2, %r13
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdi, %rbx
	pushq	%rcx
	.cfi_def_cfa_offset 64
.L23:
	cmpl	%ebp, %r12d
	jge	.L28
	leaq	(%rbx,%r13), %r15
.L24:
	movl	(%rbx), %edx
	movq	%r14, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	addq	$4, %rbx
	call	__printf_chk@PLT
	cmpq	%r15, %rbx
	jne	.L24
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incl	%r12d
	call	__printf_chk@PLT
	jmp	.L23
.L28:
	popq	%rdx
	.cfi_def_cfa_offset 56
	movl	$1, %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
.LFE2029:
	.size	_Z11printMatrixPii, .-_Z11printMatrixPii
	.globl	_Z31__device_stub__Z7add_gpuPiS_S_iPiS_S_i
	.type	_Z31__device_stub__Z7add_gpuPiS_S_iPiS_S_i, @function
_Z31__device_stub__Z7add_gpuPiS_S_iPiS_S_i:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L29
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z7add_gpuPiS_S_i(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L29:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L31
	call	__stack_chk_fail@PLT
.L31:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z31__device_stub__Z7add_gpuPiS_S_iPiS_S_i, .-_Z31__device_stub__Z7add_gpuPiS_S_iPiS_S_i
	.globl	_Z7add_gpuPiS_S_i
	.type	_Z7add_gpuPiS_S_i, @function
_Z7add_gpuPiS_S_i:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z7add_gpuPiS_S_iPiS_S_i
	.cfi_endproc
.LFE2056:
	.size	_Z7add_gpuPiS_S_i, .-_Z7add_gpuPiS_S_i
	.section	.rodata.str1.1
.LC4:
	.string	"Select one of the following options for vector addition: \n"
.LC5:
	.string	"Press a for matrix size 8 * 8 \n"
.LC6:
	.string	"Press b for matrix size 64 * 64 \n"
.LC7:
	.string	"Press c for matrix size 128 * 128 \n"
.LC8:
	.string	"Press d for matrix size 500 * 500 \n"
.LC9:
	.string	"Press e for matrix size 1000 * 1000 \n"
.LC10:
	.string	"Press any other key for exiting \n"
.LC11:
	.string	"%c"
.LC12:
	.string	"Array size is 8 * 8 \n"
.LC13:
	.string	"Array size is 64 * 64 \n"
.LC14:
	.string	"Array size is 128 * 128 \n"
.LC15:
	.string	"Array size is 500 * 500 \n"
.LC16:
	.string	"Array size is 1000 * 1000 \n"
.LC17:
	.string	"memory was not allocated on device \n"
.LC19:
	.string	"the number of the thread blocks in x direction will be %d\n"
.LC20:
	.string	"the number of the thread blocks in y direction will be %d\n"
.LC21:
	.string	"calculating results for CPU vector addition \n"
.LC22:
	.string	"---------\n"
.LC23:
	.string	"calling kernel for gpu computations for vector addition \n"
.LC24:
	.string	"calculating results for gpu \n"
.LC25:
	.string	"comparing results for CPU and GPU computations \n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	23(%rsp), %rsi
	leaq	.LC11(%rip), %rdi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movb	23(%rsp), %al
	subl	$97, %eax
	cmpb	$4, %al
	ja	.L35
	leaq	.L37(%rip), %rdx
	movzbl	%al, %eax
	movslq	(%rdx,%rax,4), %rax
	addq	%rdx, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L37:
	.long	.L41-.L37
	.long	.L40-.L37
	.long	.L39-.L37
	.long	.L38-.L37
	.long	.L36-.L37
	.section	.text.startup
.L41:
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$8, N(%rip)
	movl	$4, blocksize(%rip)
	movl	$8, %r12d
	call	__printf_chk@PLT
	jmp	.L42
.L40:
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$64, N(%rip)
	movl	$16, blocksize(%rip)
	movl	$64, %r12d
	call	__printf_chk@PLT
	jmp	.L42
.L39:
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$128, N(%rip)
	movl	$16, blocksize(%rip)
	movl	$128, %r12d
	call	__printf_chk@PLT
	jmp	.L42
.L38:
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$500, N(%rip)
	movl	$16, blocksize(%rip)
	movl	$500, %r12d
	call	__printf_chk@PLT
	jmp	.L42
.L36:
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$1000, N(%rip)
	movl	$16, blocksize(%rip)
	movl	$1000, %r12d
	call	__printf_chk@PLT
	jmp	.L42
.L35:
	movl	$1, %edi
	call	exit@PLT
.L42:
	movl	%r12d, %r13d
	movl	$1000, %r14d
	imull	%r12d, %r13d
	movslq	%r13d, %r13
	salq	$2, %r13
	movq	%r13, %rdi
	call	malloc@PLT
	movq	%r13, %rdi
	movq	%rax, %rbx
	call	malloc@PLT
	movq	%r13, %rdi
	movq	%rax, %rbp
	call	malloc@PLT
	movq	%r13, %rdi
	movl	$16384, %r13d
	movq	%rax, 8(%rsp)
	call	malloc@PLT
	movslq	%r12d, %rcx
	xorl	%r10d, %r10d
	movl	$1325, %edi
	movl	N(%rip), %esi
	movq	%rax, %r15
	salq	$2, %rcx
	xorl	%r9d, %r9d
	movl	$65536, %r12d
.L43:
	cmpl	%r9d, %esi
	jle	.L54
	movq	%r10, %r8
	xorl	%r11d, %r11d
.L44:
	imull	$3125, %edi, %eax
	incl	%r11d
	cltd
	idivl	%r12d
	leal	-32768(%rdx), %eax
	movl	%edx, %edi
	cltd
	idivl	%r13d
	movl	%eax, (%rbx,%r8)
	movl	%edi, %eax
	cltd
	idivl	%r14d
	movl	%edx, 0(%rbp,%r8)
	addq	$4, %r8
	cmpl	%r11d, %esi
	jne	.L44
	incl	%r9d
	addq	%rcx, %r10
	jmp	.L43
.L54:
	imull	%esi, %esi
	leaq	24(%rsp), %rdi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L45
	leaq	.LC17(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L45:
	movl	N(%rip), %esi
	leaq	32(%rsp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	N(%rip), %esi
	leaq	40(%rsp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	N(%rip), %edx
	movq	24(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	N(%rip), %edx
	movq	32(%rsp), %rdi
	movq	%rbp, %rsi
	movl	$1, %ecx
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	N(%rip), %ecx
	movl	blocksize(%rip), %eax
	leaq	.LC19(%rip), %rsi
	movl	$2, %edi
	movl	$1, 56(%rsp)
	cvtsi2sdl	%eax, %xmm1
	leal	-1(%rcx), %edx
	movl	%eax, 48(%rsp)
	cvtsi2sdl	%edx, %xmm0
	movl	%eax, 52(%rsp)
	xorl	%eax, %eax
	divsd	%xmm1, %xmm0
	addsd	.LC18(%rip), %xmm0
	cvttsd2sil	%xmm0, %r12d
	movl	%r12d, %edx
	call	__printf_chk@PLT
	movl	%r12d, %edx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC20(%rip), %rsi
	call	__printf_chk@PLT
	movl	%r12d, 60(%rsp)
	movl	$2, %edi
	xorl	%eax, %eax
	movl	%r12d, 64(%rsp)
	leaq	.LC21(%rip), %rsi
	leaq	.LC22(%rip), %r12
	movl	$1, 68(%rsp)
	call	__printf_chk@PLT
	movq	%r12, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r15, %rdx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	movl	N(%rip), %ecx
	call	_Z14add_matrix_cpuPiS_S_i
	leaq	.LC23(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r12, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	56(%rsp), %ecx
	movl	68(%rsp), %esi
	xorl	%r9d, %r9d
	movq	48(%rsp), %rdx
	movq	60(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L46
	movl	N(%rip), %ecx
	movq	40(%rsp), %rdx
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdi
	call	_Z31__device_stub__Z7add_gpuPiS_S_iPiS_S_i
.L46:
	leaq	.LC24(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r12, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	N(%rip), %edx
	movq	40(%rsp), %rsi
	movl	$2, %ecx
	movq	8(%rsp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	leaq	.LC25(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r12, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	N(%rip), %edx
	movq	8(%rsp), %rdi
	movq	%r15, %rsi
	call	_Z6verifyPiS_i
	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	.L47
	call	__stack_chk_fail@PLT
.L47:
	addq	$88, %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
.LFE2030:
	.size	main, .-main
	.section	.rodata.str1.1
.LC26:
	.string	"_Z7add_gpuPiS_S_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2058:
	.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	.LC26(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z7add_gpuPiS_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
.LFE2058:
	.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
	.globl	blocksize
	.bss
	.align 4
	.type	blocksize, @object
	.size	blocksize, 4
blocksize:
	.zero	4
	.globl	N
	.align 4
	.type	N, @object
	.size	N, 4
N:
	.zero	4
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC18:
	.long	0
	.long	1072693248
	.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:

//CS-4370 Parallel Programming for many core GPUs
//Name: Gesu Bal
/*
this is a simple cuda program calculating vector add for 2 dimensions on GPU device
I added two two-dimensional matrices A, B on the device GPU. 
After the device matrix addition kernel function is invoked, and the addition result is transferred back to the CPU. 
The program will also compute the  sum matrix of matrices A and B using the CPU.  
Then the program compares the device-computed result with the CPU-computed result. 
If it matches, it prints out Test PASSED to the screen before exiting.
*/
#include<stdio.h>
#include<hip/hip_runtime.h>
int N,blocksize;

//gpu function for addition
__global__ void add_gpu(int *d_a, int *d_b, int *d_c, int N)
{
     
     int row=blockIdx.y*blockDim.y+threadIdx.y;
     int col=blockIdx.x*blockDim.x+threadIdx.x;
     //int index =i+(j*N); 
     if((row <N) && (col <N))
     {
       
           d_c[row*N+col]=d_a[row*N+col]+d_b[row*N+col];
	   
      }
      
      
}

//cpu function for addition
void add_matrix_cpu(int *a, int *b, int *cpu_c, int N)
{ 
int i, j; 
for (i=0;i<N;i++) { 
       for (j=0;j<N;j++) {
             
             cpu_c[i*N+j]=a[i*N+j]+b[i*N+j];
     } 
  } 
} 

//match cpu and gpu results
int verify(int * a, int * b, int N)
{   
    int i,j;
    int error=0;
	for(i=0;i<N;i++)
	{
		for(j=0;j<N;j++)
		{
		    if(a[i*N+j]!=b[i*N+j])
		    {
		     error++; 
		    }
		}
	}
	
	if(error==0)
	{
	  printf("CPU and GPU results matched: Test Passed \n");
	}
	else
	{
	  printf("CPU and GPU results did not match");
	}
    return 1;
    
}

//print matrix fucntion
int printMatrix(int *a,int N)
{
  int i,j;
  for (i=0;i<N;i++)
    {
        for (j=0;j<N;j++)
        {
          printf("%d\t",a[i*N+j]);
        }
        printf("\n");
    }
return 1;
  
}

int main()
{
    //user input    
    int r, col;
	printf("Select one of the following options for vector addition: \n");
	printf("Press a for matrix size 8 * 8 \n");
	printf("Press b for matrix size 64 * 64 \n");
	printf("Press c for matrix size 128 * 128 \n");
	printf("Press d for matrix size 500 * 500 \n");
	printf("Press e for matrix size 1000 * 1000 \n");
    printf("Press any other key for exiting \n");
	char ch;
	scanf("%c",&ch);
	switch(ch)
        {
            case 'a':
                r=8;
		col=8;
		N=8;
		blocksize=4;
		printf("Array size is 8 * 8 \n");
		
                break;
            case 'b':
                r=64;
		col=64;
		N=64;
		blocksize=16;
		printf("Array size is 64 * 64 \n");
		
                break;
            case 'c':
                r=128;
		col=128;
		N=128;
		blocksize=16;
		printf("Array size is 128 * 128 \n");
		
                break;
	    case 'd':
                r=500;
		col=500;
		N=500;
		blocksize=16;
		printf("Array size is 500 * 500 \n");
		
                break;
            case 'e':
                r=1000;
		col=1000;
		N=1000;
		blocksize=16;
		printf("Array size is 1000 * 1000 \n");
		
                break;
	    default:
		exit(1);
                break;            
	}
  
    //vector initialization
	int *a, *b, *c, *cpu_c, *d_a, *d_b, *d_c;
	
	int a_size=r*col;
	int b_size=r*col;
	int c_size=r*col;
	int cpu_c_size=r*col;
	
    
    //memory allocation for vectors on host	
	a=(int*)malloc(sizeof(int)*a_size);
	b=(int*)malloc(sizeof(int)*b_size);
	c=(int*)malloc(sizeof(int)*c_size);
	cpu_c=(int*)malloc(sizeof(int)*cpu_c_size);
		
	
	//matrix initialization
    int i,j;
	int init=1325;
        for (i=0;i<N;i++)
	{
		for (j=0;j<N;j++)
		{
		    init=3125*init%65536;
		    a[i*col+j]=((init-32768)/16384);
		    b[i*col+j]=(init%1000);
		}
	}

	int cudaret=hipMalloc((void **)(&d_a),(N*N)*sizeof(int));
	if(cudaret!=hipSuccess)
	{printf("memory was not allocated on device \n");}
	
	hipMalloc((void **)(&d_b),(N*N)*sizeof(int));
	hipMalloc((void **)(&d_c),(N*N)*sizeof(int));
	

	//copying contents of a and b to device arrays
	hipMemcpy(d_a,a,(N*N)*sizeof(int),hipMemcpyHostToDevice);
    hipMemcpy(d_b,b,(N*N)*sizeof(int),hipMemcpyHostToDevice);
	
	//Initializing block count and block size
	dim3 dimBlock(blocksize,blocksize,1); 
	int blockCount_x = (N - 1)/(double(blocksize))+1;//Get number of blocks needed per direction.
	int blockCount_y = (N - 1)/(double(blocksize))+1;
	printf("the number of the thread blocks in x direction will be %d\n", blockCount_x);
	printf("the number of the thread blocks in y direction will be %d\n", blockCount_y);
	dim3 dimGrid(blockCount_x,blockCount_y,1);
        
	//calling CPU program
	printf("calculating results for CPU vector addition \n");
	printf("---------\n");
	add_matrix_cpu(a,b,cpu_c,N);
    
    //printMatrix(a,N);
    //pritnMatrix(b,N);
    //printMatrix(cpu_c,N);
	
	//call kernel for gpu functioning
	printf("calling kernel for gpu computations for vector addition \n");
	printf("---------\n");
	add_gpu<<<dimGrid,dimBlock>>>(d_a,d_b,d_c,N);
	printf("calculating results for gpu \n");
	printf("---------\n");
	
    //copying resulting back to cpu from gpu
    hipMemcpy(c,d_c,(N*N)*sizeof(int),hipMemcpyDeviceToHost);
        
	//matching cpu and gpu results
	printf("comparing results for CPU and GPU computations \n");
	printf("---------\n");
	verify(c,cpu_c,N);
	//printMatrix(c,N);
	
    //Deallocating memory
    hipFree(d_a);
    hipFree(d_b);
    hipFree(d_c);
	
    return 0;
}

	.text
	.file	"GesuVecAdd.hip"
	.globl	_Z22__device_stub__add_gpuPiS_S_i # -- Begin function _Z22__device_stub__add_gpuPiS_S_i
	.type	_Z22__device_stub__add_gpuPiS_S_i,@function
_Z22__device_stub__add_gpuPiS_S_i:      # @_Z22__device_stub__add_gpuPiS_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	$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), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%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	$_Z7add_gpuPiS_S_i, %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	_Z22__device_stub__add_gpuPiS_S_i, .Lfunc_end0-_Z22__device_stub__add_gpuPiS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z14add_matrix_cpuPiS_S_i       # -- Begin function _Z14add_matrix_cpuPiS_S_i
	.type	_Z14add_matrix_cpuPiS_S_i,@function
_Z14add_matrix_cpuPiS_S_i:              # @_Z14add_matrix_cpuPiS_S_i
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB1_6
# %bb.1:                                # %.preheader.lr.ph
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movl	%ecx, %eax
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
.LBB1_2:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_3 Depth 2
	movl	%r8d, %ebx
	leaq	(%rdx,%rbx,4), %r10
	leaq	(%rsi,%rbx,4), %r11
	leaq	(%rdi,%rbx,4), %rbx
	xorl	%r14d, %r14d
.LBB1_3:                                #   Parent Loop BB1_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%r11,%r14,4), %ebp
	addl	(%rbx,%r14,4), %ebp
	movl	%ebp, (%r10,%r14,4)
	incq	%r14
	cmpq	%r14, %rax
	jne	.LBB1_3
# %bb.4:                                # %._crit_edge
                                        #   in Loop: Header=BB1_2 Depth=1
	incq	%r9
	addl	%ecx, %r8d
	cmpq	%rax, %r9
	jne	.LBB1_2
# %bb.5:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %rbp
.LBB1_6:                                # %._crit_edge21
	retq
.Lfunc_end1:
	.size	_Z14add_matrix_cpuPiS_S_i, .Lfunc_end1-_Z14add_matrix_cpuPiS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z6verifyPiS_i                  # -- Begin function _Z6verifyPiS_i
	.type	_Z6verifyPiS_i,@function
_Z6verifyPiS_i:                         # @_Z6verifyPiS_i
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	testl	%edx, %edx
	jle	.LBB2_6
# %bb.1:                                # %.preheader.lr.ph
	movl	%edx, %eax
	xorl	%ecx, %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
.LBB2_2:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_3 Depth 2
	movl	%ecx, %r11d
	leaq	(%rsi,%r11,4), %r10
	leaq	(%rdi,%r11,4), %r11
	xorl	%ebx, %ebx
.LBB2_3:                                #   Parent Loop BB2_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%r11,%rbx,4), %ebp
	xorl	%r14d, %r14d
	cmpl	(%r10,%rbx,4), %ebp
	setne	%r14b
	addl	%r14d, %r8d
	incq	%rbx
	cmpq	%rbx, %rax
	jne	.LBB2_3
# %bb.4:                                # %._crit_edge
                                        #   in Loop: Header=BB2_2 Depth=1
	incq	%r9
	addl	%edx, %ecx
	cmpq	%rax, %r9
	jne	.LBB2_2
# %bb.5:                                # %._crit_edge22
	testl	%r8d, %r8d
	je	.LBB2_6
# %bb.7:
	movl	$.L.str.1, %edi
	xorl	%eax, %eax
	callq	printf
	jmp	.LBB2_8
.LBB2_6:                                # %._crit_edge22.thread
	movl	$.Lstr, %edi
	callq	puts@PLT
.LBB2_8:
	movl	$1, %eax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z6verifyPiS_i, .Lfunc_end2-_Z6verifyPiS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z11printMatrixPii              # -- Begin function _Z11printMatrixPii
	.type	_Z11printMatrixPii,@function
_Z11printMatrixPii:                     # @_Z11printMatrixPii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, (%rsp)                    # 8-byte Spill
	testl	%esi, %esi
	jle	.LBB3_5
# %bb.1:                                # %.preheader.lr.ph
	movl	%esi, %ebx
	movl	%esi, %r15d
	xorl	%ebp, %ebp
	xorl	%r12d, %r12d
.LBB3_2:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_3 Depth 2
	movl	%ebp, %eax
	movq	(%rsp), %rcx                    # 8-byte Reload
	leaq	(%rcx,%rax,4), %r13
	xorl	%r14d, %r14d
.LBB3_3:                                #   Parent Loop BB3_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%r13,%r14,4), %esi
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
	incq	%r14
	cmpq	%r14, %r15
	jne	.LBB3_3
# %bb.4:                                # %._crit_edge
                                        #   in Loop: Header=BB3_2 Depth=1
	movl	$10, %edi
	callq	putchar@PLT
	incq	%r12
	addl	%ebx, %ebp
	cmpq	%r15, %r12
	jne	.LBB3_2
.LBB3_5:                                # %._crit_edge13
	movl	$1, %eax
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z11printMatrixPii, .Lfunc_end3-_Z11printMatrixPii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI4_0:
	.quad	0x3ff0000000000000              # double 1
	.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	$56, %rsp
	.cfi_def_cfa_offset 112
	.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	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	movl	$.Lstr.5, %edi
	callq	puts@PLT
	movl	$.Lstr.6, %edi
	callq	puts@PLT
	movl	$.Lstr.7, %edi
	callq	puts@PLT
	leaq	15(%rsp), %rbx
	movl	$.L.str.11, %edi
	movq	%rbx, %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movb	(%rbx), %al
	addb	$-97, %al
	cmpb	$5, %al
	jae	.LBB4_16
# %bb.1:                                # %switch.lookup
	movzbl	%al, %eax
	movl	.Lswitch.table.main(,%rax,4), %ebx
	movl	.Lswitch.table.main.22(,%rax,4), %ecx
	movq	.Lswitch.table.main.23(,%rax,8), %rdi
	movl	%ebx, N(%rip)
	movl	%ecx, blocksize(%rip)
	callq	puts@PLT
	movl	%ebx, %r14d
	imull	%r14d, %r14d
	shll	$2, %r14d
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, %r15
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, %r12
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, %r14
	movl	N(%rip), %esi
	testl	%esi, %esi
	jle	.LBB4_6
# %bb.2:                                # %.preheader.lr.ph
	movl	$1325, %ecx                     # imm = 0x52D
	xorl	%eax, %eax
	movl	$65535, %edx                    # imm = 0xFFFF
	xorl	%edi, %edi
.LBB4_3:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_4 Depth 2
	movl	%eax, %r9d
	leaq	(%r12,%r9,4), %r8
	leaq	(%r15,%r9,4), %r9
	xorl	%r10d, %r10d
.LBB4_4:                                #   Parent Loop BB4_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	imull	$3125, %ecx, %r11d              # imm = 0xC35
	movl	%r11d, %ecx
	andl	%edx, %ecx
	movl	%r11d, %ebp
	xorl	$32768, %ebp                    # imm = 0x8000
	leal	-16385(%r11), %r13d
	testw	%bp, %bp
	cmovnsl	%ebp, %r13d
	movswl	%r13w, %ebp
	sarl	$14, %ebp
	movl	%ebp, (%r9,%r10,4)
	movl	%ecx, %ebp
	shrl	$3, %ebp
	imull	$8389, %ebp, %ebp               # imm = 0x20C5
	shrl	$20, %ebp
	imull	$1000, %ebp, %ebp               # imm = 0x3E8
	subl	%ebp, %r11d
	andl	%edx, %r11d
	movl	%r11d, (%r8,%r10,4)
	incq	%r10
	cmpq	%r10, %rsi
	jne	.LBB4_4
# %bb.5:                                # %._crit_edge
                                        #   in Loop: Header=BB4_3 Depth=1
	incq	%rdi
	addl	%ebx, %eax
	cmpq	%rsi, %rdi
	jne	.LBB4_3
.LBB4_6:                                # %._crit_edge77
	imull	%esi, %esi
	shlq	$2, %rsi
	leaq	24(%rsp), %rdi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB4_8
# %bb.7:
	movl	$.Lstr.13, %edi
	callq	puts@PLT
.LBB4_8:
	movl	N(%rip), %esi
	imull	%esi, %esi
	shlq	$2, %rsi
	leaq	32(%rsp), %r13
	movq	%r13, %rdi
	callq	hipMalloc
	movl	N(%rip), %esi
	imull	%esi, %esi
	shlq	$2, %rsi
	leaq	16(%rsp), %rdi
	callq	hipMalloc
	movq	24(%rsp), %rdi
	movl	N(%rip), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%r13), %rdi
	movl	N(%rip), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	blocksize(%rip), %eax
	movq	%rax, %rdx
	shlq	$32, %rdx
	movl	N(%rip), %ecx
	decl	%ecx
	cvtsi2sd	%ecx, %xmm0
	orq	%rax, %rdx
	movq	%rdx, 40(%rsp)                  # 8-byte Spill
	cvtsi2sd	%eax, %xmm1
	divsd	%xmm1, %xmm0
	addsd	.LCPI4_0(%rip), %xmm0
	cvttsd2si	%xmm0, %ebp
	xorl	%r13d, %r13d
	movl	$.L.str.18, %edi
	movl	%ebp, %esi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.19, %edi
	movl	%ebp, %esi
	xorl	%eax, %eax
	callq	printf
	movq	%rbp, %rbx
	shlq	$32, %rbx
	orq	%rbp, %rbx
	movl	$.Lstr.14, %edi
	callq	puts@PLT
	movl	$.Lstr.21, %edi
	callq	puts@PLT
	movl	N(%rip), %eax
	testl	%eax, %eax
	jle	.LBB4_13
# %bb.9:                                # %.preheader.lr.ph.i
	xorl	%ecx, %ecx
.LBB4_10:                               # %.lr.ph.i
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_11 Depth 2
	movl	%r13d, %edi
	leaq	(%r14,%rdi,4), %rdx
	leaq	(%r12,%rdi,4), %rsi
	leaq	(%r15,%rdi,4), %rdi
	xorl	%r8d, %r8d
.LBB4_11:                               #   Parent Loop BB4_10 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%rsi,%r8,4), %r9d
	addl	(%rdi,%r8,4), %r9d
	movl	%r9d, (%rdx,%r8,4)
	incq	%r8
	cmpq	%r8, %rax
	jne	.LBB4_11
# %bb.12:                               # %._crit_edge.i
                                        #   in Loop: Header=BB4_10 Depth=1
	incq	%rcx
	addl	%eax, %r13d
	cmpq	%rax, %rcx
	jne	.LBB4_10
.LBB4_13:                               # %_Z14add_matrix_cpuPiS_S_i.exit
	movl	$.Lstr.16, %edi
	callq	puts@PLT
	movl	$.Lstr.21, %edi
	callq	puts@PLT
	movq	%rbx, %rdi
	movl	$1, %esi
	movq	40(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_15
# %bb.14:
	movq	24(%rsp), %rdi
	movq	32(%rsp), %rsi
	movq	16(%rsp), %rdx
	movl	N(%rip), %ecx
	callq	_Z22__device_stub__add_gpuPiS_S_i
.LBB4_15:
	movl	$.Lstr.18, %edi
	callq	puts@PLT
	movl	$.Lstr.21, %edi
	callq	puts@PLT
	movq	16(%rsp), %rsi
	movl	N(%rip), %edx
	imull	%edx, %edx
	shlq	$2, %rdx
	movq	48(%rsp), %rbx                  # 8-byte Reload
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.Lstr.20, %edi
	callq	puts@PLT
	movl	$.Lstr.21, %edi
	callq	puts@PLT
	movl	N(%rip), %edx
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	_Z6verifyPiS_i
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$56, %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
.LBB4_16:
	.cfi_def_cfa_offset 112
	movl	$1, %edi
	callq	exit
.Lfunc_end4:
	.size	main, .Lfunc_end4-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	.LBB5_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB5_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z7add_gpuPiS_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_end5:
	.size	__hip_module_ctor, .Lfunc_end5-__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	.LBB6_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
.LBB6_2:
	retq
.Lfunc_end6:
	.size	__hip_module_dtor, .Lfunc_end6-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	N,@object                       # @N
	.bss
	.globl	N
	.p2align	2, 0x0
N:
	.long	0                               # 0x0
	.size	N, 4

	.type	blocksize,@object               # @blocksize
	.globl	blocksize
	.p2align	2, 0x0
blocksize:
	.long	0                               # 0x0
	.size	blocksize, 4

	.type	_Z7add_gpuPiS_S_i,@object       # @_Z7add_gpuPiS_S_i
	.section	.rodata,"a",@progbits
	.globl	_Z7add_gpuPiS_S_i
	.p2align	3, 0x0
_Z7add_gpuPiS_S_i:
	.quad	_Z22__device_stub__add_gpuPiS_S_i
	.size	_Z7add_gpuPiS_S_i, 8

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"CPU and GPU results did not match"
	.size	.L.str.1, 34

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"%d\t"
	.size	.L.str.2, 4

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"%c"
	.size	.L.str.11, 3

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"the number of the thread blocks in x direction will be %d\n"
	.size	.L.str.18, 59

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"the number of the thread blocks in y direction will be %d\n"
	.size	.L.str.19, 59

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z7add_gpuPiS_S_i"
	.size	.L__unnamed_1, 18

	.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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"CPU and GPU results matched: Test Passed "
	.size	.Lstr, 42

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"Select one of the following options for vector addition: "
	.size	.Lstr.1, 58

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"Press a for matrix size 8 * 8 "
	.size	.Lstr.2, 31

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"Press b for matrix size 64 * 64 "
	.size	.Lstr.3, 33

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"Press c for matrix size 128 * 128 "
	.size	.Lstr.4, 35

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"Press d for matrix size 500 * 500 "
	.size	.Lstr.5, 35

	.type	.Lstr.6,@object                 # @str.6
.Lstr.6:
	.asciz	"Press e for matrix size 1000 * 1000 "
	.size	.Lstr.6, 37

	.type	.Lstr.7,@object                 # @str.7
.Lstr.7:
	.asciz	"Press any other key for exiting "
	.size	.Lstr.7, 33

	.type	.Lstr.8,@object                 # @str.8
.Lstr.8:
	.asciz	"Array size is 1000 * 1000 "
	.size	.Lstr.8, 27

	.type	.Lstr.9,@object                 # @str.9
.Lstr.9:
	.asciz	"Array size is 500 * 500 "
	.size	.Lstr.9, 25

	.type	.Lstr.10,@object                # @str.10
.Lstr.10:
	.asciz	"Array size is 128 * 128 "
	.size	.Lstr.10, 25

	.type	.Lstr.11,@object                # @str.11
.Lstr.11:
	.asciz	"Array size is 64 * 64 "
	.size	.Lstr.11, 23

	.type	.Lstr.12,@object                # @str.12
.Lstr.12:
	.asciz	"Array size is 8 * 8 "
	.size	.Lstr.12, 21

	.type	.Lstr.13,@object                # @str.13
.Lstr.13:
	.asciz	"memory was not allocated on device "
	.size	.Lstr.13, 36

	.type	.Lstr.14,@object                # @str.14
.Lstr.14:
	.asciz	"calculating results for CPU vector addition "
	.size	.Lstr.14, 45

	.type	.Lstr.16,@object                # @str.16
.Lstr.16:
	.asciz	"calling kernel for gpu computations for vector addition "
	.size	.Lstr.16, 57

	.type	.Lstr.18,@object                # @str.18
.Lstr.18:
	.asciz	"calculating results for gpu "
	.size	.Lstr.18, 29

	.type	.Lstr.20,@object                # @str.20
.Lstr.20:
	.asciz	"comparing results for CPU and GPU computations "
	.size	.Lstr.20, 48

	.type	.Lstr.21,@object                # @str.21
.Lstr.21:
	.asciz	"---------"
	.size	.Lstr.21, 10

	.type	.Lswitch.table.main,@object     # @switch.table.main
	.section	.rodata,"a",@progbits
	.p2align	2, 0x0
.Lswitch.table.main:
	.long	8                               # 0x8
	.long	64                              # 0x40
	.long	128                             # 0x80
	.long	500                             # 0x1f4
	.long	1000                            # 0x3e8
	.size	.Lswitch.table.main, 20

	.type	.Lswitch.table.main.22,@object  # @switch.table.main.22
	.p2align	2, 0x0
.Lswitch.table.main.22:
	.long	4                               # 0x4
	.long	16                              # 0x10
	.long	16                              # 0x10
	.long	16                              # 0x10
	.long	16                              # 0x10
	.size	.Lswitch.table.main.22, 20

	.type	.Lswitch.table.main.23,@object  # @switch.table.main.23
	.p2align	3, 0x0
.Lswitch.table.main.23:
	.quad	.Lstr.12
	.quad	.Lstr.11
	.quad	.Lstr.10
	.quad	.Lstr.9
	.quad	.Lstr.8
	.size	.Lswitch.table.main.23, 40

	.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 _Z22__device_stub__add_gpuPiS_S_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z7add_gpuPiS_S_i
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z7add_gpuPiS_S_i       ; -- Begin function _Z7add_gpuPiS_S_i
	.globl	_Z7add_gpuPiS_S_i
	.p2align	8
	.type	_Z7add_gpuPiS_S_i,@function
_Z7add_gpuPiS_S_i:                      ; @_Z7add_gpuPiS_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x2c
	s_load_b32 s2, s[0:1], 0x18
	v_bfe_u32 v2, v0, 10, 10
	v_and_b32_e32 v3, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s4, s3, 16
	s_and_b32 s3, s3, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[0:1], null, s15, s4, v[2:3]
	v_mad_u64_u32 v[1:2], null, s14, s3, v[3:4]
	s_mov_b32 s3, exec_lo
	v_max_i32_e32 v2, v0, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e64 s2, v2
	s_cbranch_execz .LBB0_2
; %bb.1:
	s_load_b128 s[4:7], s[0:1], 0x0
	v_mad_u64_u32 v[2:3], null, v0, s2, v[1:2]
	s_load_b64 s[0:1], s[0:1], 0x10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[0:1], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s4, v0
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v1, vcc_lo
	v_add_co_u32 v4, vcc_lo, s6, v0
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v1, vcc_lo
	v_add_co_u32 v0, vcc_lo, s0, v0
	global_load_b32 v2, v[2:3], off
	global_load_b32 v3, v[4:5], off
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v2, v3, v2
	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 _Z7add_gpuPiS_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.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 6
		.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	_Z7add_gpuPiS_S_i, .Lfunc_end0-_Z7add_gpuPiS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 224
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; 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
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         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:           _Z7add_gpuPiS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z7add_gpuPiS_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
