#include<stdio.h>
#include<sys/time.h>
#include<stdlib.h>
#include<iostream>

using namespace std;

//----------------------------------- Structures and Globals---------------------------------------------

typedef struct {
	int dimension1;
	int dimension2;	
} ArrayMetadata2D;

// metadata variables describing dimensionalities of all data structures involved in the computation
ArrayMetadata2D A_MD, B_MD, C_MD;
// pointers for input and output arrays in the host memory  
float *A, *B, *C, *C_CPU;
// pointers for input and output arrays in the device memory (NVIDIA DRAM)
float *A_GPU, *B_GPU, *C_GPU;

//----------------------------------- host function definitions -----------------------------------------

void allocateAndInitializeAB();
void computeCpuMMM();
void computeGpuMMM();
void copyMatricesToGPU();
void copyResultFromGPU();
void compareHostAndGpuOutput();
void die(const char *error); 
void check_error(cudaError e);

//----------------------------------- CUDA function definitions -----------------------------------------


//-------------------------------------------------------------------------------------------------------
int main(int argc, char **argv) {
	
	A_MD.dimension1 = (argc > 1) ? atoi(argv[1]) : 100;
	A_MD.dimension2 = (argc > 2) ? atoi(argv[2]) : A_MD.dimension1;
	B_MD.dimension1 = (argc > 3) ? atoi(argv[3]) : A_MD.dimension2;
	B_MD.dimension2 = (argc > 4) ? atoi(argv[4]) : B_MD.dimension1;
	C_MD.dimension1 = A_MD.dimension1;
	C_MD.dimension2 = B_MD.dimension2;

	printf("Matrix A is %d-by-%d\n", A_MD.dimension1, A_MD.dimension2);
	printf("Matrix B is %d-by-%d\n", B_MD.dimension1, B_MD.dimension2);
	printf("Matrix C is %d-by-%d\n", C_MD.dimension1, C_MD.dimension2);

	allocateAndInitializeAB();

	// matrix matrix multiplication in the CPU
	clock_t start = clock();	
	computeCpuMMM();
	clock_t end = clock();
    
    double elapsed = (end - start) / (double) CLOCKS_PER_SEC;
    printf("Computation time in the CPU: %f seconds\n", elapsed);

    start = clock();
    computeGpuMMM();
    end = clock();

    elapsed = (end - start) / (double) CLOCKS_PER_SEC;
    printf("Computation time in the GPU: %f seconds\n", elapsed);

    compareHostAndGpuOutput();

	return 0;
}

// allocate and initialize A and B using a random number generator
void allocateAndInitializeAB() {
	
    float val;
	
    size_t sizeofA = A_MD.dimension1 * A_MD.dimension2 * sizeof(float);
	A = (float*) malloc(sizeofA);
    A_GPU = (float*) malloc(sizeofA);

	srand(time(NULL));
  	for (int i = 0; i < A_MD.dimension1; i++) {
		for (int j = 0; j < A_MD.dimension2; j++) {
			int index = i * A_MD.dimension2 + j;
            val = (rand() % 1000) * 0.001;
			A[index] = val; 
            A_GPU[index] = val;
		}
	}
	
	size_t sizeofB = B_MD.dimension1 * B_MD.dimension2 * sizeof(float);
	B = (float*) malloc(sizeofB);
	B_GPU = (float*) malloc(sizeofB);
  	for (int i = 0; i < B_MD.dimension1; i++) {
		for (int j = 0; j < B_MD.dimension2; j++) {
			int index = i * B_MD.dimension2 + j;
            val = (rand() % 1000) * 0.001;
			B[index] = val; 
            B_GPU[index] = val;
		}
	}
}

// allocate memory in the GPU for all matrices, and copy A and B content from the host CPU memory to the GPU memory
void copyMatricesToGPU() {
	
	size_t sizeofA = A_MD.dimension1 * A_MD.dimension2 * sizeof(float);
	check_error(cudaMalloc((void **) &A_GPU, sizeofA));
	check_error(cudaMemcpy(A_GPU, A, sizeofA, cudaMemcpyHostToDevice));
	
	size_t sizeofB = B_MD.dimension1 * B_MD.dimension2 * sizeof(float);
	check_error(cudaMalloc((void **) &B_GPU, sizeofB));
	check_error(cudaMemcpy(B_GPU, B, sizeofB, cudaMemcpyHostToDevice));
	
	size_t sizeofC = C_MD.dimension1 * C_MD.dimension2 * sizeof(float);
	check_error(cudaMalloc((void **) &C_GPU, sizeofC));
}

// copy results from C_GPU which is in GPU card memory to C_CPU which is in the host CPU for result comparison
void copyResultFromGPU() {
	size_t sizeofC = C_MD.dimension1 * C_MD.dimension2 * sizeof(float);
	C_CPU = (float*) malloc(sizeofC);
	check_error(cudaMemcpy(C_CPU, C_GPU, sizeofC, cudaMemcpyDeviceToHost));
}

// do a straightforward matrix-matrix multiplication in the CPU
// notice that this implementation can be massively improved in the CPU by doing proper cache blocking but we are
// not providing you the efficient CPU implementation as that reveals too much about the ideal GPU implementation
void computeCpuMMM() {
	
	// allocate the result matrix for the CPU computation
	size_t sizeofC = C_MD.dimension1 * C_MD.dimension2 * sizeof(float);
	C = (float*) malloc(sizeofC);
	
	// compute C[i][j] as the sum of A[i][k] * B[k][j] for all columns k of A
	for (int i = 0; i < A_MD.dimension1; i++) {
		int a_i = i * A_MD.dimension2;
		int c_i = i * C_MD.dimension2;
		for (int j = 0; j < B_MD.dimension2; j++) {
			int c_index = c_i + j;
			C[c_index] = 0;
			for (int k = 0; k < B_MD.dimension1; k++) {
				int a_index = a_i + k;
				int b_index = k * B_MD.dimension2 + j;
				C[c_index] += A[a_index] * B[b_index];
			}
		}
	}
}

// do a straightforward matrix-matrix multiplication in the CPU
// notice that this implementation can be massively improved in the CPU by doing proper cache blocking but we are
// not providing you the efficient CPU implementation as that reveals too much about the ideal GPU implementation
void computeGpuMMM() {
	
	// allocate the result matrix for the CPU computation
	size_t sizeofC = C_MD.dimension1 * C_MD.dimension2 * sizeof(float);
	C_GPU = (float*) malloc(sizeofC);
	
	// compute C[i][j] as the sum of A[i][k] * B[k][j] for all columns k of A
	for (int i = 0; i < A_MD.dimension1; i++) {
		int a_i = i * A_MD.dimension2;
		int c_i = i * C_MD.dimension2;
		for (int j = 0; j < B_MD.dimension2; j++) {
			int c_index = c_i + j;
			C_GPU[c_index] = 0;
			for (int k = 0; k < B_MD.dimension1; k++) {
				int a_index = a_i + k;
				int b_index = k * B_MD.dimension2 + j;
				C_GPU[c_index] += A_GPU[a_index] * B_GPU[b_index];
			}
		}
	}
}


