//Parallel programming for many core GPUs
//Name: Gesu Bal
//Instructor name: Meilin Liu
/*
this is a simple cuda program calculating Tiled Matrix vector multiplication for 2 dimensions on GPU device
I multiplied two double two-dimensional matrices A, B on the device GPU. 
After the device matrix multiplication kernel function is invoked, and the multiplication result is transferred back to the CPU. 
The program will also compute the  multiplication 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 (within a certain tolerance, i.e., 0.000001), then it will print out "Test PASSED" to the screen before exiting.

This case is for all matrix sizes and blocksize/tilewidth: 8*8
*/

#include<stdio.h>
#include<cuda.h>
#include <time.h>
int N,blocksize;

//gpu function for multiplication
__global__ void mul_matrix_gpu(double *d_a, double *d_b, double *d_c, int width)
{  
    int TILE_WIDTH=8;
    __shared__ double ds_M[8][8];
    __shared__ double ds_N[8][8];

     int bx = blockIdx.x;  int by = blockIdx.y;
     int tx = threadIdx.x; int ty = threadIdx.y;

    // Identify the row and column of the Pd element to work on
    int Row = by * TILE_WIDTH + ty;
    int Col = bx * TILE_WIDTH + tx;
    double Pvalue = 0;
    // Loop over the Md and Nd tiles required to compute the Pd element
    for (int m = 0; m < width/TILE_WIDTH; m++) 
    {
    // Coolaborative loading of Md and Nd tiles into shared memory
        ds_M[ty][tx] = d_a[Row*width + m*TILE_WIDTH +tx];
        ds_N[ty][tx] = d_b[Col+(m*TILE_WIDTH+ty)*width];
        __syncthreads();
        for (int k = 0; k < TILE_WIDTH; k++)
             Pvalue += ds_M[ty][k] * ds_N[k][tx];
        __syncthreads();
    }   
     d_c[Row*width+Col] = Pvalue;
	 
	 
}

//cpu function for multiplication
void mul_matrix_cpu(double *a, double *b, double *cpu_c, int N)
{ 
int i, j,k; 
for (i=0;i<N;i++) { 
       for (j=0;j<N;j++) {
             double sum=0;
                for (k=0;k<N;k++) 
                {
                 double p=a[i*N+k];
                 double q=b[k*N+j];
                 sum=sum+(p*q);
                }
	    cpu_c[i*N+j]=sum;
     } 
  } 
} 

//cpu and gpu result matching function
bool verify(double *A, double *B, double *C, int  width) {
     const double relativeTolerance = 0.000001; 
     for(int row = 0; row < width; row++) {
        for(int col = 0; col < width; col++) {
            double sum = 0;
            for(unsigned int k = 0; k < width; k++) {
                    sum += A[row*width + k]*B[k*width + col];
            }
            double relativeError = (sum - C[row*width + col])/sum;
	    //printf("%f \t",relativeError);
	    //printf("\n");
            if (relativeError >= relativeTolerance
             || relativeError <= -relativeTolerance) 
            {
                    printf("TEST FAILED\n\n");
                    return false;
            }
           }
        }
      printf("TEST PASSED\n\n");
      return true; 
}

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


int main()
{
    //user input
  
    int r, col;
	printf("Select one of the following options: \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 512 * 512 \n");
	printf("Press e for matrix size 1024 * 1024 \n");
	printf("Press f for matrix size 4096 * 4096 \n");
	printf("Press any other key for exit \n");
	char ch;
	scanf("%c",&ch);
	switch(ch)
        {
            case 'a':
                r=8;
		col=8;
		N=8;
		printf("Matrix size is 8 * 8 \n");
		
		break;
            case 'b':
                r=64;
		col=64;
		N=64;
		
		printf("Matrix size is 64 * 64 \n");
		
		break;
            case 'c':
                r=128;
		col=128;
		N=128;
		
		printf("Matrix size is 128 * 128 \n");
		
		break;
	    case 'd':
                r=512;
		col=512;
		N=512;
		
		printf("Matrix size is 512 * 512 \n");
		
		break;
		
            case 'e':
                r=1024;
		col=1024;
		N=1024;
		
		printf("Matrix size is 1024 * 1024 \n");
		break;
		
	    case 'f':
		r=4096;
		col=4096;
		N=4096;
		
		printf("Matrix size is 4096 * 4096 \n");
		break;
	    
	    default:
		exit(1);
                break;            
	} 
	
	//initializing the block size/tile width
	blocksize=8;
       
	//memory allocation for vectors
	double *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;
		
	a=(double*)malloc(sizeof(double)*a_size);
	b=(double*)malloc(sizeof(double)*b_size);
	c=(double*)malloc(sizeof(double)*c_size);
	cpu_c=(double*)malloc(sizeof(double)*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.0)/16384.0;
		    init=3125*init%65536;
		    b[i*col+j]=(init-32768.0)/16384.0;
		}
	}
	//printMatrix(a,N);
	//printf("\n");
	//printMatrix(b,N);
	//printf("\n");
	
	//allocating memory on device
	int cudaret=cudaMalloc((void **)(&d_a),(N*N)*sizeof(double));
	if(cudaret!=cudaSuccess)
	{printf("memory was not allocated on device \n");}
	
	cudaMalloc((void **)(&d_b),(N*N)*sizeof(double));
	cudaMalloc((void **)(&d_c),(N*N)*sizeof(double));
	
	//calculating cpu time
	clock_t startCPU, end;
	float cpu_time_used;
	
	//calling CPU program
	printf("Calculating results for CPU vector multiplication \n");
	printf("---------\n");
	startCPU = clock();
	mul_matrix_cpu(a,b,cpu_c,N);
	end = clock();
	cpu_time_used = ((float) (end - startCPU))*1000;
	cpu_time_used= cpu_time_used/ CLOCKS_PER_SEC;
	
	printf("CPU computation time (milliseconds) \n");
	printf("%f \t",cpu_time_used);  
	printf("\n");
	printf("\n");
	//printMatrix(cpu_c,N);
	//printf("\n");
	//time execution calculation
	cudaEvent_t start,stop;
	float elapsedTime; 
	float timeTransfer;
	float timeBack;
   
	
	
	//memory transfer time
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start,0);

	//copying contents of a and b to device arrays
	cudaMemcpy(d_a,a,(N*N)*sizeof(double),cudaMemcpyHostToDevice);
	cudaMemcpy(d_b,b,(N*N)*sizeof(double),cudaMemcpyHostToDevice);
	
	cudaDeviceSynchronize(); 
	cudaEventRecord(stop,0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&timeTransfer,start,stop);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);
	
	
	
	//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;
	
	dim3 dimGrid(blockCount_x,blockCount_y,1);
	printf("Block size and tile width for the program is %d\n ",blocksize);
	//call kernel for gpu functioning
	printf("Calling kernel for gpu computations for vector multiplication and calculating results\n");
	printf("---------\n");
    
    
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start,0);
	mul_matrix_gpu<<<dimGrid,dimBlock>>>(d_a,d_b,d_c,N);
	cudaEventRecord(stop,0);
     	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime,start,stop);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);
	
	printf("GPU computation time (milliseconds) \n");
	printf("%f \t",elapsedTime);  
	printf("\n");
    
	
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start,0);
	//copying resulting back to cpu
	cudaMemcpy(c,d_c,(N*N)*sizeof(double),cudaMemcpyDeviceToHost);
	cudaDeviceSynchronize(); 
	cudaEventRecord(stop,0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&timeBack,start,stop);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);
	
	timeTransfer += timeBack;
	printf("Total Memory transfer time between CPU and GPU (milliseconds)\n");
	printf("%f \t",timeTransfer);  
	printf("\n");
	
	float speedup;
	speedup=cpu_time_used/elapsedTime;
	printf("Speedup: \n");
	printf("%f \t",speedup);  
	printf("\n");

	printf("Comparing results for CPU and GPU computations \n");
	printf("---------\n");
	verify(a,b,c,N);
	    
	//deallocating memory 
	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);
	
    return 0;
}

	.file	"tmpxft_002ec0ee_00000000-6_Tiled_Matrix_Mul.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	_Z14mul_matrix_cpuPdS_S_i
	.type	_Z14mul_matrix_cpuPdS_S_i, @function
