/**************************************************************************************************************************************************************************************************************************************************
** Rezwan Matin
** EE5321  HW#5_2
** Filename: HW5_2.cu
** Due: 4/5/19
**
** Objective:
** To multiply two 2000 x 2000 matrices whose elements are randomly generated, to run the multiplication 50 times and find the average of each element of the output, to print out the resultant matrix, and to time the program - all using CUDA C.
**
***************************************************************************************************************************************************************************************************************************************************/

#include <stdio.h>
#include <time.h>
#include <cuda.h>
#include <stdlib.h>
#include <math.h>

#define Tile_size 32
#define dimension 2000
#define iterations 50

int numARows = dimension;   //Number of rows in the matrix A.
int numAColumns = dimension;  //Number of columns in the matrix A.
int numBRows = dimension;   //Number of rows in the matrix B.
int numBColumns = dimension;  //Number of columns in the matrix B.
int numCRows;  //Number of rows in the matrix C.
int numCColumns;  //Number of columns in the matrix C.


/*************************************************************
*
* 	Matrix Multiplication Kernel Function
*
**************************************************************/
__global__ void matrixMultiply(float * A, float * B, float * C,
                               int numARows, int numAColumns,
                               int numBRows, int numBColumns,
                               int numCRows, int numCColumns)  {

        //Tile size arrays.
        __shared__ float sA[Tile_size][Tile_size];
    	__shared__ float sB[Tile_size][Tile_size];

        //IDs of threads.
    	int Row = blockDim.y*blockIdx.y + threadIdx.y;
    	int Col = blockDim.x*blockIdx.x + threadIdx.x;
    	float Cvalue = 0.0;

        //Initialized to zero.
        sA[threadIdx.y][threadIdx.x] = 0.0;
    	sB[threadIdx.y][threadIdx.x] = 0.0;

    	for (int k = 0; k < (((numAColumns - 1)/ Tile_size) + 1); k++)  {

            //Copy data to tile from matrix A.
            if ( (Row < numARows) && (threadIdx.x + (k*Tile_size)) < numAColumns) {
                sA[threadIdx.y][threadIdx.x] = A[(Row*numAColumns) + threadIdx.x + (k*Tile_size)];
            }
        	else {
                sA[threadIdx.y][threadIdx.x] = 0.0; //If outside the boundaries of the matrix.
        	}

            //Copy data to tile from matrix B.
            if ( Col < numBColumns && (threadIdx.y + k*Tile_size) < numBRows) {
                sB[threadIdx.y][threadIdx.x] = B[(threadIdx.y + k*Tile_size)*numBColumns + Col];
        	}
        	else {
                sB[threadIdx.y][threadIdx.x] = 0.0; //If outside the boundary of the matrix.
        	}

            //Making sure all threads have arrived before multiplications.
            __syncthreads();

            //Multiplying elements present in the current tile.
        	for (int j = 0; j < Tile_size; ++j) {
                Cvalue += sA[threadIdx.y][j] * sB[j][threadIdx.x];  //Each thread's C-value in the tile.
        	}
    	}

        //Saving final result to Matrix C.
        if (Row < numCRows && Col < numCColumns) {
            C[Row*numCColumns + Col] += Cvalue;
    	}
}

/*************************************************************
*
* 	Matrix Print function
*
**************************************************************/

void Print_Mat(int Row,int Col,float * Mat) {

	for(int i=0;i<Row*Col;i++) {

		printf("%f  ",*(Mat+i)/50);

   		if(( i % Col) == 0 ) {
            printf("\n");
        }
   	}
}