// function to determine if the GPU computation is done correctly by comparing the output from the GPU with that
// from the CPU
void compareHostAndGpuOutput() {
	int totalElements = C_MD.dimension1 * C_MD.dimension2;
	int missmatchCount = 0;
	for (int i = 0; i < totalElements; i++) {
        printf("%d\n", i);
        if (fabs(C[i] - C_GPU[i]) > 0.01) {
			missmatchCount++;
			printf("mismatch at index %i: %f\t%f\n", i, C[i], C_CPU[i]);
		}
	}
	if (missmatchCount > 0) {
		printf("Computation is incorrect: outputs do not match in %d indexes\n", missmatchCount);
	} else {
		printf("Computation is correct: CPU and GPU outputs match\n");
	}
}

// Prints the specified error message and then exits
void die(const char *error) {
        printf("%s", error);
        exit(1);
}

// If the specified error code refers to a real error, report it and quit the program
void check_error(cudaError e) {
        if (e != cudaSuccess) {
                printf("\nCUDA error: %s\n", cudaGetErrorString(e));
                exit(1);
        }
}


	.file	"tmpxft_0031ec86_00000000-6_mmm_template.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3646:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3646:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z23allocateAndInitializeABv
	.type	_Z23allocateAndInitializeABv, @function
_Z23allocateAndInitializeABv:
.LFB3636:
	.cfi_startproc
	endbr64
	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
	movl	A_MD(%rip), %ebx
	imull	4+A_MD(%rip), %ebx
	movslq	%ebx, %rbx
	salq	$2, %rbx
	movq	%rbx, %rdi
	call	malloc@PLT
	movq	%rbx, %rdi
	xorl	%ebx, %ebx
	movq	%rax, A(%rip)
	call	malloc@PLT
	xorl	%edi, %edi
	movq	%rax, A_GPU(%rip)
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
.L3:
	cmpl	%ebx, A_MD(%rip)
	jle	.L13
	xorl	%ebp, %ebp
.L5:
	movl	4+A_MD(%rip), %eax
	cmpl	%ebp, %eax
	jle	.L14
	imull	%ebx, %eax
	leal	(%rax,%rbp), %r12d
	incl	%ebp
	call	rand@PLT
	movl	$1000, %ecx
	cltd
	idivl	%ecx
	movslq	%r12d, %rax
	cvtsi2sdl	%edx, %xmm0
	mulsd	.LC0(%rip), %xmm0
	movq	A(%rip), %rdx
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, (%rdx,%rax,4)
	movq	A_GPU(%rip), %rdx
	movss	%xmm0, (%rdx,%rax,4)
	jmp	.L5
.L14:
	incl	%ebx
	jmp	.L3
.L13:
	movl	B_MD(%rip), %ebx
	imull	4+B_MD(%rip), %ebx
	movslq	%ebx, %rbx
	salq	$2, %rbx
	movq	%rbx, %rdi
	call	malloc@PLT
	movq	%rbx, %rdi
	xorl	%ebx, %ebx
	movq	%rax, B(%rip)
	call	malloc@PLT
	movq	%rax, B_GPU(%rip)
.L6:
	cmpl	%ebx, B_MD(%rip)
	jle	.L15
	xorl	%ebp, %ebp
.L8:
	movl	4+B_MD(%rip), %eax
	cmpl	%ebp, %eax
	jle	.L16
	imull	%ebx, %eax
	leal	(%rax,%rbp), %r12d
	incl	%ebp
	call	rand@PLT
	movl	$1000, %ecx
	cltd
	idivl	%ecx
	movslq	%r12d, %rax
	cvtsi2sdl	%edx, %xmm0
	mulsd	.LC0(%rip), %xmm0
	movq	B(%rip), %rdx
	cvtsd2ss	%xmm0, %xmm0
	movss	%xmm0, (%rdx,%rax,4)
	movq	B_GPU(%rip), %rdx
	movss	%xmm0, (%rdx,%rax,4)
	jmp	.L8
.L16:
	incl	%ebx
	jmp	.L6
.L15:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3636:
	.size	_Z23allocateAndInitializeABv, .-_Z23allocateAndInitializeABv
	.globl	_Z13computeCpuMMMv
	.type	_Z13computeCpuMMMv, @function
_Z13computeCpuMMMv:
.LFB3639:
	.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
	pushq	%rdx
	.cfi_def_cfa_offset 64
	movslq	4+C_MD(%rip), %rbx
	movl	C_MD(%rip), %edi
	imull	%ebx, %edi
	salq	$2, %rbx
	movslq	%edi, %rdi
	salq	$2, %rdi
	call	malloc@PLT
	movslq	4+B_MD(%rip), %rsi
	xorl	%ecx, %ecx
	movslq	4+A_MD(%rip), %r8
	movq	%rax, C(%rip)
	movl	A_MD(%rip), %r10d
	movl	B_MD(%rip), %ebp
	movq	B(%rip), %r11
	movq	%rsi, %r9
	salq	$2, %r8
	movq	A(%rip), %rdi
	salq	$2, %rsi
.L18:
	cmpl	%ecx, %r10d
	jle	.L25
	movq	%r11, %r13
	movq	%rax, %r14
	xorl	%r12d, %r12d
.L19:
	cmpl	%r12d, %r9d
	jle	.L26
	movq	%r13, %r15
	xorl	%edx, %edx
	xorps	%xmm0, %xmm0
.L21:
	movss	%xmm0, (%r14)
	cmpl	%edx, %ebp
	jle	.L27
	movss	(%rdi,%rdx,4), %xmm1
	mulss	(%r15), %xmm1
	incq	%rdx
	addq	%rsi, %r15
	addss	%xmm1, %xmm0
	jmp	.L21
.L27:
	incl	%r12d
	addq	$4, %r13
	addq	$4, %r14
	jmp	.L19
.L26:
	incl	%ecx
	addq	%rbx, %rax
	addq	%r8, %rdi
	jmp	.L18
.L25:
	popq	%rax
	.cfi_def_cfa_offset 56
	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
.LFE3639:
	.size	_Z13computeCpuMMMv, .-_Z13computeCpuMMMv
	.globl	_Z13computeGpuMMMv
	.type	_Z13computeGpuMMMv, @function
_Z13computeGpuMMMv:
.LFB3640:
	.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
	pushq	%rdx
	.cfi_def_cfa_offset 64
	movslq	4+C_MD(%rip), %rbx
	movl	C_MD(%rip), %edi
	imull	%ebx, %edi
	salq	$2, %rbx
	movslq	%edi, %rdi
	salq	$2, %rdi
	call	malloc@PLT
	movslq	4+B_MD(%rip), %rsi
	xorl	%ecx, %ecx
	movslq	4+A_MD(%rip), %r8
	movq	%rax, C_GPU(%rip)
	movl	A_MD(%rip), %r10d
	movl	B_MD(%rip), %ebp
	movq	B_GPU(%rip), %r11
	movq	%rsi, %r9
	salq	$2, %r8
	movq	A_GPU(%rip), %rdi
	salq	$2, %rsi