_Z14mul_matrix_cpuPdS_S_i:
.LFB2027:
	.cfi_startproc
	endbr64
	movq	%rdi, %r10
	xorl	%r9d, %r9d
	movslq	%ecx, %rdi
	movq	%rsi, %r11
	salq	$3, %rdi
	xorl	%esi, %esi
	cmpl	%ecx, %r9d
	jl	.L11
	ret
.L11:
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
.L7:
	xorl	%eax, %eax
	leaq	(%r10,%rsi), %rbp
	leaq	(%rdx,%rsi), %r12
.L6:
	leaq	(%r11,%rax,8), %rbx
	xorl	%r8d, %r8d
	xorps	%xmm0, %xmm0
.L4:
	movsd	0(%rbp,%r8,8), %xmm1
	mulsd	(%rbx), %xmm1
	incq	%r8
	addq	%rdi, %rbx
	addsd	%xmm1, %xmm0
	cmpl	%r8d, %ecx
	jg	.L4
	movsd	%xmm0, (%r12,%rax,8)
	incq	%rax
	cmpl	%eax, %ecx
	jg	.L6
	incl	%r9d
	addq	%rdi, %rsi
	cmpl	%ecx, %r9d
	jl	.L7
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z14mul_matrix_cpuPdS_S_i, .-_Z14mul_matrix_cpuPdS_S_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"TEST FAILED\n\n"
.LC4:
	.string	"TEST PASSED\n\n"
	.text
	.globl	_Z6verifyPdS_S_i
	.type	_Z6verifyPdS_S_i, @function
_Z6verifyPdS_S_i:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movslq	%ecx, %r10
	movsd	.LC1(%rip), %xmm2
	movq	%rsi, %r11
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movsd	.LC2(%rip), %xmm3
	movq	%r10, %rax
	xorl	%r8d, %r8d
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	salq	$3, %r10
	xorl	%r9d, %r9d
.L14:
	cmpl	%eax, %r9d
	jge	.L21
	movl	%r8d, %ecx
	leaq	(%rdi,%rcx,8), %rbp
	xorl	%ecx, %ecx
.L22:
	movl	%ecx, %ebx
	xorl	%esi, %esi
	xorps	%xmm0, %xmm0
.L15:
	movl	%ebx, %r12d
	addl	%eax, %ebx
	movsd	(%r11,%r12,8), %xmm1
	mulsd	0(%rbp,%rsi,8), %xmm1
	incq	%rsi
	addsd	%xmm1, %xmm0
	cmpl	%eax, %esi
	jb	.L15
	movaps	%xmm0, %xmm1
	subsd	(%rdx,%rcx,8), %xmm1
	divsd	%xmm0, %xmm1
	comisd	%xmm2, %xmm1
	jnb	.L16
	comisd	%xmm1, %xmm3
	jb	.L24
.L16:
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%eax, %eax
	jmp	.L13
.L24:
	incq	%rcx
	cmpl	%ecx, %eax
	jg	.L22
	incl	%r9d
	addl	%eax, %r8d
	addq	%r10, %rdx
	jmp	.L14
.L21:
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movb	$1, %al
.L13:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z6verifyPdS_S_i, .-_Z6verifyPdS_S_i
	.section	.rodata.str1.1
.LC5:
	.string	"%f\t"
.LC6:
	.string	"\n"
	.text
	.globl	_Z11printMatrixPdi
	.type	_Z11printMatrixPdi, @function