/*************************************************************
*
* 	Main Function
*
**************************************************************/
int main(int argc, char ** argv) {

        float * hostA; // The A matrix
    	float * hostB; // The B matrix
    	float * hostC; // The output C matrix
    	float * deviceA;
    	float * deviceB;
    	float * deviceC;

        clock_t start_whole, finish_whole;
        float time_for_program = 0;

        //Start the timer for the entire program.
        start_whole = clock();

        //Seed the rand function.
        srand(time(0));

    	hostA = (float *) malloc(sizeof(float)*numARows*numAColumns);
    	hostB = (float *) malloc(sizeof(float)*numBRows*numBColumns);

        //Setting numCRows and numCColumns.
    	numCRows = numARows;
    	numCColumns = numBColumns;

    	hostC = (float *) malloc(sizeof(float)*numCRows*numCColumns);

        //Allocating GPU memory.
    	cudaMalloc((void **)&deviceA, sizeof(float)*numARows*numAColumns);
    	cudaMalloc((void **)&deviceB, sizeof(float)*numBRows*numBColumns);
    	cudaMalloc((void **)&deviceC, sizeof(float)*numCRows*numCColumns);

        //Initializing matrix C.
        for (int i = 0; i < numARows*numAColumns; i++) {
                hostC[i]=0.0;
            }

        //Copy matrix C to GPU.
        cudaMemcpy(deviceC, hostC, sizeof(float)*numCRows*numCColumns, cudaMemcpyHostToDevice);

        //Initialize the grid and block dimensions.
        dim3 dimGrid((numCColumns/Tile_size) + 1, (numCRows/Tile_size) + 1, 1); //Number of Blocks required.
        dim3 dimBlock(Tile_size, Tile_size, 1);	 //Number of threads in each block.

        //Defining loop for running multiple iterations.
        for (int l = 0; l < iterations; l++){

            //Generating matrix A.
            for (int i = 0; i < numARows*numAColumns; i++) {
                hostA[i]=(199.9 + 199.9)*((float) rand()/(float)RAND_MAX) - 199.9;
            }

            //Generating matrix B.
            for (int i = 0; i < numBRows*numBColumns; i++) {
                hostB[i]=(199.9 + 199.9)*((float) rand()/(float)RAND_MAX) - 199.9;;
            }

            //Copy memory to the GPU.
            cudaMemcpy(deviceA, hostA, sizeof(float)*numARows*numAColumns, cudaMemcpyHostToDevice);
            cudaMemcpy(deviceB, hostB, sizeof(float)*numBRows*numBColumns, cudaMemcpyHostToDevice);

            //Making the kernel call.
            matrixMultiply<<<dimGrid, dimBlock>>>(deviceA, deviceB, deviceC,
						numARows, numAColumns,
						numBRows, numBColumns,
						numCRows, numCColumns);

            //Synchronize threads from the GPU.
            cudaDeviceSynchronize();

        }

    	// Copy the results in GPU memory back to the CPU.
    	cudaMemcpy(hostC, deviceC, sizeof(float)*numCRows*numCColumns, cudaMemcpyDeviceToHost);

    	printf("\n Average Matrix From Device\n");
    	Print_Mat(numCRows,numCColumns,hostC);


    	printf("\n Number of Blocks Created:%d \n",((numCColumns/Tile_size) + 1)*((numCColumns/Tile_size) + 1));
    	printf("\n Number of Threads Per Block: %d \n",(Tile_size*Tile_size));

        //Ending the timer for the entire program.
        finish_whole = clock();

        //Calculate time taken to run the entire program.
        time_for_program = ((float) (finish_whole - start_whole)) / CLOCKS_PER_SEC;

        //Print time for entire program.
        printf("\n Time for running the entire program = %f \n", time_for_program);

        //Free the GPU memory.
    	cudaFree(deviceA);
    	cudaFree(deviceB);
    	cudaFree(deviceC);

        //Free the Pointer Memory.
    	free(hostA);
    	free(hostB);
    	free(hostC);

        printf("\n Deallocations done!...Goodnight!\n");

    	return 0;
}

	.file	"tmpxft_0031d6e6_00000000-6_HW5_2.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"%f  "
.LC2:
	.string	"\n"
	.text
	.globl	_Z9Print_MatiiPf
	.type	_Z9Print_MatiiPf, @function
_Z9Print_MatiiPf:
.LFB2027:
	.cfi_startproc
	endbr64
	imull	%esi, %edi
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movl	%edi, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	xorl	%ebx, %ebx
.L3:
	movl	%ebx, %r13d
	cmpl	%ebx, %r12d
	jle	.L8
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	movss	(%r14,%rbx,4), %xmm0
	divss	.LC0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movl	%r13d, %eax
	cltd
	idivl	%ebp
	testl	%edx, %edx
	jne	.L4
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L4:
	incq	%rbx
	jmp	.L3
.L8:
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z9Print_MatiiPf, .-_Z9Print_MatiiPf
	.globl	_Z44__device_stub__Z14matrixMultiplyPfS_S_iiiiiiPfS_S_iiiiii
	.type	_Z44__device_stub__Z14matrixMultiplyPfS_S_iiiiiiPfS_S_iiiiii, @function