.L29:
	cmpl	%ecx, %r10d
	jle	.L36
	movq	%r11, %r13
	movq	%rax, %r14
	xorl	%r12d, %r12d
.L30:
	cmpl	%r12d, %r9d
	jle	.L37
	movq	%r13, %r15
	xorl	%edx, %edx
	xorps	%xmm0, %xmm0
.L32:
	movss	%xmm0, (%r14)
	cmpl	%edx, %ebp
	jle	.L38
	movss	(%rdi,%rdx,4), %xmm1
	mulss	(%r15), %xmm1
	incq	%rdx
	addq	%rsi, %r15
	addss	%xmm1, %xmm0
	jmp	.L32
.L38:
	incl	%r12d
	addq	$4, %r13
	addq	$4, %r14
	jmp	.L30
.L37:
	incl	%ecx
	addq	%rbx, %rax
	addq	%r8, %rdi
	jmp	.L29
.L36:
	popq	%rax
	.cfi_def_cfa_offset 56
	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
.LFE3640:
	.size	_Z13computeGpuMMMv, .-_Z13computeGpuMMMv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"%d\n"
.LC5:
	.string	"mismatch at index %i: %f\t%f\n"
.LC6:
	.string	"Computation is incorrect: outputs do not match in %d indexes\n"
.LC7:
	.string	"Computation is correct: CPU and GPU outputs match\n"
	.text
	.globl	_Z23compareHostAndGpuOutputv
	.type	_Z23compareHostAndGpuOutputv, @function
_Z23compareHostAndGpuOutputv:
.LFB3641:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	xorl	%ebx, %ebx
	pushq	%rsi
	.cfi_def_cfa_offset 48
	movl	C_MD(%rip), %r12d
	imull	4+C_MD(%rip), %r12d
.L40:
	movl	%ebx, %r13d
	cmpl	%ebx, %r12d
	jle	.L47
	xorl	%eax, %eax
	movl	%r13d, %edx
	movl	$2, %edi
	leaq	.LC2(%rip), %rsi
	call	__printf_chk@PLT
	movq	C(%rip), %rax
	movss	(%rax,%rbx,4), %xmm0
	movq	C_GPU(%rip), %rax
	movaps	%xmm0, %xmm1
	subss	(%rax,%rbx,4), %xmm1
	andps	.LC3(%rip), %xmm1
	cvtss2sd	%xmm1, %xmm1
	comisd	.LC4(%rip), %xmm1
	jbe	.L41
	movq	C_CPU(%rip), %rax
	movl	%r13d, %edx
	movl	$2, %edi
	incl	%ebp
	leaq	.LC5(%rip), %rsi
	cvtss2sd	%xmm0, %xmm0
	cvtss2sd	(%rax,%rbx,4), %xmm1
	movb	$2, %al
	call	__printf_chk@PLT
.L41:
	incq	%rbx
	jmp	.L40
.L47:
	testl	%ebp, %ebp
	jle	.L44
	popq	%rcx
	.cfi_remember_state
	.cfi_def_cfa_offset 40
	movl	%ebp, %edx
	popq	%rbx
	.cfi_def_cfa_offset 32
	movl	$2, %edi
	popq	%rbp
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rsi
	popq	%r12
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
.L44:
	.cfi_restore_state
	popq	%rdx
	.cfi_def_cfa_offset 40
	leaq	.LC7(%rip), %rsi
	popq	%rbx
	.cfi_def_cfa_offset 32
	movl	$2, %edi
	popq	%rbp
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE3641:
	.size	_Z23compareHostAndGpuOutputv, .-_Z23compareHostAndGpuOutputv
	.section	.rodata.str1.1
.LC8:
	.string	"Matrix A is %d-by-%d\n"
.LC9:
	.string	"Matrix B is %d-by-%d\n"
.LC10:
	.string	"Matrix C is %d-by-%d\n"
.LC12:
	.string	"Computation time in the CPU: %f seconds\n"
.LC13:
	.string	"Computation time in the GPU: %f seconds\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	%edi, %r13d
	decl	%edi
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	$100, %ebx
	pushq	%rcx
	.cfi_def_cfa_offset 48
	jle	.L49
	movq	8(%rsi), %rdi
	call	atoi@PLT
	movl	%eax, %ebx
.L49:
	movl	%ebx, A_MD(%rip)
	movl	%ebx, %ebp
	cmpl	$2, %r13d
	jle	.L50
	movq	16(%r12), %rdi
	call	atoi@PLT
	movl	%eax, %ebp
.L50:
	movl	%ebp, 4+A_MD(%rip)
	movl	%ebp, %eax
	cmpl	$3, %r13d
	jle	.L51
	movq	24(%r12), %rdi
	call	atoi@PLT
.L51:
	movl	%eax, B_MD(%rip)
	cmpl	$4, %r13d
	jle	.L52
	movq	32(%r12), %rdi
	call	atoi@PLT
.L52:
	movl	%ebp, %ecx
	movl	%ebx, %edx
	movl	%eax, 4+B_MD(%rip)
	movl	$2, %edi
	movl	%eax, 4+C_MD(%rip)
	leaq	.LC8(%rip), %rsi
	xorl	%eax, %eax
	movl	%ebx, C_MD(%rip)
	call	__printf_chk@PLT
	movl	4+B_MD(%rip), %ecx
	movl	B_MD(%rip), %edx
	xorl	%eax, %eax
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	C_MD(%rip), %edx
	movl	4+C_MD(%rip), %ecx
	xorl	%eax, %eax
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	call	_Z23allocateAndInitializeABv
	call	clock@PLT
	movq	%rax, %rbx
	call	_Z13computeCpuMMMv
	call	clock@PLT
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	subq	%rbx, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC11(%rip), %xmm0
	call	__printf_chk@PLT
	call	clock@PLT
	movq	%rax, %rbx
	call	_Z13computeGpuMMMv
	call	clock@PLT
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	subq	%rbx, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	divsd	.LC11(%rip), %xmm0
	call	__printf_chk@PLT
	call	_Z23compareHostAndGpuOutputv
	popq	%rdx
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3635:
	.size	main, .-main
	.section	.rodata.str1.1
.LC14:
	.string	"%s"
	.text
	.globl	_Z3diePKc
	.type	_Z3diePKc, @function
_Z3diePKc:
.LFB3642:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	leaq	.LC14(%rip), %rsi
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rdi, %rdx
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
	.cfi_endproc
.LFE3642:
	.size	_Z3diePKc, .-_Z3diePKc
	.section	.rodata.str1.1
.LC15:
	.string	"\nCUDA error: %s\n"
	.text
	.globl	_Z11check_error9cudaError
	.type	_Z11check_error9cudaError, @function