_Z11printMatrixPdi:
.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	.LC5(%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	$3, %r13
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdi, %rbx
	pushq	%rcx
	.cfi_def_cfa_offset 64
.L27:
	cmpl	%ebp, %r12d
	jge	.L32
	leaq	(%rbx,%r13), %r15
.L28:
	movsd	(%rbx), %xmm0
	movq	%r14, %rsi
	movl	$2, %edi
	movb	$1, %al
	addq	$8, %rbx
	call	__printf_chk@PLT
	cmpq	%r15, %rbx
	jne	.L28
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incl	%r12d
	call	__printf_chk@PLT
	jmp	.L27
.L32:
	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	_Z11printMatrixPdi, .-_Z11printMatrixPdi
	.globl	_Z39__device_stub__Z14mul_matrix_gpuPdS_S_iPdS_S_i
	.type	_Z39__device_stub__Z14mul_matrix_gpuPdS_S_iPdS_S_i, @function
_Z39__device_stub__Z14mul_matrix_gpuPdS_S_iPdS_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	.L33
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z14mul_matrix_gpuPdS_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
.L33:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L35
	call	__stack_chk_fail@PLT
.L35:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z39__device_stub__Z14mul_matrix_gpuPdS_S_iPdS_S_i, .-_Z39__device_stub__Z14mul_matrix_gpuPdS_S_iPdS_S_i
	.globl	_Z14mul_matrix_gpuPdS_S_i
	.type	_Z14mul_matrix_gpuPdS_S_i, @function
_Z14mul_matrix_gpuPdS_S_i:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z14mul_matrix_gpuPdS_S_iPdS_S_i
	.cfi_endproc
.LFE2056:
	.size	_Z14mul_matrix_gpuPdS_S_i, .-_Z14mul_matrix_gpuPdS_S_i
	.section	.rodata.str1.1
.LC7:
	.string	"Select one of the following options: \n"
.LC8:
	.string	"Press a for matrix size 8 * 8 \n"
.LC9:
	.string	"Press b for matrix size 64 * 64 \n"
.LC10:
	.string	"Press c for matrix size 128 * 128 \n"
.LC11:
	.string	"Press d for matrix size 512 * 512 \n"
.LC12:
	.string	"Press e for matrix size 1024 * 1024 \n"
.LC13:
	.string	"Press f for matrix size 4096 * 4096 \n"
.LC14:
	.string	"Press any other key for exit \n"
.LC15:
	.string	"%c"
.LC16:
	.string	"Matrix size is 8 * 8 \n"
.LC17:
	.string	"Matrix size is 64 * 64 \n"
.LC18:
	.string	"Matrix size is 128 * 128 \n"
.LC19:
	.string	"Matrix size is 512 * 512 \n"
.LC20:
	.string	"Matrix size is 1024 * 1024 \n"
.LC21:
	.string	"Matrix size is 4096 * 4096 \n"
.LC24:
	.string	"memory was not allocated on device \n"
.LC25:
	.string	"Calculating results for CPU vector multiplication \n"
.LC26:
	.string	"---------\n"
.LC29:
	.string	"CPU computation time (milliseconds) \n"
.LC30:
	.string	"%f \t"
.LC32:
	.string	"Block size and tile width for the program is %d\n "
.LC33:
	.string	"Calling kernel for gpu computations for vector multiplication and calculating results\n"
.LC34:
	.string	"GPU computation time (milliseconds) \n"
.LC35:
	.string	"Total Memory transfer time between CPU and GPU (milliseconds)\n"
.LC36:
	.string	"Speedup: \n"
.LC37:
	.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	.LC7(%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	$120, %rsp
	.cfi_def_cfa_offset 176
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	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	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	27(%rsp), %rsi
	leaq	.LC15(%rip), %rdi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movb	27(%rsp), %al
	subl	$97, %eax
	cmpb	$5, %al
	ja	.L39
	leaq	.L41(%rip), %rdx
	movzbl	%al, %eax
	movslq	(%rdx,%rax,4), %rax
	addq	%rdx, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L41:
	.long	.L46-.L41
	.long	.L45-.L41
	.long	.L44-.L41
	.long	.L43-.L41
	.long	.L42-.L41
	.long	.L40-.L41
	.section	.text.startup
.L46:
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$8, N(%rip)
	call	__printf_chk@PLT
	movl	$8, %r13d
	jmp	.L47
.L45:
	leaq	.LC17(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$64, N(%rip)
	call	__printf_chk@PLT
	movl	$64, %r13d
	jmp	.L47
.L44:
	leaq	.LC18(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$128, N(%rip)
	call	__printf_chk@PLT
	movl	$128, %r13d
	jmp	.L47
.L43:
	leaq	.LC19(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$512, N(%rip)
	call	__printf_chk@PLT
	movl	$512, %r13d
	jmp	.L47
.L42:
	leaq	.LC20(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$1024, N(%rip)
	call	__printf_chk@PLT
	movl	$1024, %r13d
	jmp	.L47
.L40:
	leaq	.LC21(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$4096, N(%rip)
	call	__printf_chk@PLT
	movl	$4096, %r13d
	jmp	.L47
.L39:
	movl	$1, %edi
	call	exit@PLT
.L47:
	movl	$8, blocksize(%rip)
	movl	%r13d, %r12d
	imull	%r13d, %r12d
	salq	$3, %r13
	movslq	%r12d, %r12
	salq	$3, %r12
	movq	%r12, %rdi
	call	malloc@PLT
	movq	%r12, %rdi
	movq	%rax, %rbx
	call	malloc@PLT
	movq	%r12, %rdi
	movq	%rax, %rbp
	call	malloc@PLT
	movq	%r12, %rdi
	movq	%rax, (%rsp)
	call	malloc@PLT
	xorl	%r9d, %r9d
	movl	$1325, %edx
	xorl	%r8d, %r8d
	movl	N(%rip), %esi
	movsd	.LC22(%rip), %xmm1
	movq	%rax, %r12
	movl	$65536, %ecx
	movsd	.LC23(%rip), %xmm2
.L48:
	cmpl	%r8d, %esi
	jle	.L59
	movq	%r9, %rdi
	xorl	%r10d, %r10d
.L49:
	imull	$3125, %edx, %eax
	incl	%r10d
	cltd
	idivl	%ecx
	imull	$3125, %edx, %eax
	cvtsi2sdl	%edx, %xmm0
	cltd
	idivl	%ecx
	subsd	%xmm1, %xmm0
	mulsd	%xmm2, %xmm0
	movsd	%xmm0, (%rbx,%rdi)
	cvtsi2sdl	%edx, %xmm0
	subsd	%xmm1, %xmm0
	mulsd	%xmm2, %xmm0
	movsd	%xmm0, 0(%rbp,%rdi)
	addq	$8, %rdi
	cmpl	%r10d, %esi
	jne	.L49
	incl	%r8d
	addq	%r13, %r9
	jmp	.L48
.L59:
	imull	%esi, %esi
	leaq	40(%rsp), %rdi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L50
	leaq	.LC24(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L50:
	movl	N(%rip), %esi
	leaq	48(%rsp), %rdi
	leaq	64(%rsp), %r15
	leaq	72(%rsp), %r14
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	cudaMalloc@PLT
	movl	N(%rip), %esi
	leaq	56(%rsp), %rdi
	imull	%esi, %esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	call	cudaMalloc@PLT
	leaq	.LC25(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC26(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	call	clock@PLT
	movq	%r12, %rdx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	movl	N(%rip), %ecx
	movq	%rax, %r13
	leaq	.LC6(%rip), %r12
	call	_Z14mul_matrix_cpuPdS_S_i
	call	clock@PLT
	leaq	.LC29(%rip), %rsi
	movl	$2, %edi
	subq	%r13, %rax
	leaq	.LC30(%rip), %r13
	cvtsi2ssq	%rax, %xmm0
	mulss	.LC27(%rip), %xmm0
	xorl	%eax, %eax
	divss	.LC28(%rip), %xmm0
	movss	%xmm0, 12(%rsp)
	call	__printf_chk@PLT
	movq	%r13, %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	12(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	%r12, %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%r12, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r15, %rdi
	call	cudaEventCreate@PLT
	movq	%r14, %rdi
	call	cudaEventCreate@PLT
	movq	64(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	N(%rip), %edx
	movq	40(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movl	N(%rip), %edx
	movq	48(%rsp), %rdi
	movq	%rbp, %rsi
	movl	$1, %ecx
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movq	72(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	72(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	72(%rsp), %rdx
	movq	64(%rsp), %rsi
	leaq	32(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	64(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	72(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	N(%rip), %eax
	movl	blocksize(%rip), %edx
	leaq	.LC32(%rip), %rsi
	movl	$2, %edi
	movl	$1, 88(%rsp)
	decl	%eax
	cvtsi2sdl	%edx, %xmm1
	movl	%edx, 80(%rsp)
	cvtsi2sdl	%eax, %xmm0
	movl	%edx, 84(%rsp)
	movl	$1, 100(%rsp)
	divsd	%xmm1, %xmm0
	addsd	.LC31(%rip), %xmm0
	cvttsd2sil	%xmm0, %eax
	movl	%eax, 92(%rsp)
	movl	%eax, 96(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC33(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC26(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r15, %rdi
	call	cudaEventCreate@PLT
	movq	%r14, %rdi
	call	cudaEventCreate@PLT
	movq	64(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	88(%rsp), %ecx
	movl	100(%rsp), %esi
	xorl	%r9d, %r9d
	movq	80(%rsp), %rdx
	movq	92(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L51
	movl	N(%rip), %ecx
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	movq	40(%rsp), %rdi
	call	_Z39__device_stub__Z14mul_matrix_gpuPdS_S_iPdS_S_i
.L51:
	movq	72(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	72(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	72(%rsp), %rdx
	movq	64(%rsp), %rsi
	leaq	28(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	64(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	72(%rsp), %rdi
	call	cudaEventDestroy@PLT
	leaq	.LC34(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r13, %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	28(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	%r12, %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	%r15, %rdi
	call	cudaEventCreate@PLT
	movq	%r14, %rdi
	call	cudaEventCreate@PLT
	movq	64(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	N(%rip), %edx
	movq	56(%rsp), %rsi
	movl	$2, %ecx
	movq	(%rsp), %rdi
	imull	%edx, %edx
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movq	72(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	72(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	72(%rsp), %rdx
	movq	64(%rsp), %rsi
	leaq	36(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movq	64(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movq	72(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movss	32(%rsp), %xmm0
	addss	36(%rsp), %xmm0
	xorl	%eax, %eax
	leaq	.LC35(%rip), %rsi
	movl	$2, %edi
	movss	%xmm0, 32(%rsp)
	call	__printf_chk@PLT
	movq	%r13, %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	32(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	%r12, %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	leaq	.LC36(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movss	12(%rsp), %xmm0
	divss	28(%rsp), %xmm0
	movss	%xmm0, 12(%rsp)
	call	__printf_chk@PLT
	movss	12(%rsp), %xmm0
	movq	%r13, %rsi
	movb	$1, %al
	movl	$2, %edi
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	%r12, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC37(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC26(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	(%rsp), %rdx
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	movl	N(%rip), %ecx
	call	_Z6verifyPdS_S_i
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L52
	call	__stack_chk_fail@PLT
.L52:
	addq	$120, %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
.LC38:
	.string	"_Z14mul_matrix_gpuPdS_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	.LC38(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z14mul_matrix_gpuPdS_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
.LC1:
	.long	-1598689907
	.long	1051772663
	.align 8
.LC2:
	.long	-1598689907
	.long	-1095710985
	.align 8
.LC22:
	.long	0
	.long	1088421888
	.align 8
.LC23:
	.long	0
	.long	1058013184
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC27:
	.long	1148846080
	.align 4
.LC28:
	.long	1232348160
	.section	.rodata.cst8
	.align 8
.LC31:
	.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:

//Parallel programming for many core GPUs
//Name: Gesu Bal
//Instructor name: Meilin Liu
/*
this is a simple cuda program calculating Tiled Matrix vector multiplication for 2 dimensions on GPU device
I multiplied two double two-dimensional matrices A, B on the device GPU. 
After the device matrix multiplication kernel function is invoked, and the multiplication result is transferred back to the CPU. 
The program will also compute the  multiplication 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 (within a certain tolerance, i.e., 0.000001), then it will print out "Test PASSED" to the screen before exiting.

This case is for all matrix sizes and blocksize/tilewidth: 8*8
*/

#include<stdio.h>
#include<hip/hip_runtime.h>
#include <time.h>
int N,blocksize;

//gpu function for multiplication
__global__ void mul_matrix_gpu(double *d_a, double *d_b, double *d_c, int width)
{  
    int TILE_WIDTH=8;
    __shared__ double ds_M[8][8];
    __shared__ double ds_N[8][8];

     int bx = blockIdx.x;  int by = blockIdx.y;
     int tx = threadIdx.x; int ty = threadIdx.y;

    // Identify the row and column of the Pd element to work on
    int Row = by * TILE_WIDTH + ty;
    int Col = bx * TILE_WIDTH + tx;
    double Pvalue = 0;
    // Loop over the Md and Nd tiles required to compute the Pd element
    for (int m = 0; m < width/TILE_WIDTH; m++) 
    {
    // Coolaborative loading of Md and Nd tiles into shared memory
        ds_M[ty][tx] = d_a[Row*width + m*TILE_WIDTH +tx];
        ds_N[ty][tx] = d_b[Col+(m*TILE_WIDTH+ty)*width];
        __syncthreads();
        for (int k = 0; k < TILE_WIDTH; k++)
             Pvalue += ds_M[ty][k] * ds_N[k][tx];
        __syncthreads();
    }   
     d_c[Row*width+Col] = Pvalue;
	 
	 
}

//cpu function for multiplication
void mul_matrix_cpu(double *a, double *b, double *cpu_c, int N)
{ 
int i, j,k; 
for (i=0;i<N;i++) { 
       for (j=0;j<N;j++) {
             double sum=0;
                for (k=0;k<N;k++) 
                {
                 double p=a[i*N+k];
                 double q=b[k*N+j];
                 sum=sum+(p*q);
                }
	    cpu_c[i*N+j]=sum;
     } 
  } 
} 

//cpu and gpu result matching function
bool verify(double *A, double *B, double *C, int  width) {
     const double relativeTolerance = 0.000001; 
     for(int row = 0; row < width; row++) {
        for(int col = 0; col < width; col++) {
            double sum = 0;
            for(unsigned int k = 0; k < width; k++) {
                    sum += A[row*width + k]*B[k*width + col];
            }
            double relativeError = (sum - C[row*width + col])/sum;
	    //printf("%f \t",relativeError);
	    //printf("\n");
            if (relativeError >= relativeTolerance
             || relativeError <= -relativeTolerance) 
            {
                    printf("TEST FAILED\n\n");
                    return false;
            }
           }
        }
      printf("TEST PASSED\n\n");
      return true; 
}

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


int main()
{
    //user input
  
    int r, col;
	printf("Select one of the following options: \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 512 * 512 \n");
	printf("Press e for matrix size 1024 * 1024 \n");
	printf("Press f for matrix size 4096 * 4096 \n");
	printf("Press any other key for exit \n");
	char ch;
	scanf("%c",&ch);
	switch(ch)
        {
            case 'a':
                r=8;
		col=8;
		N=8;
		printf("Matrix size is 8 * 8 \n");
		
		break;
            case 'b':
                r=64;
		col=64;
		N=64;
		
		printf("Matrix size is 64 * 64 \n");
		
		break;
            case 'c':
                r=128;
		col=128;
		N=128;
		
		printf("Matrix size is 128 * 128 \n");
		
		break;
	    case 'd':
                r=512;
		col=512;
		N=512;
		
		printf("Matrix size is 512 * 512 \n");
		
		break;
		
            case 'e':
                r=1024;
		col=1024;
		N=1024;
		
		printf("Matrix size is 1024 * 1024 \n");
		break;
		
	    case 'f':
		r=4096;
		col=4096;
		N=4096;
		
		printf("Matrix size is 4096 * 4096 \n");
		break;
	    
	    default:
		exit(1);
                break;            
	} 
	
	//initializing the block size/tile width
	blocksize=8;
       
	//memory allocation for vectors
	double *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;
		
	a=(double*)malloc(sizeof(double)*a_size);
	b=(double*)malloc(sizeof(double)*b_size);
	c=(double*)malloc(sizeof(double)*c_size);
	cpu_c=(double*)malloc(sizeof(double)*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.0)/16384.0;
		    init=3125*init%65536;
		    b[i*col+j]=(init-32768.0)/16384.0;
		}
	}
	//printMatrix(a,N);
	//printf("\n");
	//printMatrix(b,N);
	//printf("\n");
	
	//allocating memory on device
	int cudaret=hipMalloc((void **)(&d_a),(N*N)*sizeof(double));
	if(cudaret!=hipSuccess)
	{printf("memory was not allocated on device \n");}
	
	hipMalloc((void **)(&d_b),(N*N)*sizeof(double));
	hipMalloc((void **)(&d_c),(N*N)*sizeof(double));
	
	//calculating cpu time
	clock_t startCPU, end;
	float cpu_time_used;
	
	//calling CPU program
	printf("Calculating results for CPU vector multiplication \n");
	printf("---------\n");
	startCPU = clock();
	mul_matrix_cpu(a,b,cpu_c,N);
	end = clock();
	cpu_time_used = ((float) (end - startCPU))*1000;
	cpu_time_used= cpu_time_used/ CLOCKS_PER_SEC;
	
	printf("CPU computation time (milliseconds) \n");
	printf("%f \t",cpu_time_used);  
	printf("\n");
	printf("\n");
	//printMatrix(cpu_c,N);
	//printf("\n");
	//time execution calculation
	hipEvent_t start,stop;
	float elapsedTime; 
	float timeTransfer;
	float timeBack;
   
	
	
	//memory transfer time
	hipEventCreate(&start);
	hipEventCreate(&stop);
	hipEventRecord(start,0);

	//copying contents of a and b to device arrays
	hipMemcpy(d_a,a,(N*N)*sizeof(double),hipMemcpyHostToDevice);
	hipMemcpy(d_b,b,(N*N)*sizeof(double),hipMemcpyHostToDevice);
	
	hipDeviceSynchronize(); 
	hipEventRecord(stop,0);
	hipEventSynchronize(stop);
	hipEventElapsedTime(&timeTransfer,start,stop);
	hipEventDestroy(start);
	hipEventDestroy(stop);
	
	
	
	//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;
	
	dim3 dimGrid(blockCount_x,blockCount_y,1);
	printf("Block size and tile width for the program is %d\n ",blocksize);
	//call kernel for gpu functioning
	printf("Calling kernel for gpu computations for vector multiplication and calculating results\n");
	printf("---------\n");
    
    
	hipEventCreate(&start);
	hipEventCreate(&stop);
	hipEventRecord(start,0);
	mul_matrix_gpu<<<dimGrid,dimBlock>>>(d_a,d_b,d_c,N);
	hipEventRecord(stop,0);
     	hipEventSynchronize(stop);
	hipEventElapsedTime(&elapsedTime,start,stop);
	hipEventDestroy(start);
	hipEventDestroy(stop);
	
	printf("GPU computation time (milliseconds) \n");
	printf("%f \t",elapsedTime);  
	printf("\n");
    
	
	hipEventCreate(&start);
	hipEventCreate(&stop);
	hipEventRecord(start,0);
	//copying resulting back to cpu
	hipMemcpy(c,d_c,(N*N)*sizeof(double),hipMemcpyDeviceToHost);
	hipDeviceSynchronize(); 
	hipEventRecord(stop,0);
	hipEventSynchronize(stop);
	hipEventElapsedTime(&timeBack,start,stop);
	hipEventDestroy(start);
	hipEventDestroy(stop);
	
	timeTransfer += timeBack;
	printf("Total Memory transfer time between CPU and GPU (milliseconds)\n");
	printf("%f \t",timeTransfer);  
	printf("\n");
	
	float speedup;
	speedup=cpu_time_used/elapsedTime;
	printf("Speedup: \n");
	printf("%f \t",speedup);  
	printf("\n");

	printf("Comparing results for CPU and GPU computations \n");
	printf("---------\n");
	verify(a,b,c,N);
	    
	//deallocating memory 
	hipFree(d_a);
	hipFree(d_b);
	hipFree(d_c);
	
    return 0;
}

	.text
	.file	"Tiled_Matrix_Mul.hip"
	.globl	_Z29__device_stub__mul_matrix_gpuPdS_S_i # -- Begin function _Z29__device_stub__mul_matrix_gpuPdS_S_i
	.type	_Z29__device_stub__mul_matrix_gpuPdS_S_i,@function
_Z29__device_stub__mul_matrix_gpuPdS_S_i: # @_Z29__device_stub__mul_matrix_gpuPdS_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	$_Z14mul_matrix_gpuPdS_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	_Z29__device_stub__mul_matrix_gpuPdS_S_i, .Lfunc_end0-_Z29__device_stub__mul_matrix_gpuPdS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z14mul_matrix_cpuPdS_S_i       # -- Begin function _Z14mul_matrix_cpuPdS_S_i
	.type	_Z14mul_matrix_cpuPdS_S_i,@function
_Z14mul_matrix_cpuPdS_S_i:              # @_Z14mul_matrix_cpuPdS_S_i
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB1_8
# %bb.1:                                # %.preheader28.lr.ph
	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
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%ecx, %eax
	leaq	(,%rax,8), %r8
	xorl	%r9d, %r9d
	xorl	%r10d, %r10d
.LBB1_2:                                # %.preheader.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_3 Depth 2
                                        #       Child Loop BB1_4 Depth 3
	movl	%r9d, %r11d
	leaq	(%rdi,%r11,8), %r11
	movq	%r10, %rbx
	imulq	%rax, %rbx
	leaq	(%rdx,%rbx,8), %rbx
	movq	%rsi, %r14
	xorl	%r15d, %r15d
.LBB1_3:                                # %.preheader
                                        #   Parent Loop BB1_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB1_4 Depth 3
	xorpd	%xmm0, %xmm0
	movq	%r14, %r12
	xorl	%r13d, %r13d
.LBB1_4:                                #   Parent Loop BB1_2 Depth=1
                                        #     Parent Loop BB1_3 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movsd	(%r11,%r13,8), %xmm1            # xmm1 = mem[0],zero
	mulsd	(%r12), %xmm1
	addsd	%xmm1, %xmm0
	incq	%r13
	addq	%r8, %r12
	cmpq	%r13, %rax
	jne	.LBB1_4
# %bb.5:                                # %._crit_edge
                                        #   in Loop: Header=BB1_3 Depth=2
	movsd	%xmm0, (%rbx,%r15,8)
	incq	%r15
	addq	$8, %r14
	cmpq	%rax, %r15
	jne	.LBB1_3
# %bb.6:                                # %._crit_edge32
                                        #   in Loop: Header=BB1_2 Depth=1
	incq	%r10
	addl	%ecx, %r9d
	cmpq	%rax, %r10
	jne	.LBB1_2
# %bb.7:
	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
	.cfi_restore %rbx
	.cfi_restore %r12
	.cfi_restore %r13
	.cfi_restore %r14
	.cfi_restore %r15
.LBB1_8:                                # %._crit_edge34
	retq
.Lfunc_end1:
	.size	_Z14mul_matrix_cpuPdS_S_i, .Lfunc_end1-_Z14mul_matrix_cpuPdS_S_i
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z6verifyPdS_S_i
.LCPI2_0:
	.quad	0x3eb0c6f7a0b5ed8d              # double 9.9999999999999995E-7
.LCPI2_1:
	.quad	0xbeb0c6f7a0b5ed8d              # double -9.9999999999999995E-7
	.text
	.globl	_Z6verifyPdS_S_i
	.type	_Z6verifyPdS_S_i,@function
_Z6verifyPdS_S_i:                       # @_Z6verifyPdS_S_i
	.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, %rax
	movl	$.Lstr.1, %edi
	testl	%ecx, %ecx
	jle	.LBB2_1
# %bb.2:                                # %.preheader47.lr.ph
	movl	%ecx, %r8d
	xorl	%r9d, %r9d
	movsd	.LCPI2_0(%rip), %xmm0           # xmm0 = mem[0],zero
	movsd	.LCPI2_1(%rip), %xmm1           # xmm1 = mem[0],zero
	xorl	%r10d, %r10d
	xorl	%ebx, %ebx
.LBB2_3:                                # %.preheader.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_6 Depth 2
                                        #       Child Loop BB2_7 Depth 3
	movl	%r10d, %r11d
	imull	%ecx, %r11d
	leaq	(%rdx,%r11,8), %r11
	xorl	%r14d, %r14d
.LBB2_6:                                # %.preheader
                                        #   Parent Loop BB2_3 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB2_7 Depth 3
	xorpd	%xmm2, %xmm2
	movq	%r8, %r15
	movl	%r9d, %ebp
	movl	%r14d, %r12d
.LBB2_7:                                #   Parent Loop BB2_3 Depth=1
                                        #     Parent Loop BB2_6 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movl	%ebp, %r13d
	movsd	(%rax,%r13,8), %xmm3            # xmm3 = mem[0],zero
	movl	%r12d, %r13d
	mulsd	(%rsi,%r13,8), %xmm3
	addsd	%xmm3, %xmm2
	addl	%ecx, %r12d
	incl	%ebp
	decq	%r15
	jne	.LBB2_7
# %bb.8:                                # %.critedge
                                        #   in Loop: Header=BB2_6 Depth=2
	movapd	%xmm2, %xmm3
	subsd	(%r11,%r14,8), %xmm3
	divsd	%xmm2, %xmm3
	ucomisd	%xmm0, %xmm3
	jae	.LBB2_9
# %bb.4:                                # %.critedge
                                        #   in Loop: Header=BB2_6 Depth=2
	ucomisd	%xmm3, %xmm1
	jae	.LBB2_9
# %bb.5:                                #   in Loop: Header=BB2_6 Depth=2
	incq	%r14
	cmpq	%r8, %r14
	jne	.LBB2_6
# %bb.10:                               # %.critedge44
                                        #   in Loop: Header=BB2_3 Depth=1
	incq	%r10
	cmpq	%r8, %r10
	setae	%bl
	addl	%ecx, %r9d
	cmpq	%r8, %r10
	jne	.LBB2_3
	jmp	.LBB2_11
.LBB2_9:
	movl	$.Lstr, %edi
	jmp	.LBB2_11
.LBB2_1:
	movb	$1, %bl
.LBB2_11:                               # %.critedge46
	callq	puts@PLT
	andb	$1, %bl
	movl	%ebx, %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_end2:
	.size	_Z6verifyPdS_S_i, .Lfunc_end2-_Z6verifyPdS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z11printMatrixPdi              # -- Begin function _Z11printMatrixPdi
	.type	_Z11printMatrixPdi,@function
_Z11printMatrixPdi:                     # @_Z11printMatrixPdi
	.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	%r12d, %r12d
	xorl	%r13d, %r13d
.LBB3_2:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_3 Depth 2
	movl	%r12d, %eax
	movq	(%rsp), %rcx                    # 8-byte Reload
	leaq	(%rcx,%rax,8), %rbp
	xorl	%r14d, %r14d
.LBB3_3:                                #   Parent Loop BB3_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rbp,%r14,8), %xmm0            # xmm0 = mem[0],zero
	movl	$.L.str.2, %edi
	movb	$1, %al
	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	%r13
	addl	%ebx, %r12d
	cmpq	%r15, %r13
	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	_Z11printMatrixPdi, .Lfunc_end3-_Z11printMatrixPdi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI4_0:
	.quad	0x3f10000000000000              # double 6.103515625E-5
.LCPI4_3:
	.quad	0x3ff0000000000000              # double 1
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI4_1:
	.long	0x447a0000                      # float 1000
.LCPI4_2:
	.long	0x49742400                      # float 1.0E+6
	.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	$88, %rsp
	.cfi_def_cfa_offset 144
	.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.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
	movl	$.Lstr.8, %edi
	callq	puts@PLT
	movl	$.Lstr.9, %edi
	callq	puts@PLT
	leaq	23(%rsp), %rbx
	movl	$.L.str.12, %edi
	movq	%rbx, %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movb	(%rbx), %al
	addb	$-97, %al
	cmpb	$6, %al
	jae	.LBB4_11
# %bb.1:                                # %switch.lookup
	movzbl	%al, %eax
	movl	.Lswitch.table.main(,%rax,4), %ebp
	movq	.Lswitch.table.main.27(,%rax,8), %rdi
	movl	%ebp, N(%rip)
	callq	puts@PLT
	movl	$8, blocksize(%rip)
	movl	%ebp, %r15d
	imull	%r15d, %r15d
	shll	$3, %r15d
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %rbx
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, %r14
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, 56(%rsp)                  # 8-byte Spill
	movl	N(%rip), %esi
	testl	%esi, %esi
	jle	.LBB4_6
# %bb.2:                                # %.preheader.lr.ph
	movl	$1325, %eax                     # imm = 0x52D
	xorl	%ecx, %ecx
	movl	$65535, %edx                    # imm = 0xFFFF
	movl	$-32768, %edi                   # imm = 0x8000
	movsd	.LCPI4_0(%rip), %xmm0           # xmm0 = mem[0],zero
	xorl	%r8d, %r8d
.LBB4_3:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_4 Depth 2
	movl	%ecx, %r10d
	leaq	(%r14,%r10,8), %r9
	leaq	(%rbx,%r10,8), %r10
	xorl	%r11d, %r11d
.LBB4_4:                                #   Parent Loop BB4_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	imull	$3125, %eax, %r15d              # imm = 0xC35
	andl	%edx, %r15d
	addl	%edi, %r15d
	xorps	%xmm1, %xmm1
	cvtsi2sd	%r15d, %xmm1
	mulsd	%xmm0, %xmm1
	movsd	%xmm1, (%r10,%r11,8)
	imull	$761, %eax, %eax                # imm = 0x2F9
	andl	%edx, %eax
	leal	(%rax,%rdi), %r15d
	xorps	%xmm1, %xmm1
	cvtsi2sd	%r15d, %xmm1
	mulsd	%xmm0, %xmm1
	movsd	%xmm1, (%r9,%r11,8)
	incq	%r11
	cmpq	%r11, %rsi
	jne	.LBB4_4
# %bb.5:                                # %._crit_edge
                                        #   in Loop: Header=BB4_3 Depth=1
	incq	%r8
	addl	%ebp, %ecx
	cmpq	%rsi, %r8
	jne	.LBB4_3
.LBB4_6:                                # %._crit_edge92
	imull	%esi, %esi
	shlq	$3, %rsi
	leaq	32(%rsp), %rdi
	callq	hipMalloc
	testl	%eax, %eax
	je	.LBB4_8
# %bb.7:
	movl	$.Lstr.16, %edi
	callq	puts@PLT
.LBB4_8:                                # %_Z14mul_matrix_cpuPdS_S_i.exit
	movl	N(%rip), %esi
	imull	%esi, %esi
	shlq	$3, %rsi
	leaq	48(%rsp), %rbp
	movq	%rbp, %rdi
	callq	hipMalloc
	movl	N(%rip), %esi
	imull	%esi, %esi
	shlq	$3, %rsi
	leaq	24(%rsp), %rdi
	callq	hipMalloc
	movl	$.Lstr.17, %edi
	callq	puts@PLT
	movl	$.Lstr.26, %edi
	callq	puts@PLT
	callq	clock
	movq	%rax, %r15
	callq	clock
	subq	%r15, %rax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rax, %xmm0
	mulss	.LCPI4_1(%rip), %xmm0
	divss	.LCPI4_2(%rip), %xmm0
	movss	%xmm0, 12(%rsp)                 # 4-byte Spill
	movl	$.Lstr.19, %edi
	callq	puts@PLT
	xorps	%xmm0, %xmm0
	cvtss2sd	12(%rsp), %xmm0                 # 4-byte Folded Reload
	movl	$.L.str.23, %edi
	movb	$1, %al
	callq	printf
	movl	$10, %edi
	callq	putchar@PLT
	movl	$10, %edi
	callq	putchar@PLT
	leaq	72(%rsp), %r12
	movq	%r12, %rdi
	callq	hipEventCreate
	leaq	64(%rsp), %r13
	movq	%r13, %rdi
	callq	hipEventCreate
	movq	(%r12), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	32(%rsp), %rdi
	movl	N(%rip), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbp), %rdi
	movl	N(%rip), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	callq	hipDeviceSynchronize
	movq	(%r13), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%r13), %rdi
	callq	hipEventSynchronize
	movq	(%r12), %rsi
	movq	(%r13), %rdx
	leaq	16(%rsp), %rdi
	callq	hipEventElapsedTime
	movq	(%r12), %rdi
	callq	hipEventDestroy
	movq	(%r13), %rdi
	callq	hipEventDestroy
	movl	blocksize(%rip), %esi
	movq	%rsi, %rbp
	movl	N(%rip), %eax
	decl	%eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	shlq	$32, %rbp
	cvtsi2sd	%esi, %xmm1
	divsd	%xmm1, %xmm0
	addsd	.LCPI4_3(%rip), %xmm0
	cvttsd2si	%xmm0, %eax
	orq	%rsi, %rbp
	movq	%rax, %r15
	shlq	$32, %r15
	orq	%rax, %r15
	movl	$.L.str.24, %edi
                                        # kill: def $esi killed $esi killed $rsi
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr.20, %edi
	callq	puts@PLT
	movl	$.Lstr.26, %edi
	callq	puts@PLT
	movq	%r12, %rdi
	callq	hipEventCreate
	movq	%r13, %rdi
	callq	hipEventCreate
	movq	(%r12), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%rbp, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_10
# %bb.9:
	movq	32(%rsp), %rdi
	movq	48(%rsp), %rsi
	movq	24(%rsp), %rdx
	movl	N(%rip), %ecx
	callq	_Z29__device_stub__mul_matrix_gpuPdS_S_i
.LBB4_10:
	leaq	64(%rsp), %r13
	movq	(%r13), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%r13), %rdi
	callq	hipEventSynchronize
	leaq	72(%rsp), %rbp
	movq	(%rbp), %rsi
	movq	(%r13), %rdx
	leaq	44(%rsp), %rdi
	movq	%rdi, %r15
	callq	hipEventElapsedTime
	movq	(%rbp), %rdi
	callq	hipEventDestroy
	movq	(%r13), %rdi
	callq	hipEventDestroy
	movl	$.Lstr.22, %edi
	callq	puts@PLT
	xorps	%xmm0, %xmm0
	cvtss2sd	(%r15), %xmm0
	movl	$.L.str.23, %edi
	movb	$1, %al
	callq	printf
	movl	$10, %edi
	callq	putchar@PLT
	movq	%rbp, %rdi
	callq	hipEventCreate
	movq	%r13, %rdi
	callq	hipEventCreate
	movq	(%rbp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	24(%rsp), %rsi
	movl	N(%rip), %edx
	imull	%edx, %edx
	shlq	$3, %rdx
	movq	56(%rsp), %r12                  # 8-byte Reload
	movq	%r12, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	callq	hipDeviceSynchronize
	movq	(%r13), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%r13), %rdi
	callq	hipEventSynchronize
	movq	(%rbp), %rsi
	movq	(%r13), %rdx
	leaq	84(%rsp), %r15
	movq	%r15, %rdi
	callq	hipEventElapsedTime
	movq	(%rbp), %rdi
	callq	hipEventDestroy
	movq	(%r13), %rdi
	callq	hipEventDestroy
	movss	(%r15), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	addss	16(%rsp), %xmm0
	movss	%xmm0, 16(%rsp)
	movl	$.Lstr.23, %edi
	callq	puts@PLT
	xorps	%xmm0, %xmm0
	cvtss2sd	16(%rsp), %xmm0
	movl	$.L.str.23, %edi
	movb	$1, %al
	callq	printf
	movl	$10, %edi
	callq	putchar@PLT
	leaq	44(%rsp), %rax
	movss	12(%rsp), %xmm0                 # 4-byte Reload
                                        # xmm0 = mem[0],zero,zero,zero
	divss	(%rax), %xmm0
	movss	%xmm0, 12(%rsp)                 # 4-byte Spill
	movl	$.Lstr.24, %edi
	callq	puts@PLT
	xorps	%xmm0, %xmm0
	cvtss2sd	12(%rsp), %xmm0                 # 4-byte Folded Reload
	movl	$.L.str.23, %edi
	movb	$1, %al
	callq	printf
	movl	$10, %edi
	callq	putchar@PLT
	movl	$.Lstr.25, %edi
	callq	puts@PLT
	movl	$.Lstr.26, %edi
	callq	puts@PLT
	movl	N(%rip), %ecx
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r12, %rdx
	callq	_Z6verifyPdS_S_i
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$88, %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_11:
	.cfi_def_cfa_offset 144
	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	$_Z14mul_matrix_gpuPdS_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	_Z14mul_matrix_gpuPdS_S_i,@object # @_Z14mul_matrix_gpuPdS_S_i
	.section	.rodata,"a",@progbits
	.globl	_Z14mul_matrix_gpuPdS_S_i
	.p2align	3, 0x0
_Z14mul_matrix_gpuPdS_S_i:
	.quad	_Z29__device_stub__mul_matrix_gpuPdS_S_i
	.size	_Z14mul_matrix_gpuPdS_S_i, 8

	.type	.L.str.2,@object                # @.str.2
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.2:
	.asciz	"%f\t"
	.size	.L.str.2, 4

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

	.type	.L.str.23,@object               # @.str.23
.L.str.23:
	.asciz	"%f \t"
	.size	.L.str.23, 5

	.type	.L.str.24,@object               # @.str.24
.L.str.24:
	.asciz	"Block size and tile width for the program is %d\n "
	.size	.L.str.24, 50

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z14mul_matrix_gpuPdS_S_i"
	.size	.L__unnamed_1, 26

	.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	"TEST FAILED\n"
	.size	.Lstr, 13

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"TEST PASSED\n"
	.size	.Lstr.1, 13

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"Select one of the following options: "
	.size	.Lstr.2, 38

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

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

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

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

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

	.type	.Lstr.8,@object                 # @str.8
.Lstr.8:
	.asciz	"Press f for matrix size 4096 * 4096 "
	.size	.Lstr.8, 37

	.type	.Lstr.9,@object                 # @str.9
.Lstr.9:
	.asciz	"Press any other key for exit "
	.size	.Lstr.9, 30

	.type	.Lstr.10,@object                # @str.10
.Lstr.10:
	.asciz	"Matrix size is 4096 * 4096 "
	.size	.Lstr.10, 28

	.type	.Lstr.11,@object                # @str.11
.Lstr.11:
	.asciz	"Matrix size is 1024 * 1024 "
	.size	.Lstr.11, 28

	.type	.Lstr.12,@object                # @str.12
.Lstr.12:
	.asciz	"Matrix size is 512 * 512 "
	.size	.Lstr.12, 26

	.type	.Lstr.13,@object                # @str.13
.Lstr.13:
	.asciz	"Matrix size is 128 * 128 "
	.size	.Lstr.13, 26

	.type	.Lstr.14,@object                # @str.14
.Lstr.14:
	.asciz	"Matrix size is 64 * 64 "
	.size	.Lstr.14, 24

	.type	.Lstr.15,@object                # @str.15
.Lstr.15:
	.asciz	"Matrix size is 8 * 8 "
	.size	.Lstr.15, 22

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

	.type	.Lstr.17,@object                # @str.17
.Lstr.17:
	.asciz	"Calculating results for CPU vector multiplication "
	.size	.Lstr.17, 51

	.type	.Lstr.19,@object                # @str.19
.Lstr.19:
	.asciz	"CPU computation time (milliseconds) "
	.size	.Lstr.19, 37

	.type	.Lstr.20,@object                # @str.20
.Lstr.20:
	.asciz	"Calling kernel for gpu computations for vector multiplication and calculating results"
	.size	.Lstr.20, 86

	.type	.Lstr.22,@object                # @str.22
.Lstr.22:
	.asciz	"GPU computation time (milliseconds) "
	.size	.Lstr.22, 37

	.type	.Lstr.23,@object                # @str.23
.Lstr.23:
	.asciz	"Total Memory transfer time between CPU and GPU (milliseconds)"
	.size	.Lstr.23, 62

	.type	.Lstr.24,@object                # @str.24
.Lstr.24:
	.asciz	"Speedup: "
	.size	.Lstr.24, 10

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

	.type	.Lstr.26,@object                # @str.26
.Lstr.26:
	.asciz	"---------"
	.size	.Lstr.26, 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	512                             # 0x200
	.long	1024                            # 0x400
	.long	4096                            # 0x1000
	.size	.Lswitch.table.main, 24

	.type	.Lswitch.table.main.27,@object  # @switch.table.main.27
	.p2align	3, 0x0
.Lswitch.table.main.27:
	.quad	.Lstr.15
	.quad	.Lstr.14
	.quad	.Lstr.13
	.quad	.Lstr.12
	.quad	.Lstr.11
	.quad	.Lstr.10
	.size	.Lswitch.table.main.27, 48

	.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 _Z29__device_stub__mul_matrix_gpuPdS_S_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z14mul_matrix_gpuPdS_S_i
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z14mul_matrix_gpuPdS_S_i ; -- Begin function _Z14mul_matrix_gpuPdS_S_i
	.globl	_Z14mul_matrix_gpuPdS_S_i
	.p2align	8
	.type	_Z14mul_matrix_gpuPdS_S_i,@function
_Z14mul_matrix_gpuPdS_S_i:              ; @_Z14mul_matrix_gpuPdS_S_i
; %bb.0:
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x18
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_bfe_u32 v4, v0, 10, 10
	v_and_b32_e32 v7, 0x3ff, v0
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshl_add_u32 v1, s15, 3, v4
	s_waitcnt lgkmcnt(0)
	v_mul_lo_u32 v3, v1, s2
	v_mov_b32_e32 v1, 0
	v_mov_b32_e32 v2, 0
	v_lshl_add_u32 v0, s14, 3, v7
	s_cmp_lt_i32 s2, 8
	s_cbranch_scc1 .LBB0_5
; %bb.1:                                ; %.lr.ph
	v_lshlrev_b32_e32 v1, 3, v7
	v_lshlrev_b32_e32 v5, 6, v4
	s_ashr_i32 s3, s2, 31
	v_add_nc_u32_e32 v7, v3, v7
	s_lshr_b32 s3, s3, 29
	v_add_nc_u32_e32 v6, 0x200, v1
	v_dual_mov_b32 v1, 0 :: v_dual_add_nc_u32 v8, v5, v1
	v_mov_b32_e32 v2, 0
	s_add_i32 s3, s2, s3
	s_delay_alu instid0(VALU_DEP_3)
	v_add_nc_u32_e32 v9, v6, v5
	s_ashr_i32 s3, s3, 3
	s_mov_b32 s8, 0
.LBB0_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_3 Depth 2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b32 s9, s8, 3
	v_add_nc_u32_e32 v11, s9, v4
	v_add_nc_u32_e32 v10, s9, v7
	s_mov_b32 s9, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[12:13], null, v11, s2, v[0:1]
	v_ashrrev_i32_e32 v11, 31, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_lshlrev_b64 v[10:11], 3, v[10:11]
	v_ashrrev_i32_e32 v13, 31, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v10, vcc_lo, s4, v10
	v_lshlrev_b64 v[12:13], 3, v[12:13]
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_ci_u32_e32 v11, vcc_lo, s5, v11, vcc_lo
	v_add_co_u32 v12, vcc_lo, s6, v12
	s_delay_alu instid0(VALU_DEP_3)
	v_add_co_ci_u32_e32 v13, vcc_lo, s7, v13, vcc_lo
	global_load_b64 v[14:15], v[10:11], off
	global_load_b64 v[11:12], v[12:13], off
	v_mov_b32_e32 v10, v6
	s_waitcnt vmcnt(1)
	ds_store_b64 v8, v[14:15]
	s_waitcnt vmcnt(0)
	ds_store_b64 v9, v[11:12]
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB0_3:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_add_nc_u32_e32 v13, s9, v5
	s_add_i32 s9, s9, 8
	ds_load_b64 v[11:12], v10
	ds_load_b64 v[13:14], v13
	v_add_nc_u32_e32 v10, 64, v10
	s_cmp_eq_u32 s9, 64
	s_waitcnt lgkmcnt(0)
	v_fma_f64 v[1:2], v[13:14], v[11:12], v[1:2]
	s_cbranch_scc0 .LBB0_3
; %bb.4:                                ;   in Loop: Header=BB0_2 Depth=1
	s_add_i32 s8, s8, 1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_eq_u32 s8, s3
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc0 .LBB0_2
.LBB0_5:                                ; %Flow62
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v3, v0
	v_ashrrev_i32_e32 v4, 31, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[3:4], 3, v[3:4]
	v_add_co_u32 v3, vcc_lo, s0, v3
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v4, vcc_lo, s1, v4, vcc_lo
	global_store_b64 v[3:4], v[1:2], off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14mul_matrix_gpuPdS_S_i
		.amdhsa_group_segment_fixed_size 1024
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 28
		.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 16
		.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	_Z14mul_matrix_gpuPdS_S_i, .Lfunc_end0-_Z14mul_matrix_gpuPdS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 444
; NumSgprs: 18
; NumVgprs: 16
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 1024 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 16
; 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
    .group_segment_fixed_size: 1024
    .kernarg_segment_align: 8
    .kernarg_segment_size: 28
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z14mul_matrix_gpuPdS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z14mul_matrix_gpuPdS_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     16
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