_Z44__device_stub__Z14matrixMultiplyPfS_S_iiiiiiPfS_S_iiiiii:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$200, %rsp
	.cfi_def_cfa_offset 208
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 16(%rsp)
	movl	%r9d, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 184(%rsp)
	xorl	%eax, %eax
	leaq	40(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	208(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	216(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	224(%rsp), %rax
	movq	%rax, 176(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	pushq	56(%rsp)
	.cfi_def_cfa_offset 216
	leaq	_Z14matrixMultiplyPfS_S_iiiiii(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 224
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 216
	popq	%rdx
	.cfi_def_cfa_offset 208
.L9:
	movq	184(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$200, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z44__device_stub__Z14matrixMultiplyPfS_S_iiiiiiPfS_S_iiiiii, .-_Z44__device_stub__Z14matrixMultiplyPfS_S_iiiiiiPfS_S_iiiiii
	.globl	_Z14matrixMultiplyPfS_S_iiiiii
	.type	_Z14matrixMultiplyPfS_S_iiiiii, @function
_Z14matrixMultiplyPfS_S_iiiiii:
.LFB2054:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z14matrixMultiplyPfS_S_iiiiiiPfS_S_iiiiii
	.cfi_endproc
.LFE2054:
	.size	_Z14matrixMultiplyPfS_S_iiiiii, .-_Z14matrixMultiplyPfS_S_iiiiii
	.section	.rodata.str1.1
.LC7:
	.string	"\n Average Matrix From Device\n"
.LC8:
	.string	"\n Number of Blocks Created:%d \n"
.LC9:
	.string	"\n Number of Threads Per Block: %d \n"
.LC11:
	.string	"\n Time for running the entire program = %f \n"
.LC12:
	.string	"\n Deallocations done!...Goodnight!\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.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
	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	clock@PLT
	xorl	%edi, %edi
	movq	%rax, %r13
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	movslq	numARows(%rip), %rbx
	movslq	numAColumns(%rip), %r14
	imulq	%rbx, %r14
	movl	%ebx, 12(%rsp)
	leaq	0(,%r14,4), %rsi
	movq	%rsi, %rdi
	movq	%rsi, (%rsp)
	call	malloc@PLT
	movslq	numBColumns(%rip), %r15
	movslq	numBRows(%rip), %rdi
	movq	%rax, %r12
	imulq	%r15, %rdi
	imulq	%r15, %rbx
	salq	$2, %rdi
	call	malloc@PLT
	movl	12(%rsp), %ecx
	leaq	0(,%rbx,4), %rdi
	movl	%r15d, numCColumns(%rip)
	movq	%rax, %rbp
	movl	%ecx, numCRows(%rip)
	call	malloc@PLT
	movq	(%rsp), %rsi
	leaq	24(%rsp), %rdi
	movq	%rax, %rbx
	call	cudaMalloc@PLT
	movslq	numBRows(%rip), %rsi
	leaq	32(%rsp), %rdi
	movslq	numBColumns(%rip), %rax
	imulq	%rax, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movslq	numCRows(%rip), %rsi
	leaq	40(%rsp), %rdi
	movslq	numCColumns(%rip), %rax
	imulq	%rax, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	numARows(%rip), %edx
	xorl	%eax, %eax
	imull	numAColumns(%rip), %edx
.L15:
	cmpl	%eax, %edx
	jle	.L26
	movl	$0x00000000, (%rbx,%rax,4)
	incq	%rax
	jmp	.L15
.L26:
	movslq	numCRows(%rip), %rdx
	movslq	numCColumns(%rip), %rax
	movl	$1, %ecx
	movq	%rbx, %rsi
	movq	40(%rsp), %rdi
	movl	$50, %r14d
	imulq	%rax, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	numCRows(%rip), %eax
	movl	$32, %esi
	movl	$1, 56(%rsp)
	movl	$1, 68(%rsp)
	cltd
	idivl	%esi
	leal	1(%rax), %ecx
	movl	numCColumns(%rip), %eax
	movl	%ecx, 52(%rsp)
	cltd
	idivl	%esi
	incl	%eax
	movl	%eax, 48(%rsp)
	movabsq	$137438953504, %rax
	movq	%rax, 60(%rsp)
.L17:
	xorl	%r15d, %r15d
.L23:
	movl	numARows(%rip), %eax
	imull	numAColumns(%rip), %eax
	cmpl	%r15d, %eax
	jle	.L27
	call	rand@PLT
	cvtsi2ssl	%eax, %xmm0
	mulss	.LC4(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	mulsd	.LC5(%rip), %xmm0
	subsd	.LC6(%rip), %xmm0
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, (%r12,%r15,4)
	incq	%r15
	jmp	.L23
.L27:
	xorl	%r15d, %r15d
.L19:
	movl	numBRows(%rip), %eax
	imull	numBColumns(%rip), %eax
	cmpl	%r15d, %eax
	jle	.L28
	call	rand@PLT
	cvtsi2ssl	%eax, %xmm0
	mulss	.LC4(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	mulsd	.LC5(%rip), %xmm0
	subsd	.LC6(%rip), %xmm0
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, 0(%rbp,%r15,4)
	incq	%r15
	jmp	.L19
.L28:
	movslq	numARows(%rip), %rdx
	movslq	numAColumns(%rip), %rax
	movl	$1, %ecx
	movq	%r12, %rsi
	movq	24(%rsp), %rdi
	imulq	%rax, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movslq	numBRows(%rip), %rdx
	movq	32(%rsp), %rdi
	movq	%rbp, %rsi
	movslq	numBColumns(%rip), %rax
	movl	$1, %ecx
	imulq	%rax, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	68(%rsp), %ecx
	movl	56(%rsp), %esi
	xorl	%r9d, %r9d
	movq	60(%rsp), %rdx
	movq	48(%rsp), %rdi
	xorl	%r8d, %r8d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	movl	numBRows(%rip), %r9d
	movl	numAColumns(%rip), %r8d
	pushq	%rax
	.cfi_def_cfa_offset 152
	movl	numCColumns(%rip), %eax
	movl	numARows(%rip), %ecx
	pushq	%rax
	.cfi_def_cfa_offset 160
	movl	numCRows(%rip), %eax
	pushq	%rax
	.cfi_def_cfa_offset 168
	movl	numBColumns(%rip), %eax
	pushq	%rax
	.cfi_def_cfa_offset 176
	movq	72(%rsp), %rdx
	movq	64(%rsp), %rsi
	movq	56(%rsp), %rdi
	call	_Z44__device_stub__Z14matrixMultiplyPfS_S_iiiiiiPfS_S_iiiiii
	addq	$32, %rsp
	.cfi_def_cfa_offset 144
.L21:
	call	cudaDeviceSynchronize@PLT
	decl	%r14d
	jne	.L17
	movslq	numCRows(%rip), %rdx
	movslq	numCColumns(%rip), %rax
	movl	$2, %ecx
	movq	%rbx, %rdi
	movq	40(%rsp), %rsi
	imulq	%rax, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	numCColumns(%rip), %esi
	movl	numCRows(%rip), %edi
	movq	%rbx, %rdx
	call	_Z9Print_MatiiPf
	movl	numCColumns(%rip), %eax
	movl	$32, %ecx
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	cltd
	idivl	%ecx
	leal	1(%rax), %edx
	xorl	%eax, %eax
	imull	%edx, %edx
	call	__printf_chk@PLT
	movl	$1024, %edx
	leaq	.LC9(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	call	clock@PLT
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	subq	%r13, %rax
	cvtsi2ssq	%rax, %xmm0
	movb	$1, %al
	divss	.LC10(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	xorl	%eax, %eax
	movl	$2, %edi
	leaq	.LC12(%rip), %rsi
	call	__printf_chk@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L24
	call	__stack_chk_fail@PLT
.L24:
	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
.LFE2028:
	.size	main, .-main
	.section	.rodata.str1.1
.LC13:
	.string	"_Z14matrixMultiplyPfS_S_iiiiii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2056:
	.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	.LC13(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z14matrixMultiplyPfS_S_iiiiii(%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
.LFE2056:
	.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	numCColumns
	.bss
	.align 4
	.type	numCColumns, @object
	.size	numCColumns, 4
numCColumns:
	.zero	4
	.globl	numCRows
	.align 4
	.type	numCRows, @object
	.size	numCRows, 4
numCRows:
	.zero	4
	.globl	numBColumns
	.data
	.align 4
	.type	numBColumns, @object
	.size	numBColumns, 4
numBColumns:
	.long	2000
	.globl	numBRows
	.align 4
	.type	numBRows, @object
	.size	numBRows, 4
numBRows:
	.long	2000
	.globl	numAColumns
	.align 4
	.type	numAColumns, @object
	.size	numAColumns, 4
numAColumns:
	.long	2000
	.globl	numARows
	.align 4
	.type	numARows, @object
	.size	numARows, 4
numARows:
	.long	2000
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC0:
	.long	1112014848
	.align 4
.LC4:
	.long	805306368
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC5:
	.long	-858993459
	.long	1081670860
	.align 8
.LC6:
	.long	-858993459
	.long	1080622284
	.section	.rodata.cst4
	.align 4
.LC10:
	.long	1232348160
	.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:

/**************************************************************************************************************************************************************************************************************************************************
** Rezwan Matin
** EE5321  HW#5_2
** Filename: HW5_2.cu
** Due: 4/5/19
**
** Objective:
** To multiply two 2000 x 2000 matrices whose elements are randomly generated, to run the multiplication 50 times and find the average of each element of the output, to print out the resultant matrix, and to time the program - all using CUDA C.
**
***************************************************************************************************************************************************************************************************************************************************/

#include <stdio.h>
#include <time.h>
#include <hip/hip_runtime.h>
#include <stdlib.h>
#include <math.h>

#define Tile_size 32
#define dimension 2000
#define iterations 50

int numARows = dimension;   //Number of rows in the matrix A.
int numAColumns = dimension;  //Number of columns in the matrix A.
int numBRows = dimension;   //Number of rows in the matrix B.
int numBColumns = dimension;  //Number of columns in the matrix B.
int numCRows;  //Number of rows in the matrix C.
int numCColumns;  //Number of columns in the matrix C.


/*************************************************************
*
* 	Matrix Multiplication Kernel Function
*
**************************************************************/
__global__ void matrixMultiply(float * A, float * B, float * C,
                               int numARows, int numAColumns,
                               int numBRows, int numBColumns,
                               int numCRows, int numCColumns)  {

        //Tile size arrays.
        __shared__ float sA[Tile_size][Tile_size];
    	__shared__ float sB[Tile_size][Tile_size];

        //IDs of threads.
    	int Row = blockDim.y*blockIdx.y + threadIdx.y;
    	int Col = blockDim.x*blockIdx.x + threadIdx.x;
    	float Cvalue = 0.0;

        //Initialized to zero.
        sA[threadIdx.y][threadIdx.x] = 0.0;
    	sB[threadIdx.y][threadIdx.x] = 0.0;

    	for (int k = 0; k < (((numAColumns - 1)/ Tile_size) + 1); k++)  {

            //Copy data to tile from matrix A.
            if ( (Row < numARows) && (threadIdx.x + (k*Tile_size)) < numAColumns) {
                sA[threadIdx.y][threadIdx.x] = A[(Row*numAColumns) + threadIdx.x + (k*Tile_size)];
            }
        	else {
                sA[threadIdx.y][threadIdx.x] = 0.0; //If outside the boundaries of the matrix.
        	}

            //Copy data to tile from matrix B.
            if ( Col < numBColumns && (threadIdx.y + k*Tile_size) < numBRows) {
                sB[threadIdx.y][threadIdx.x] = B[(threadIdx.y + k*Tile_size)*numBColumns + Col];
        	}
        	else {
                sB[threadIdx.y][threadIdx.x] = 0.0; //If outside the boundary of the matrix.
        	}

            //Making sure all threads have arrived before multiplications.
            __syncthreads();

            //Multiplying elements present in the current tile.
        	for (int j = 0; j < Tile_size; ++j) {
                Cvalue += sA[threadIdx.y][j] * sB[j][threadIdx.x];  //Each thread's C-value in the tile.
        	}
    	}

        //Saving final result to Matrix C.
        if (Row < numCRows && Col < numCColumns) {
            C[Row*numCColumns + Col] += Cvalue;
    	}
}

/*************************************************************
*
* 	Matrix Print function
*
**************************************************************/

void Print_Mat(int Row,int Col,float * Mat) {

	for(int i=0;i<Row*Col;i++) {

		printf("%f  ",*(Mat+i)/50);

   		if(( i % Col) == 0 ) {
            printf("\n");
        }
   	}
}

/*************************************************************
*
* 	Main Function
*
**************************************************************/
int main(int argc, char ** argv) {

        float * hostA; // The A matrix
    	float * hostB; // The B matrix
    	float * hostC; // The output C matrix
    	float * deviceA;
    	float * deviceB;
    	float * deviceC;

        clock_t start_whole, finish_whole;
        float time_for_program = 0;

        //Start the timer for the entire program.
        start_whole = clock();

        //Seed the rand function.
        srand(time(0));

    	hostA = (float *) malloc(sizeof(float)*numARows*numAColumns);
    	hostB = (float *) malloc(sizeof(float)*numBRows*numBColumns);

        //Setting numCRows and numCColumns.
    	numCRows = numARows;
    	numCColumns = numBColumns;

    	hostC = (float *) malloc(sizeof(float)*numCRows*numCColumns);

        //Allocating GPU memory.
    	hipMalloc((void **)&deviceA, sizeof(float)*numARows*numAColumns);
    	hipMalloc((void **)&deviceB, sizeof(float)*numBRows*numBColumns);
    	hipMalloc((void **)&deviceC, sizeof(float)*numCRows*numCColumns);

        //Initializing matrix C.
        for (int i = 0; i < numARows*numAColumns; i++) {
                hostC[i]=0.0;
            }

        //Copy matrix C to GPU.
        hipMemcpy(deviceC, hostC, sizeof(float)*numCRows*numCColumns, hipMemcpyHostToDevice);

        //Initialize the grid and block dimensions.
        dim3 dimGrid((numCColumns/Tile_size) + 1, (numCRows/Tile_size) + 1, 1); //Number of Blocks required.
        dim3 dimBlock(Tile_size, Tile_size, 1);	 //Number of threads in each block.

        //Defining loop for running multiple iterations.
        for (int l = 0; l < iterations; l++){

            //Generating matrix A.
            for (int i = 0; i < numARows*numAColumns; i++) {
                hostA[i]=(199.9 + 199.9)*((float) rand()/(float)RAND_MAX) - 199.9;
            }

            //Generating matrix B.
            for (int i = 0; i < numBRows*numBColumns; i++) {
                hostB[i]=(199.9 + 199.9)*((float) rand()/(float)RAND_MAX) - 199.9;;
            }

            //Copy memory to the GPU.
            hipMemcpy(deviceA, hostA, sizeof(float)*numARows*numAColumns, hipMemcpyHostToDevice);
            hipMemcpy(deviceB, hostB, sizeof(float)*numBRows*numBColumns, hipMemcpyHostToDevice);

            //Making the kernel call.
            matrixMultiply<<<dimGrid, dimBlock>>>(deviceA, deviceB, deviceC,
						numARows, numAColumns,
						numBRows, numBColumns,
						numCRows, numCColumns);

            //Synchronize threads from the GPU.
            hipDeviceSynchronize();

        }

    	// Copy the results in GPU memory back to the CPU.
    	hipMemcpy(hostC, deviceC, sizeof(float)*numCRows*numCColumns, hipMemcpyDeviceToHost);

    	printf("\n Average Matrix From Device\n");
    	Print_Mat(numCRows,numCColumns,hostC);


    	printf("\n Number of Blocks Created:%d \n",((numCColumns/Tile_size) + 1)*((numCColumns/Tile_size) + 1));
    	printf("\n Number of Threads Per Block: %d \n",(Tile_size*Tile_size));

        //Ending the timer for the entire program.
        finish_whole = clock();

        //Calculate time taken to run the entire program.
        time_for_program = ((float) (finish_whole - start_whole)) / CLOCKS_PER_SEC;

        //Print time for entire program.
        printf("\n Time for running the entire program = %f \n", time_for_program);

        //Free the GPU memory.
    	hipFree(deviceA);
    	hipFree(deviceB);
    	hipFree(deviceC);

        //Free the Pointer Memory.
    	free(hostA);
    	free(hostB);
    	free(hostC);

        printf("\n Deallocations done!...Goodnight!\n");

    	return 0;
}

	.text
	.file	"HW5_2.hip"
	.globl	_Z29__device_stub__matrixMultiplyPfS_S_iiiiii # -- Begin function _Z29__device_stub__matrixMultiplyPfS_S_iiiiii
	.type	_Z29__device_stub__matrixMultiplyPfS_S_iiiiii,@function
_Z29__device_stub__matrixMultiplyPfS_S_iiiiii: # @_Z29__device_stub__matrixMultiplyPfS_S_iiiiii
	.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	$176, %rsp
	.cfi_def_cfa_offset 224
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	20(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	16(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	12(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	224(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	232(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	240(%rsp), %rax
	movq	%rax, 64(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%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	$_Z14matrixMultiplyPfS_S_iiiiii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$192, %rsp
	.cfi_adjust_cfa_offset -192
	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__matrixMultiplyPfS_S_iiiiii, .Lfunc_end0-_Z29__device_stub__matrixMultiplyPfS_S_iiiiii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z9Print_MatiiPf
.LCPI1_0:
	.long	0x42480000                      # float 50
	.text
	.globl	_Z9Print_MatiiPf
	.type	_Z9Print_MatiiPf,@function
_Z9Print_MatiiPf:                       # @_Z9Print_MatiiPf
	.cfi_startproc
# %bb.0:
	imull	%esi, %edi
	testl	%edi, %edi
	jle	.LBB1_6
# %bb.1:                                # %.lr.ph.preheader
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdx, %rbx
	movl	%esi, %ebp
	movl	%edi, %r15d
	xorl	%r14d, %r14d
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movss	(%rbx,%r14,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	divss	.LCPI1_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str, %edi
	movb	$1, %al
	callq	printf
	movl	%r14d, %eax
	cltd
	idivl	%ebp
	testl	%edx, %edx
	jne	.LBB1_4
# %bb.3:                                #   in Loop: Header=BB1_2 Depth=1
	movl	$10, %edi
	callq	putchar@PLT
.LBB1_4:                                #   in Loop: Header=BB1_2 Depth=1
	incq	%r14
	cmpq	%r14, %r15
	jne	.LBB1_2
# %bb.5:
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.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
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %r15
	.cfi_restore %rbp
.LBB1_6:                                # %._crit_edge
	retq
.Lfunc_end1:
	.size	_Z9Print_MatiiPf, .Lfunc_end1-_Z9Print_MatiiPf
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI2_0:
	.long	0x30000000                      # float 4.65661287E-10
.LCPI2_3:
	.long	0x49742400                      # float 1.0E+6
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI2_1:
	.quad	0x4078fccccccccccd              # double 399.80000000000001
.LCPI2_2:
	.quad	0xc068fccccccccccd              # double -199.90000000000001
	.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	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	callq	clock
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	xorl	%r14d, %r14d
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	movslq	numARows(%rip), %r12
	movq	%r12, 8(%rsp)                   # 8-byte Spill
	shlq	$2, %r12
	movslq	numAColumns(%rip), %r13
	imulq	%r12, %r13
	movq	%r13, %rdi
	callq	malloc
	movq	%rax, %rbx
	movslq	numBRows(%rip), %rdi
	movslq	numBColumns(%rip), %rbp
	imulq	%rbp, %rdi
	shlq	$2, %rdi
	callq	malloc
	movq	%rax, %r15
	movq	8(%rsp), %rax                   # 8-byte Reload
	movl	%eax, numCRows(%rip)
	movl	%ebp, numCColumns(%rip)
	imulq	%rbp, %r12
	movq	%r12, %rdi
	callq	malloc
	movq	%rax, %r12
	leaq	24(%rsp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	movslq	numBRows(%rip), %rax
	movslq	numBColumns(%rip), %rsi
	imulq	%rax, %rsi
	shlq	$2, %rsi
	leaq	16(%rsp), %rdi
	callq	hipMalloc
	movslq	numCRows(%rip), %rax
	movslq	numCColumns(%rip), %rsi
	imulq	%rax, %rsi
	shlq	$2, %rsi
	movq	%rsp, %rdi
	callq	hipMalloc
	movl	numAColumns(%rip), %eax
	imull	numARows(%rip), %eax
	testl	%eax, %eax
	jle	.LBB2_2
# %bb.1:                                # %.lr.ph.preheader
	movl	%eax, %edx
	shlq	$2, %rdx
	movq	%r12, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
.LBB2_2:                                # %._crit_edge
	movq	(%rsp), %rdi
	movslq	numCRows(%rip), %rax
	movslq	numCColumns(%rip), %rdx
	imulq	%rax, %rdx
	shlq	$2, %rdx
	movq	%r12, 8(%rsp)                   # 8-byte Spill
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	numCColumns(%rip), %eax
	leal	31(%rax), %ecx
	testl	%eax, %eax
	cmovnsl	%eax, %ecx
	sarl	$5, %ecx
	incl	%ecx
	movl	numCRows(%rip), %eax
	leal	31(%rax), %r13d
	testl	%eax, %eax
	cmovnsl	%eax, %r13d
	sarl	$5, %r13d
	incl	%r13d
	shlq	$32, %r13
	orq	%rcx, %r13
	movabsq	$137438953504, %rbp             # imm = 0x2000000020
.LBB2_3:                                # %.preheader37
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_5 Depth 2
                                        #     Child Loop BB2_8 Depth 2
	movl	numARows(%rip), %eax
	movl	numAColumns(%rip), %ecx
	movl	%ecx, %edx
	imull	%eax, %edx
	testl	%edx, %edx
	jle	.LBB2_6
# %bb.4:                                # %.lr.ph40.preheader
                                        #   in Loop: Header=BB2_3 Depth=1
	xorl	%r12d, %r12d
.LBB2_5:                                # %.lr.ph40
                                        #   Parent Loop BB2_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	mulsd	.LCPI2_1(%rip), %xmm0
	addsd	.LCPI2_2(%rip), %xmm0
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, (%rbx,%r12,4)
	incq	%r12
	movl	numARows(%rip), %eax
	movl	numAColumns(%rip), %ecx
	movl	%ecx, %edx
	imull	%eax, %edx
	movslq	%edx, %rdx
	cmpq	%rdx, %r12
	jl	.LBB2_5
.LBB2_6:                                # %.preheader
                                        #   in Loop: Header=BB2_3 Depth=1
	movl	numBColumns(%rip), %edx
	imull	numBRows(%rip), %edx
	testl	%edx, %edx
	jle	.LBB2_10
# %bb.7:                                # %.lr.ph42.preheader
                                        #   in Loop: Header=BB2_3 Depth=1
	xorl	%r12d, %r12d
.LBB2_8:                                # %.lr.ph42
                                        #   Parent Loop BB2_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	mulsd	.LCPI2_1(%rip), %xmm0
	addsd	.LCPI2_2(%rip), %xmm0
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, (%r15,%r12,4)
	incq	%r12
	movslq	numBRows(%rip), %rax
	movslq	numBColumns(%rip), %rcx
	imulq	%rax, %rcx
	cmpq	%rcx, %r12
	jl	.LBB2_8
# %bb.9:                                # %._crit_edge43.loopexit
                                        #   in Loop: Header=BB2_3 Depth=1
	movl	numARows(%rip), %eax
	movl	numAColumns(%rip), %ecx
.LBB2_10:                               # %._crit_edge43
                                        #   in Loop: Header=BB2_3 Depth=1
	movq	24(%rsp), %rdi
	cltq
	movslq	%ecx, %rdx
	imulq	%rax, %rdx
	shlq	$2, %rdx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rdi
	movslq	numBRows(%rip), %rax
	movslq	numBColumns(%rip), %rdx
	imulq	%rax, %rdx
	shlq	$2, %rdx
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%rbp, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_12
# %bb.11:                               #   in Loop: Header=BB2_3 Depth=1
	movq	24(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	(%rsp), %rdx
	movl	numARows(%rip), %ecx
	movl	numAColumns(%rip), %r8d
	movl	numBRows(%rip), %r9d
	movl	numBColumns(%rip), %eax
	movl	numCRows(%rip), %r10d
	movl	numCColumns(%rip), %r11d
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%rax
	.cfi_adjust_cfa_offset 8
	callq	_Z29__device_stub__matrixMultiplyPfS_S_iiiiii
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
.LBB2_12:                               #   in Loop: Header=BB2_3 Depth=1
	callq	hipDeviceSynchronize
	incl	%r14d
	cmpl	$50, %r14d
	jne	.LBB2_3
# %bb.13:
	movq	(%rsp), %rsi
	movslq	numCRows(%rip), %rax
	movslq	numCColumns(%rip), %rdx
	imulq	%rax, %rdx
	shlq	$2, %rdx
	movq	8(%rsp), %r14                   # 8-byte Reload
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	numCRows(%rip), %edi
	movl	numCColumns(%rip), %esi
	movq	%r14, %rdx
	callq	_Z9Print_MatiiPf
	movl	numCColumns(%rip), %eax
	leal	31(%rax), %esi
	testl	%eax, %eax
	cmovnsl	%eax, %esi
	sarl	$5, %esi
	incl	%esi
	imull	%esi, %esi
	movl	$.L.str.3, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.4, %edi
	movl	$1024, %esi                     # imm = 0x400
	xorl	%eax, %eax
	callq	printf
	callq	clock
	subq	32(%rsp), %rax                  # 8-byte Folded Reload
	cvtsi2ss	%rax, %xmm0
	divss	.LCPI2_3(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.5, %edi
	movb	$1, %al
	callq	printf
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	movq	%rbx, %rdi
	callq	free
	movq	%r15, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	xorl	%eax, %eax
	addq	$40, %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	main, .Lfunc_end2-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	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z14matrixMultiplyPfS_S_iiiiii, %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_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__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	.LBB4_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
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	numARows,@object                # @numARows
	.data
	.globl	numARows
	.p2align	2, 0x0
numARows:
	.long	2000                            # 0x7d0
	.size	numARows, 4

	.type	numAColumns,@object             # @numAColumns
	.globl	numAColumns
	.p2align	2, 0x0
numAColumns:
	.long	2000                            # 0x7d0
	.size	numAColumns, 4

	.type	numBRows,@object                # @numBRows
	.globl	numBRows
	.p2align	2, 0x0
numBRows:
	.long	2000                            # 0x7d0
	.size	numBRows, 4

	.type	numBColumns,@object             # @numBColumns
	.globl	numBColumns
	.p2align	2, 0x0
numBColumns:
	.long	2000                            # 0x7d0
	.size	numBColumns, 4

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

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

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

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

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"\n Number of Blocks Created:%d \n"
	.size	.L.str.3, 32

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"\n Number of Threads Per Block: %d \n"
	.size	.L.str.4, 36

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"\n Time for running the entire program = %f \n"
	.size	.L.str.5, 45

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z14matrixMultiplyPfS_S_iiiiii"
	.size	.L__unnamed_1, 31

	.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	"\n Average Matrix From Device"
	.size	.Lstr, 29

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"\n Deallocations done!...Goodnight!"
	.size	.Lstr.1, 35

	.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__matrixMultiplyPfS_S_iiiiii
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z14matrixMultiplyPfS_S_iiiiii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z14matrixMultiplyPfS_S_iiiiii ; -- Begin function _Z14matrixMultiplyPfS_S_iiiiii
	.globl	_Z14matrixMultiplyPfS_S_iiiiii
	.p2align	8
	.type	_Z14matrixMultiplyPfS_S_iiiiii,@function
_Z14matrixMultiplyPfS_S_iiiiii:         ; @_Z14matrixMultiplyPfS_S_iiiiii
; %bb.0:
	s_clause 0x4
	s_load_b32 s16, s[0:1], 0x3c
	s_load_b128 s[4:7], s[0:1], 0x18
	s_load_b128 s[8:11], s[0:1], 0x0
	s_load_b64 s[2:3], s[0:1], 0x10
	s_load_b64 s[12:13], s[0:1], 0x28
	v_and_b32_e32 v3, 0x3ff, v0
	v_bfe_u32 v4, v0, 10, 10
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b32_e32 v5, 2, v3
	v_lshl_add_u32 v7, v4, 7, v5
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s0, s16, 16
	s_and_b32 s1, s16, 0xffff
	v_mad_u64_u32 v[0:1], null, s15, s0, v[4:5]
	v_mad_u64_u32 v[1:2], null, s14, s1, v[3:4]
	v_mov_b32_e32 v2, 0
	s_cmpk_lt_i32 s5, 0xffe2
	ds_store_2addr_stride64_b32 v7, v2, v2 offset1:16
	s_cbranch_scc1 .LBB0_14
; %bb.1:                                ; %.lr.ph
	s_add_i32 s0, s5, -1
	v_mul_lo_u32 v9, v0, s5
	s_ashr_i32 s1, s0, 31
	v_add_nc_u32_e32 v8, 0x1000, v7
	s_lshr_b32 s1, s1, 27
	v_cmp_gt_i32_e32 vcc_lo, s4, v0
	s_add_i32 s1, s0, s1
	v_cmp_gt_i32_e64 s0, s7, v1
	v_lshlrev_b32_e32 v10, 7, v4
	v_lshl_or_b32 v11, v3, 2, 0x1000
	v_mov_b32_e32 v6, 0
	v_mov_b32_e32 v2, 0
	s_ashr_i32 s4, s1, 5
	s_mov_b32 s14, 0
.LBB0_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_11 Depth 2
	v_mov_b32_e32 v5, 0
	s_and_saveexec_b32 s15, vcc_lo
	s_cbranch_execz .LBB0_6
; %bb.3:                                ;   in Loop: Header=BB0_2 Depth=1
	v_lshl_add_u32 v12, s14, 5, v3
	v_mov_b32_e32 v5, 0
	s_mov_b32 s16, exec_lo
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_gt_u32_e64 s5, v12
	s_cbranch_execz .LBB0_5
; %bb.4:                                ;   in Loop: Header=BB0_2 Depth=1
	v_add_nc_u32_e32 v5, v12, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[12:13], 2, v[5:6]
	v_add_co_u32 v12, s1, s8, v12
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v13, s1, s9, v13, s1
	global_load_b32 v5, v[12:13], off
.LBB0_5:                                ; %Flow84
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s16
.LBB0_6:                                ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s15
	v_mov_b32_e32 v12, 0
	s_waitcnt vmcnt(0)
	ds_store_b32 v7, v5
	s_and_saveexec_b32 s15, s0
	s_cbranch_execz .LBB0_10
; %bb.7:                                ;   in Loop: Header=BB0_2 Depth=1
	v_lshl_add_u32 v5, s14, 5, v4
	v_mov_b32_e32 v12, 0
	s_mov_b32 s16, exec_lo
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_gt_u32_e64 s6, v5
	s_cbranch_execz .LBB0_9
; %bb.8:                                ;   in Loop: Header=BB0_2 Depth=1
	v_mad_u64_u32 v[12:13], null, v5, s7, v[1:2]
	v_mov_b32_e32 v13, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[12:13], 2, v[12:13]
	v_add_co_u32 v12, s1, s10, v12
	s_delay_alu instid0(VALU_DEP_1)
	v_add_co_ci_u32_e64 v13, s1, s11, v13, s1
	global_load_b32 v12, v[12:13], off
.LBB0_9:                                ; %Flow
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_or_b32 exec_lo, exec_lo, s16
.LBB0_10:                               ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s15
	v_mov_b32_e32 v5, v11
	s_mov_b32 s1, 0
	s_waitcnt vmcnt(0)
	ds_store_b32 v8, v12
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
.LBB0_11:                               ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_add_nc_u32_e32 v12, s1, v10
	s_add_i32 s1, s1, 4
	ds_load_b32 v13, v5
	ds_load_b32 v12, v12
	v_add_nc_u32_e32 v5, 0x80, v5
	s_cmpk_eq_i32 s1, 0x80
	s_waitcnt lgkmcnt(0)
	v_fmac_f32_e32 v2, v12, v13
	s_cbranch_scc0 .LBB0_11
; %bb.12:                               ;   in Loop: Header=BB0_2 Depth=1
	s_add_i32 s1, s14, 1
	s_cmp_eq_u32 s14, s4
	s_cbranch_scc1 .LBB0_14
; %bb.13:                               ;   in Loop: Header=BB0_2 Depth=1
	s_mov_b32 s14, s1
	s_branch .LBB0_2
.LBB0_14:                               ; %Flow86
	v_cmp_gt_i32_e32 vcc_lo, s12, v0
	v_cmp_gt_i32_e64 s0, s13, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s0, vcc_lo, s0
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB0_16
; %bb.15:
	v_mad_u64_u32 v[3:4], null, v0, s13, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v4, 31, v3
	v_lshlrev_b64 v[0:1], 2, v[3:4]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s2, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_load_b32 v3, v[0:1], off
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v2, v2, v3
	global_store_b32 v[0:1], v2, off
.LBB0_16:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z14matrixMultiplyPfS_S_iiiiii
		.amdhsa_group_segment_fixed_size 8192
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.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 14
		.amdhsa_next_free_sgpr 17
		.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	_Z14matrixMultiplyPfS_S_iiiiii, .Lfunc_end0-_Z14matrixMultiplyPfS_S_iiiiii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 632
; NumSgprs: 19
; NumVgprs: 14
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 8192 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 19
; NumVGPRsForWavesPerEU: 14
; 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:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     by_value
      - .offset:         36
        .size:           4
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .offset:         44
        .size:           4
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 8192
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z14matrixMultiplyPfS_S_iiiiii
    .private_segment_fixed_size: 0
    .sgpr_count:     19
    .sgpr_spill_count: 0
    .symbol:         _Z14matrixMultiplyPfS_S_iiiiii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     14
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