_Z11check_error9cudaError:
.LFB3643:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L60
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	cudaGetErrorString@PLT
	movl	$2, %edi
	leaq	.LC15(%rip), %rsi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L60:
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3643:
	.size	_Z11check_error9cudaError, .-_Z11check_error9cudaError
	.globl	_Z17copyMatricesToGPUv
	.type	_Z17copyMatricesToGPUv, @function
_Z17copyMatricesToGPUv:
.LFB3637:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	A_MD(%rip), %ebx
	leaq	A_GPU(%rip), %rdi
	imull	4+A_MD(%rip), %ebx
	movslq	%ebx, %rbx
	salq	$2, %rbx
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movl	%eax, %edi
	call	_Z11check_error9cudaError
	movq	A(%rip), %rsi
	movq	%rbx, %rdx
	movq	A_GPU(%rip), %rdi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	call	_Z11check_error9cudaError
	movl	B_MD(%rip), %ebx
	imull	4+B_MD(%rip), %ebx
	leaq	B_GPU(%rip), %rdi
	movslq	%ebx, %rbx
	salq	$2, %rbx
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movl	%eax, %edi
	call	_Z11check_error9cudaError
	movq	B(%rip), %rsi
	movq	%rbx, %rdx
	movq	B_GPU(%rip), %rdi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	call	_Z11check_error9cudaError
	movl	C_MD(%rip), %esi
	imull	4+C_MD(%rip), %esi
	leaq	C_GPU(%rip), %rdi
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	popq	%rbx
	.cfi_def_cfa_offset 8
	movl	%eax, %edi
	jmp	_Z11check_error9cudaError
	.cfi_endproc
.LFE3637:
	.size	_Z17copyMatricesToGPUv, .-_Z17copyMatricesToGPUv
	.globl	_Z17copyResultFromGPUv
	.type	_Z17copyResultFromGPUv, @function
_Z17copyResultFromGPUv:
.LFB3638:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	C_MD(%rip), %ebx
	imull	4+C_MD(%rip), %ebx
	movslq	%ebx, %rbx
	salq	$2, %rbx
	movq	%rbx, %rdi
	call	malloc@PLT
	movq	C_GPU(%rip), %rsi
	movq	%rbx, %rdx
	movl	$2, %ecx
	movq	%rax, %rdi
	movq	%rax, C_CPU(%rip)
	call	cudaMemcpy@PLT
	popq	%rbx
	.cfi_def_cfa_offset 8
	movl	%eax, %edi
	jmp	_Z11check_error9cudaError
	.cfi_endproc
.LFE3638:
	.size	_Z17copyResultFromGPUv, .-_Z17copyResultFromGPUv
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3669:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3669:
	.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	C_GPU
	.bss
	.align 8
	.type	C_GPU, @object
	.size	C_GPU, 8
C_GPU:
	.zero	8
	.globl	B_GPU
	.align 8
	.type	B_GPU, @object
	.size	B_GPU, 8
B_GPU:
	.zero	8
	.globl	A_GPU
	.align 8
	.type	A_GPU, @object
	.size	A_GPU, 8
A_GPU:
	.zero	8
	.globl	C_CPU
	.align 8
	.type	C_CPU, @object
	.size	C_CPU, 8
C_CPU:
	.zero	8
	.globl	C
	.align 8
	.type	C, @object
	.size	C, 8
C:
	.zero	8
	.globl	B
	.align 8
	.type	B, @object
	.size	B, 8
B:
	.zero	8
	.globl	A
	.align 8
	.type	A, @object
	.size	A, 8
A:
	.zero	8
	.globl	C_MD
	.align 8
	.type	C_MD, @object
	.size	C_MD, 8
C_MD:
	.zero	8
	.globl	B_MD
	.align 8
	.type	B_MD, @object
	.size	B_MD, 8
B_MD:
	.zero	8
	.globl	A_MD
	.align 8
	.type	A_MD, @object
	.size	A_MD, 8
A_MD:
	.zero	8
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.long	-755914244
	.long	1062232653
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC3:
	.long	2147483647
	.long	0
	.long	0
	.long	0
	.section	.rodata.cst8
	.align 8
.LC4:
	.long	1202590843
	.long	1065646817
	.align 8
.LC11:
	.long	0
	.long	1093567616
	.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:


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

using namespace std;

//----------------------------------- Structures and Globals---------------------------------------------

typedef struct {
	int dimension1;
	int dimension2;	
} ArrayMetadata2D;

// metadata variables describing dimensionalities of all data structures involved in the computation
ArrayMetadata2D A_MD, B_MD, C_MD;
// pointers for input and output arrays in the host memory  
float *A, *B, *C, *C_CPU;
// pointers for input and output arrays in the device memory (NVIDIA DRAM)
float *A_GPU, *B_GPU, *C_GPU;

//----------------------------------- host function definitions -----------------------------------------

void allocateAndInitializeAB();
void computeCpuMMM();
void computeGpuMMM();
void copyMatricesToGPU();
void copyResultFromGPU();
void compareHostAndGpuOutput();
void die(const char *error); 
void check_error(hipError_t e);

//----------------------------------- CUDA function definitions -----------------------------------------


//-------------------------------------------------------------------------------------------------------
int main(int argc, char **argv) {
	
	A_MD.dimension1 = (argc > 1) ? atoi(argv[1]) : 100;
	A_MD.dimension2 = (argc > 2) ? atoi(argv[2]) : A_MD.dimension1;
	B_MD.dimension1 = (argc > 3) ? atoi(argv[3]) : A_MD.dimension2;
	B_MD.dimension2 = (argc > 4) ? atoi(argv[4]) : B_MD.dimension1;
	C_MD.dimension1 = A_MD.dimension1;
	C_MD.dimension2 = B_MD.dimension2;

	printf("Matrix A is %d-by-%d\n", A_MD.dimension1, A_MD.dimension2);
	printf("Matrix B is %d-by-%d\n", B_MD.dimension1, B_MD.dimension2);
	printf("Matrix C is %d-by-%d\n", C_MD.dimension1, C_MD.dimension2);

	allocateAndInitializeAB();

	// matrix matrix multiplication in the CPU
	clock_t start = clock();	
	computeCpuMMM();
	clock_t end = clock();
    
    double elapsed = (end - start) / (double) CLOCKS_PER_SEC;
    printf("Computation time in the CPU: %f seconds\n", elapsed);

    start = clock();
    computeGpuMMM();
    end = clock();

    elapsed = (end - start) / (double) CLOCKS_PER_SEC;
    printf("Computation time in the GPU: %f seconds\n", elapsed);

    compareHostAndGpuOutput();

	return 0;
}

// allocate and initialize A and B using a random number generator
void allocateAndInitializeAB() {
	
    float val;
	
    size_t sizeofA = A_MD.dimension1 * A_MD.dimension2 * sizeof(float);
	A = (float*) malloc(sizeofA);
    A_GPU = (float*) malloc(sizeofA);

	srand(time(NULL));
  	for (int i = 0; i < A_MD.dimension1; i++) {
		for (int j = 0; j < A_MD.dimension2; j++) {
			int index = i * A_MD.dimension2 + j;
            val = (rand() % 1000) * 0.001;
			A[index] = val; 
            A_GPU[index] = val;
		}
	}
	
	size_t sizeofB = B_MD.dimension1 * B_MD.dimension2 * sizeof(float);
	B = (float*) malloc(sizeofB);
	B_GPU = (float*) malloc(sizeofB);
  	for (int i = 0; i < B_MD.dimension1; i++) {
		for (int j = 0; j < B_MD.dimension2; j++) {
			int index = i * B_MD.dimension2 + j;
            val = (rand() % 1000) * 0.001;
			B[index] = val; 
            B_GPU[index] = val;
		}
	}
}

// allocate memory in the GPU for all matrices, and copy A and B content from the host CPU memory to the GPU memory
void copyMatricesToGPU() {
	
	size_t sizeofA = A_MD.dimension1 * A_MD.dimension2 * sizeof(float);
	check_error(hipMalloc((void **) &A_GPU, sizeofA));
	check_error(hipMemcpy(A_GPU, A, sizeofA, hipMemcpyHostToDevice));
	
	size_t sizeofB = B_MD.dimension1 * B_MD.dimension2 * sizeof(float);
	check_error(hipMalloc((void **) &B_GPU, sizeofB));
	check_error(hipMemcpy(B_GPU, B, sizeofB, hipMemcpyHostToDevice));
	
	size_t sizeofC = C_MD.dimension1 * C_MD.dimension2 * sizeof(float);
	check_error(hipMalloc((void **) &C_GPU, sizeofC));
}

// copy results from C_GPU which is in GPU card memory to C_CPU which is in the host CPU for result comparison
void copyResultFromGPU() {
	size_t sizeofC = C_MD.dimension1 * C_MD.dimension2 * sizeof(float);
	C_CPU = (float*) malloc(sizeofC);
	check_error(hipMemcpy(C_CPU, C_GPU, sizeofC, hipMemcpyDeviceToHost));
}

// do a straightforward matrix-matrix multiplication in the CPU
// notice that this implementation can be massively improved in the CPU by doing proper cache blocking but we are
// not providing you the efficient CPU implementation as that reveals too much about the ideal GPU implementation
void computeCpuMMM() {
	
	// allocate the result matrix for the CPU computation
	size_t sizeofC = C_MD.dimension1 * C_MD.dimension2 * sizeof(float);
	C = (float*) malloc(sizeofC);
	
	// compute C[i][j] as the sum of A[i][k] * B[k][j] for all columns k of A
	for (int i = 0; i < A_MD.dimension1; i++) {
		int a_i = i * A_MD.dimension2;
		int c_i = i * C_MD.dimension2;
		for (int j = 0; j < B_MD.dimension2; j++) {
			int c_index = c_i + j;
			C[c_index] = 0;
			for (int k = 0; k < B_MD.dimension1; k++) {
				int a_index = a_i + k;
				int b_index = k * B_MD.dimension2 + j;
				C[c_index] += A[a_index] * B[b_index];
			}
		}
	}
}

// do a straightforward matrix-matrix multiplication in the CPU
// notice that this implementation can be massively improved in the CPU by doing proper cache blocking but we are
// not providing you the efficient CPU implementation as that reveals too much about the ideal GPU implementation
void computeGpuMMM() {
	
	// allocate the result matrix for the CPU computation
	size_t sizeofC = C_MD.dimension1 * C_MD.dimension2 * sizeof(float);
	C_GPU = (float*) malloc(sizeofC);
	
	// compute C[i][j] as the sum of A[i][k] * B[k][j] for all columns k of A
	for (int i = 0; i < A_MD.dimension1; i++) {
		int a_i = i * A_MD.dimension2;
		int c_i = i * C_MD.dimension2;
		for (int j = 0; j < B_MD.dimension2; j++) {
			int c_index = c_i + j;
			C_GPU[c_index] = 0;
			for (int k = 0; k < B_MD.dimension1; k++) {
				int a_index = a_i + k;
				int b_index = k * B_MD.dimension2 + j;
				C_GPU[c_index] += A_GPU[a_index] * B_GPU[b_index];
			}
		}
	}
}


// function to determine if the GPU computation is done correctly by comparing the output from the GPU with that
// from the CPU
void compareHostAndGpuOutput() {
	int totalElements = C_MD.dimension1 * C_MD.dimension2;
	int missmatchCount = 0;
	for (int i = 0; i < totalElements; i++) {
        printf("%d\n", i);
        if (fabs(C[i] - C_GPU[i]) > 0.01) {
			missmatchCount++;
			printf("mismatch at index %i: %f\t%f\n", i, C[i], C_CPU[i]);
		}
	}
	if (missmatchCount > 0) {
		printf("Computation is incorrect: outputs do not match in %d indexes\n", missmatchCount);
	} else {
		printf("Computation is correct: CPU and GPU outputs match\n");
	}
}

// Prints the specified error message and then exits
void die(const char *error) {
        printf("%s", error);
        exit(1);
}

// If the specified error code refers to a real error, report it and quit the program
void check_error(hipError_t e) {
        if (e != hipSuccess) {
                printf("\nCUDA error: %s\n", hipGetErrorString(e));
                exit(1);
        }
}


	.text
	.file	"mmm_template.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI0_0:
	.quad	0x412e848000000000              # double 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	%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
	cmpl	$1, %edi
	jle	.LBB0_1
# %bb.2:
	movq	%rsi, %r14
	movl	%edi, %r15d
	movq	8(%rsi), %rdi
	callq	atoi
	movl	%eax, %ebx
	movl	%eax, A_MD(%rip)
	cmpl	$2, %r15d
	je	.LBB0_3
# %bb.4:
	movq	16(%r14), %rdi
	callq	atoi
	movl	%eax, %ebp
	movl	%eax, A_MD+4(%rip)
	cmpl	$3, %r15d
	jbe	.LBB0_5
# %bb.6:
	movq	24(%r14), %rdi
	callq	atoi
	movl	%eax, B_MD(%rip)
	cmpl	$4, %r15d
	je	.LBB0_8
# %bb.7:
	movq	32(%r14), %rdi
	callq	atoi
	jmp	.LBB0_8
.LBB0_1:                                # %.thread
	movl	$100, A_MD(%rip)
	movl	$100, %ebx
.LBB0_3:                                # %.thread14
	movl	%ebx, A_MD+4(%rip)
	movl	%ebx, %ebp
.LBB0_5:                                # %.thread19
	movl	%ebp, B_MD(%rip)
	movl	%ebp, %eax
.LBB0_8:
	movl	%eax, B_MD+4(%rip)
	movl	%ebx, C_MD(%rip)
	movl	%eax, C_MD+4(%rip)
	movl	$.L.str, %edi
	movl	%ebx, %esi
	movl	%ebp, %edx
	xorl	%eax, %eax
	callq	printf
	movl	B_MD(%rip), %esi
	movl	B_MD+4(%rip), %edx
	movl	$.L.str.1, %edi
	xorl	%eax, %eax
	callq	printf
	movl	C_MD(%rip), %esi
	movl	C_MD+4(%rip), %edx
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
	callq	_Z23allocateAndInitializeABv
	callq	clock
	movq	%rax, %rbx
	callq	_Z13computeCpuMMMv
	callq	clock
	subq	%rbx, %rax
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI0_0(%rip), %xmm0
	movl	$.L.str.3, %edi
	movb	$1, %al
	callq	printf
	callq	clock
	movq	%rax, %rbx
	callq	_Z13computeGpuMMMv
	callq	clock
	subq	%rbx, %rax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI0_0(%rip), %xmm0
	movl	$.L.str.4, %edi
	movb	$1, %al
	callq	printf
	callq	_Z23compareHostAndGpuOutputv
	xorl	%eax, %eax
	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
	retq
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z23allocateAndInitializeABv
.LCPI1_0:
	.quad	0x3f50624dd2f1a9fc              # double 0.001
	.text
	.globl	_Z23allocateAndInitializeABv
	.type	_Z23allocateAndInitializeABv,@function
_Z23allocateAndInitializeABv:           # @_Z23allocateAndInitializeABv
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movslq	A_MD(%rip), %rax
	movslq	A_MD+4(%rip), %rbx
	imulq	%rax, %rbx
	shlq	$2, %rbx
	movq	%rbx, %rdi
	callq	malloc
	movq	%rax, A(%rip)
	movq	%rbx, %rdi
	callq	malloc
	movq	%rax, A_GPU(%rip)
	xorl	%ebx, %ebx
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	movl	A_MD(%rip), %eax
	testl	%eax, %eax
	jle	.LBB1_7
# %bb.1:                                # %.preheader28.preheader
	movl	A_MD+4(%rip), %ebp
.LBB1_2:                                # %.preheader28
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_4 Depth 2
	testl	%ebp, %ebp
	jle	.LBB1_6
# %bb.3:                                # %.lr.ph.preheader
                                        #   in Loop: Header=BB1_2 Depth=1
	xorl	%r14d, %r14d
.LBB1_4:                                # %.lr.ph
                                        #   Parent Loop BB1_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	imull	%ebx, %ebp
	addl	%r14d, %ebp
	callq	rand
	movsd	.LCPI1_0(%rip), %xmm1           # xmm1 = mem[0],zero
	cltq
	imulq	$274877907, %rax, %rcx          # imm = 0x10624DD3
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$38, %rcx
	addl	%edx, %ecx
	imull	$1000, %ecx, %ecx               # imm = 0x3E8
	subl	%ecx, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	mulsd	%xmm1, %xmm0
	cvtsd2ss	%xmm0, %xmm0
	movq	A(%rip), %rax
	movslq	%ebp, %rcx
	movss	%xmm0, (%rax,%rcx,4)
	movq	A_GPU(%rip), %rax
	movss	%xmm0, (%rax,%rcx,4)
	incl	%r14d
	movl	A_MD+4(%rip), %ebp
	cmpl	%ebp, %r14d
	jl	.LBB1_4
# %bb.5:                                # %._crit_edge.loopexit
                                        #   in Loop: Header=BB1_2 Depth=1
	movl	A_MD(%rip), %eax
.LBB1_6:                                # %._crit_edge
                                        #   in Loop: Header=BB1_2 Depth=1
	incl	%ebx
	cmpl	%eax, %ebx
	jl	.LBB1_2
.LBB1_7:                                # %._crit_edge31
	movl	B_MD(%rip), %r14d
	movl	B_MD+4(%rip), %ebp
	movl	%ebp, %eax
	imull	%r14d, %eax
	movslq	%eax, %rbx
	shlq	$2, %rbx
	movq	%rbx, %rdi
	callq	malloc
	movq	%rax, B(%rip)
	movq	%rbx, %rdi
	callq	malloc
	movq	%rax, B_GPU(%rip)
	testl	%r14d, %r14d
	jle	.LBB1_14
# %bb.8:                                # %.preheader.preheader
	xorl	%ebx, %ebx
.LBB1_9:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_11 Depth 2
	testl	%ebp, %ebp
	jle	.LBB1_13
# %bb.10:                               # %.lr.ph33.preheader
                                        #   in Loop: Header=BB1_9 Depth=1
	xorl	%r14d, %r14d
.LBB1_11:                               # %.lr.ph33
                                        #   Parent Loop BB1_9 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	imull	%ebx, %ebp
	addl	%r14d, %ebp
	callq	rand
	movsd	.LCPI1_0(%rip), %xmm1           # xmm1 = mem[0],zero
	cltq
	imulq	$274877907, %rax, %rcx          # imm = 0x10624DD3
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$38, %rcx
	addl	%edx, %ecx
	imull	$1000, %ecx, %ecx               # imm = 0x3E8
	subl	%ecx, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	mulsd	%xmm1, %xmm0
	cvtsd2ss	%xmm0, %xmm0
	movq	B(%rip), %rax
	movslq	%ebp, %rcx
	movss	%xmm0, (%rax,%rcx,4)
	movq	B_GPU(%rip), %rax
	movss	%xmm0, (%rax,%rcx,4)
	incl	%r14d
	movl	B_MD+4(%rip), %ebp
	cmpl	%ebp, %r14d
	jl	.LBB1_11
# %bb.12:                               # %._crit_edge34.loopexit
                                        #   in Loop: Header=BB1_9 Depth=1
	movl	B_MD(%rip), %r14d
.LBB1_13:                               # %._crit_edge34
                                        #   in Loop: Header=BB1_9 Depth=1
	incl	%ebx
	cmpl	%r14d, %ebx
	jl	.LBB1_9
.LBB1_14:                               # %._crit_edge36
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z23allocateAndInitializeABv, .Lfunc_end1-_Z23allocateAndInitializeABv
	.cfi_endproc
                                        # -- End function
	.globl	_Z13computeCpuMMMv              # -- Begin function _Z13computeCpuMMMv
	.type	_Z13computeCpuMMMv,@function
_Z13computeCpuMMMv:                     # @_Z13computeCpuMMMv
	.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
	movslq	C_MD(%rip), %rdi
	movslq	C_MD+4(%rip), %rbx
	imulq	%rbx, %rdi
	shlq	$2, %rdi
	callq	malloc
	movq	%rax, C(%rip)
	movl	A_MD(%rip), %ecx
	testl	%ecx, %ecx
	jle	.LBB2_9
# %bb.1:                                # %.lr.ph28
	movslq	A_MD+4(%rip), %rdx
	movl	B_MD+4(%rip), %esi
	movl	B_MD(%rip), %edi
	movq	A(%rip), %r8
	movq	B(%rip), %r9
	movslq	%esi, %r10
	shlq	$2, %rdx
	shlq	$2, %r10
	xorl	%r11d, %r11d
.LBB2_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_4 Depth 2
                                        #       Child Loop BB2_6 Depth 3
	testl	%esi, %esi
	jle	.LBB2_8
# %bb.3:                                # %.lr.ph24.preheader
                                        #   in Loop: Header=BB2_2 Depth=1
	movq	%r11, %r14
	imulq	%rbx, %r14
	leaq	(%rax,%r14,4), %r14
	movq	%r9, %r15
	xorl	%r12d, %r12d
.LBB2_4:                                # %.lr.ph24
                                        #   Parent Loop BB2_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB2_6 Depth 3
	movl	$0, (%r14,%r12,4)
	testl	%edi, %edi
	jle	.LBB2_7
# %bb.5:                                # %.lr.ph
                                        #   in Loop: Header=BB2_4 Depth=2
	xorps	%xmm0, %xmm0
	movq	%r15, %r13
	xorl	%ebp, %ebp
.LBB2_6:                                #   Parent Loop BB2_2 Depth=1
                                        #     Parent Loop BB2_4 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movss	(%r8,%rbp,4), %xmm1             # xmm1 = mem[0],zero,zero,zero
	mulss	(%r13), %xmm1
	addss	%xmm1, %xmm0
	movss	%xmm0, (%r14,%r12,4)
	incq	%rbp
	addq	%r10, %r13
	cmpq	%rbp, %rdi
	jne	.LBB2_6
.LBB2_7:                                # %._crit_edge
                                        #   in Loop: Header=BB2_4 Depth=2
	incq	%r12
	addq	$4, %r15
	cmpq	%rsi, %r12
	jne	.LBB2_4
.LBB2_8:                                # %._crit_edge25
                                        #   in Loop: Header=BB2_2 Depth=1
	incq	%r11
	addq	%rdx, %r8
	cmpq	%rcx, %r11
	jne	.LBB2_2
.LBB2_9:                                # %._crit_edge29
	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	_Z13computeCpuMMMv, .Lfunc_end2-_Z13computeCpuMMMv
	.cfi_endproc
                                        # -- End function
	.globl	_Z13computeGpuMMMv              # -- Begin function _Z13computeGpuMMMv
	.type	_Z13computeGpuMMMv,@function
_Z13computeGpuMMMv:                     # @_Z13computeGpuMMMv
	.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
	movslq	C_MD(%rip), %rdi
	movslq	C_MD+4(%rip), %rbx
	imulq	%rbx, %rdi
	shlq	$2, %rdi
	callq	malloc
	movq	%rax, C_GPU(%rip)
	movl	A_MD(%rip), %ecx
	testl	%ecx, %ecx
	jle	.LBB3_9
# %bb.1:                                # %.lr.ph28
	movslq	A_MD+4(%rip), %rdx
	movl	B_MD+4(%rip), %esi
	movl	B_MD(%rip), %edi
	movq	A_GPU(%rip), %r8
	movq	B_GPU(%rip), %r9
	movslq	%esi, %r10
	shlq	$2, %rdx
	shlq	$2, %r10
	xorl	%r11d, %r11d
.LBB3_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_4 Depth 2
                                        #       Child Loop BB3_6 Depth 3
	testl	%esi, %esi
	jle	.LBB3_8
# %bb.3:                                # %.lr.ph24.preheader
                                        #   in Loop: Header=BB3_2 Depth=1
	movq	%r11, %r14
	imulq	%rbx, %r14
	leaq	(%rax,%r14,4), %r14
	movq	%r9, %r15
	xorl	%r12d, %r12d
.LBB3_4:                                # %.lr.ph24
                                        #   Parent Loop BB3_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB3_6 Depth 3
	movl	$0, (%r14,%r12,4)
	testl	%edi, %edi
	jle	.LBB3_7
# %bb.5:                                # %.lr.ph
                                        #   in Loop: Header=BB3_4 Depth=2
	xorps	%xmm0, %xmm0
	movq	%r15, %r13
	xorl	%ebp, %ebp
.LBB3_6:                                #   Parent Loop BB3_2 Depth=1
                                        #     Parent Loop BB3_4 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movss	(%r8,%rbp,4), %xmm1             # xmm1 = mem[0],zero,zero,zero
	mulss	(%r13), %xmm1
	addss	%xmm1, %xmm0
	movss	%xmm0, (%r14,%r12,4)
	incq	%rbp
	addq	%r10, %r13
	cmpq	%rbp, %rdi
	jne	.LBB3_6
.LBB3_7:                                # %._crit_edge
                                        #   in Loop: Header=BB3_4 Depth=2
	incq	%r12
	addq	$4, %r15
	cmpq	%rsi, %r12
	jne	.LBB3_4
.LBB3_8:                                # %._crit_edge25
                                        #   in Loop: Header=BB3_2 Depth=1
	incq	%r11
	addq	%rdx, %r8
	cmpq	%rcx, %r11
	jne	.LBB3_2
.LBB3_9:                                # %._crit_edge29
	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	_Z13computeGpuMMMv, .Lfunc_end3-_Z13computeGpuMMMv
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _Z23compareHostAndGpuOutputv
.LCPI4_0:
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI4_1:
	.quad	0x3f847ae147ae147b              # double 0.01
	.text
	.globl	_Z23compareHostAndGpuOutputv
	.type	_Z23compareHostAndGpuOutputv,@function
_Z23compareHostAndGpuOutputv:           # @_Z23compareHostAndGpuOutputv
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	C_MD+4(%rip), %eax
	imull	C_MD(%rip), %eax
	testl	%eax, %eax
	jle	.LBB4_7
# %bb.1:                                # %.lr.ph.preheader
	movl	%eax, %r15d
	xorl	%r14d, %r14d
	xorl	%ebx, %ebx
.LBB4_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	$.L.str.5, %edi
	movl	%r14d, %esi
	xorl	%eax, %eax
	callq	printf
	movq	C(%rip), %rax
	movss	(%rax,%r14,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	movq	C_GPU(%rip), %rax
	movaps	%xmm0, %xmm1
	subss	(%rax,%r14,4), %xmm1
	andps	.LCPI4_0(%rip), %xmm1
	cvtss2sd	%xmm1, %xmm1
	ucomisd	.LCPI4_1(%rip), %xmm1
	jbe	.LBB4_4
# %bb.3:                                #   in Loop: Header=BB4_2 Depth=1
	incl	%ebx
	cvtss2sd	%xmm0, %xmm0
	movq	C_CPU(%rip), %rax
	xorps	%xmm1, %xmm1
	cvtss2sd	(%rax,%r14,4), %xmm1
	movl	$.L.str.6, %edi
	movl	%r14d, %esi
	movb	$2, %al
	callq	printf
.LBB4_4:                                #   in Loop: Header=BB4_2 Depth=1
	incq	%r14
	cmpq	%r14, %r15
	jne	.LBB4_2
# %bb.5:                                # %._crit_edge
	testl	%ebx, %ebx
	jle	.LBB4_7
# %bb.6:
	movl	$.L.str.7, %edi
	movl	%ebx, %esi
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	printf                          # TAILCALL
.LBB4_7:                                # %._crit_edge.thread
	.cfi_def_cfa_offset 32
	movl	$.Lstr, %edi
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	puts@PLT                        # TAILCALL
.Lfunc_end4:
	.size	_Z23compareHostAndGpuOutputv, .Lfunc_end4-_Z23compareHostAndGpuOutputv
	.cfi_endproc
                                        # -- End function
	.globl	_Z17copyMatricesToGPUv          # -- Begin function _Z17copyMatricesToGPUv
	.type	_Z17copyMatricesToGPUv,@function
_Z17copyMatricesToGPUv:                 # @_Z17copyMatricesToGPUv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movslq	A_MD(%rip), %rax
	movslq	A_MD+4(%rip), %rbx
	imulq	%rax, %rbx
	shlq	$2, %rbx
	movl	$A_GPU, %edi
	movq	%rbx, %rsi
	callq	hipMalloc
	movl	%eax, %edi
	callq	_Z11check_error10hipError_t
	movq	A_GPU(%rip), %rdi
	movq	A(%rip), %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	%eax, %edi
	callq	_Z11check_error10hipError_t
	movslq	B_MD(%rip), %rax
	movslq	B_MD+4(%rip), %rbx
	imulq	%rax, %rbx
	shlq	$2, %rbx
	movl	$B_GPU, %edi
	movq	%rbx, %rsi
	callq	hipMalloc
	movl	%eax, %edi
	callq	_Z11check_error10hipError_t
	movq	B_GPU(%rip), %rdi
	movq	B(%rip), %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	%eax, %edi
	callq	_Z11check_error10hipError_t
	movslq	C_MD(%rip), %rax
	movslq	C_MD+4(%rip), %rsi
	imulq	%rax, %rsi
	shlq	$2, %rsi
	movl	$C_GPU, %edi
	callq	hipMalloc
	movl	%eax, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	_Z11check_error10hipError_t     # TAILCALL
.Lfunc_end5:
	.size	_Z17copyMatricesToGPUv, .Lfunc_end5-_Z17copyMatricesToGPUv
	.cfi_endproc
                                        # -- End function
	.globl	_Z11check_error10hipError_t     # -- Begin function _Z11check_error10hipError_t
	.type	_Z11check_error10hipError_t,@function
_Z11check_error10hipError_t:            # @_Z11check_error10hipError_t
	.cfi_startproc
# %bb.0:
	testl	%edi, %edi
	jne	.LBB6_2
# %bb.1:
	retq
.LBB6_2:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	hipGetErrorString
	movl	$.L.str.10, %edi
	movq	%rax, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	callq	exit
.Lfunc_end6:
	.size	_Z11check_error10hipError_t, .Lfunc_end6-_Z11check_error10hipError_t
	.cfi_endproc
                                        # -- End function
	.globl	_Z17copyResultFromGPUv          # -- Begin function _Z17copyResultFromGPUv
	.type	_Z17copyResultFromGPUv,@function
_Z17copyResultFromGPUv:                 # @_Z17copyResultFromGPUv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movslq	C_MD(%rip), %rax
	movslq	C_MD+4(%rip), %rbx
	imulq	%rax, %rbx
	shlq	$2, %rbx
	movq	%rbx, %rdi
	callq	malloc
	movq	%rax, C_CPU(%rip)
	movq	C_GPU(%rip), %rsi
	movq	%rax, %rdi
	movq	%rbx, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movl	%eax, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	_Z11check_error10hipError_t     # TAILCALL
.Lfunc_end7:
	.size	_Z17copyResultFromGPUv, .Lfunc_end7-_Z17copyResultFromGPUv
	.cfi_endproc
                                        # -- End function
	.globl	_Z3diePKc                       # -- Begin function _Z3diePKc
	.type	_Z3diePKc,@function
_Z3diePKc:                              # @_Z3diePKc
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rdi, %rsi
	movl	$.L.str.9, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	callq	exit
.Lfunc_end8:
	.size	_Z3diePKc, .Lfunc_end8-_Z3diePKc
	.cfi_endproc
                                        # -- End function
	.type	A_MD,@object                    # @A_MD
	.bss
	.globl	A_MD
	.p2align	2, 0x0
A_MD:
	.zero	8
	.size	A_MD, 8

	.type	B_MD,@object                    # @B_MD
	.globl	B_MD
	.p2align	2, 0x0
B_MD:
	.zero	8
	.size	B_MD, 8

	.type	C_MD,@object                    # @C_MD
	.globl	C_MD
	.p2align	2, 0x0
C_MD:
	.zero	8
	.size	C_MD, 8

	.type	A,@object                       # @A
	.globl	A
	.p2align	3, 0x0
A:
	.quad	0
	.size	A, 8

	.type	B,@object                       # @B
	.globl	B
	.p2align	3, 0x0
B:
	.quad	0
	.size	B, 8

	.type	C,@object                       # @C
	.globl	C
	.p2align	3, 0x0
C:
	.quad	0
	.size	C, 8

	.type	C_CPU,@object                   # @C_CPU
	.globl	C_CPU
	.p2align	3, 0x0
C_CPU:
	.quad	0
	.size	C_CPU, 8

	.type	A_GPU,@object                   # @A_GPU
	.globl	A_GPU
	.p2align	3, 0x0
A_GPU:
	.quad	0
	.size	A_GPU, 8

	.type	B_GPU,@object                   # @B_GPU
	.globl	B_GPU
	.p2align	3, 0x0
B_GPU:
	.quad	0
	.size	B_GPU, 8

	.type	C_GPU,@object                   # @C_GPU
	.globl	C_GPU
	.p2align	3, 0x0
C_GPU:
	.quad	0
	.size	C_GPU, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Matrix A is %d-by-%d\n"
	.size	.L.str, 22

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Matrix B is %d-by-%d\n"
	.size	.L.str.1, 22

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Matrix C is %d-by-%d\n"
	.size	.L.str.2, 22

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Computation time in the CPU: %f seconds\n"
	.size	.L.str.3, 41

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Computation time in the GPU: %f seconds\n"
	.size	.L.str.4, 41

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

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"mismatch at index %i: %f\t%f\n"
	.size	.L.str.6, 29

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Computation is incorrect: outputs do not match in %d indexes\n"
	.size	.L.str.7, 62

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"%s"
	.size	.L.str.9, 3

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"\nCUDA error: %s\n"
	.size	.L.str.10, 17

	.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	"Computation is correct: CPU and GPU outputs match"
	.size	.Lstr, 50

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym A_GPU
	.addrsig_sym B_GPU
	.addrsig_sym C_GPU
	.addrsig_sym __hip_cuid_

	.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_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
