#include <iostream>
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
#include <math.h>
#include <functional>
#include <stdlib.h>     /* srand, rand */
#include <time.h>       /* time */
#include <chrono>
#include <list>
#include <cuda.h>
#include <cstdlib>
#include <ctime>
#include <cmath>
using namespace std;
#define ROW_TILE_WIDTH  64
#define COL_TILE_WIDTH  64

#define EPSILON         (1e-6)

template<typename T>
__global__
void naive_matrix_multiply(T* A, T* B, T* C, int width, int cRows, int cCols)
{
    __shared__ T shATile[ROW_TILE_WIDTH][COL_TILE_WIDTH];
    __shared__ T shBTile[ROW_TILE_WIDTH][COL_TILE_WIDTH];

    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    T pValue = 0;

    // iterate for width/COL_TILE_WIDTH number of times
    // to compute the C tile
    for (int p = 0; p < width / COL_TILE_WIDTH; p++) {
        //load values to tiles from A and B
        shATile[threadIdx.y][threadIdx.x] = A[row * width + p * ROW_TILE_WIDTH + threadIdx.x];
        shBTile[threadIdx.y][threadIdx.x] = B[(p * COL_TILE_WIDTH + threadIdx.y) * cCols + col];

        // wait until all threads finish loading values
        __syncthreads();
        // update pValue for this thread
        for (int i = 0; i < COL_TILE_WIDTH; i++) pValue += shATile[threadIdx.y][i] * shBTile[i][threadIdx.x];
        // wait until all threads finish computing pValue before overwriting
        __syncthreads();

    }

    C[row * cCols + col] = pValue;

}

template<typename T>
void initialize_matrix(T* M, int rows, int cols, std::function<float()> F) {
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            M[i * cols + j] = F();
        }
    }
}

template<typename T>
void initialize_matrix(T* M, int rows, int cols, std::function<float(int, int)> F) {
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            M[i * cols + j] = F(i, j);
        }
    }
}

template<typename T>
void print_matrix(T* M, int rows, int cols) {
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            std::cout << M[i * cols + j] << " ";
        }
        std::cout << "\n";
    }
}



template<typename T>
void naive_matrix_multiply_cpu(T* A, T* B, T* C, int width, int C_rows, int C_cols) {

    for (int i = 0; i < C_rows; i++)
        for (int j = 0; j < C_cols; j++) {
            T value = 0.0f;
            for (int k = 0; k < width; k++) {
                value += A[i * width + k] * B[k * C_cols + j];
            }


            C[i * C_cols + j] = value;
        }
}

template<typename T>
T maxDiff(T* A1, T* A2, int rows, int cols) {
    T maxDiff = A1[0] - A2[0];
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            T diff = abs(A1[i * cols + j] - A2[i * cols + j]);
            if (diff > maxDiff) {
                maxDiff = diff;
            }
        }
    }


    return maxDiff;
}


int main(void)
{
    FILE* fp;
    fp = fopen("Naive_Multiplication_CPU_Mohit.csv", "w+");
    fprintf(fp, "Algorithm_Name,Input_Dimensions,Execution_Time(ms)");
    int max_limit = 800000;
    int my_list[] = {16, 32, 48, 64, 80, 96, 112, 128, 144, 160, 176, 192, 208, 224, 240, 256, 272, 288, 304, 320, 336, 352, 368, 384, 400, 416, 432, 448, 464, 480, 496, 512, 528, 544, 560, 576, 592, 608, 624, 640, 656, 672, 688, 704, 720, 736, 752, 768, 784, 800, 816, 832, 848, 864, 880, 896, 912, 928, 944, 960, 976, 992, 1008, 1040, 1072, 1104, 1136, 1168, 1200, 1232, 1264, 1296, 1328, 1360, 1392, 1424, 1456, 1488, 1520, 1552, 1584, 1616, 1648, 1680, 1712, 1744, 1776, 1808, 1840, 1872, 1904, 1936, 1968, 2000, 2032, 2064, 2096, 2128, 2160, 2192, 2224, 2256, 2288, 2320, 2352, 2384, 2416, 2448, 2480, 2512, 2544, 2576, 2608, 2640, 2672, 2704, 2736, 2768, 2800, 2832, 2864, 2896, 2928, 2960, 2992, 3024, 3056, 3088, 3120, 3152, 3184, 3216, 3248, 3280, 3312, 3344, 3376, 3408, 3440, 3472, 3504, 3536, 3568, 3600, 3632, 3664, 3696, 3728, 3760, 3792, 3824, 3856, 3888, 3920, 3952, 3984, 4016, 4080, 4144, 4208, 4272, 4336, 4400, 4464, 4528, 4592, 4656, 4720, 4784, 4848, 4912,4976, 5040, 5104, 5168, 5232, 5296, 5360, 5424, 5488, 5552, 5616, 5680, 5744, 5808, 5872, 5936, 6000, 6064, 6128, 6192, 6256, 6320, 6384, 6448, 6512, 6576, 6640, 6704, 6768, 6832, 6896, 6960, 7024, 7088, 7152, 7216, 7280, 7344, 7408, 7472, 7536, 7600, 7664, 7728, 7792, 7856, 7920, 7984, 8048, 8112, 8176, 8240, 8304, 8368, 8432, 8496, 8560, 8624, 8688, 8752, 8816, 8880, 8944, 9008, 9072, 9136, 9200, 9264, 9328, 9392, 9456, 9520, 9584, 9648, 9712, 9776, 9840, 9904, 9968, 10032, 10096, 10160, 10224, 10288, 10352, 10416, 10480, 10544, 10608, 10672, 10736, 10800, 10864, 10928, 10992};
    int length = sizeof(my_list) / sizeof(my_list[0]);
    printf("%d", length);
    char algoname[100] = "naive_matrix_cpu";
    for (int i = 0; i < length; i++){
        if(my_list[i]<=1360){
        double l[5];
        int matSize = my_list[i];
        cout<<"Matrix Size : "<<matSize<<endl;
        int A_rows = matSize;
        int A_cols = matSize;
        int B_rows = matSize;
        int B_cols = matSize;
        int C_rows = A_rows;
        int C_cols = B_cols;
        int A_size = A_rows * A_cols;
        int B_size = B_rows * B_cols;
        int C_size = C_rows * C_cols;
        float* A, * B, * C, * C_cpu;
        // // timing
        // cudaEvent_t start_gpu, stop_gpu;
        // float gpu_time_ms = 0;
        // cudaEventCreate(&start_gpu);
        // cudaEventCreate(&stop_gpu);

        // Allocate Unified Memory – accessible from CPU or GPU
        cudaMallocManaged(&A, A_size * sizeof(float));
        cudaMallocManaged(&B, B_size * sizeof(float));
        cudaMallocManaged(&C, C_size * sizeof(float));
        cudaMallocManaged(&C_cpu, C_size * sizeof(float));

        // initialize A and B matrices
        auto all_ones = []() -> float {
            return 1.0f;
        };

        srand(time(NULL));
        auto rand_numbers = []() -> float {
            return static_cast<float>(rand()) / (static_cast<float>(RAND_MAX / 1000));
        };
        
        auto index_based = [](int i, int j) -> float {
            return j;
        };

        initialize_matrix<float>(A, A_rows, A_cols, rand_numbers);
        initialize_matrix<float>(B, B_rows, B_cols, rand_numbers);


        // launch kernel

        dim3 dim_grid(C_cols / COL_TILE_WIDTH, C_rows / ROW_TILE_WIDTH, 1);
        dim3 dim_block(COL_TILE_WIDTH, ROW_TILE_WIDTH, 1);

        //cudaEventRecord(start_gpu);
        //naive_matrix_multiply<float> << <dim_grid, dim_block >> > (A, B, C, A_cols, C_rows, C_cols);
        //cudaEventRecord(stop_gpu);

        //// Wait for GPU to finish before accessing on host
        //cudaDeviceSynchronize();

        //cudaEventSynchronize(stop_gpu);
        //cudaEventElapsedTime(&gpu_time_ms, start_gpu, stop_gpu);

        if(my_list[i]<=768){
        
        for(int j=0;j<5;j++){
        // check results on CPU
        auto t1 = std::chrono::system_clock::now();
        naive_matrix_multiply_cpu<float>(A, B, C_cpu, A_cols, C_rows, C_cols);
        auto t2 = std::chrono::system_clock::now();

        if (fabsf(maxDiff<float>(C, C_cpu, C_rows, C_cols)) <= (float)EPSILON)
            std::cout << "PASS" << std::endl;
        else
            std::cout << "FAIL" << std::endl;

        auto cpu_time_ms = std::chrono::duration_cast<std::chrono::microseconds>(t2 - t1).count() / 1000.0f;
        // std::cout << "GPU time = " << gpu_time_ms << "ms" << std::endl;
        l[j]=cpu_time_ms;    
    }
    double avg;
	avg=(l[0]+l[1]+l[2]+l[3]+l[4])/5;
		cout << "Using Milliseconds Clock: (AVG)"<< endl;
		cout << " CPU time taken to execute for strassen matrices of size - " 
		<< matSize << " : " <<avg<<" ms"<< endl;
		cout << endl;
        fprintf(fp,"\n%s,%d,%lf",algoname,matSize,avg);
    }else{
        auto t1 = std::chrono::system_clock::now();
        naive_matrix_multiply_cpu<float>(A, B, C_cpu, A_cols, C_rows, C_cols);
        auto t2 = std::chrono::system_clock::now();
        if (fabsf(maxDiff<float>(C, C_cpu, C_rows, C_cols)) <= (float)EPSILON)
            std::cout << "PASS" << std::endl;
        else
            std::cout << "FAIL" << std::endl;
        auto cpu_time_ms = std::chrono::duration_cast<std::chrono::microseconds>(t2 - t1).count() / 1000.0f;
        cout << "Using Milliseconds Clock:(else above 1000) "<< endl;
		cout << " CPU time taken to execute for strassen matrices of size - " 
		<< matSize << " : " <<cpu_time_ms<<" ms"<< endl;
		cout << endl;
        fprintf(fp,"\n%s,%d,%lf",algoname,matSize,cpu_time_ms);
    }
       
        // std::cout << "Speedup = " << cpu_time_ms / gpu_time_ms << std::endl;

  
        // fprintf(fp, "\n%s,%d,%lf", algoname, matSize, cpu_time_ms);

        
        // Free memory
        cudaFree(A);
        cudaFree(B);
        cudaFree(C);
    }
    else
      {
           cout<<"MAX LIMIT"<<endl;
           fprintf(fp, "\n%s,%d,%d", algoname, my_list[i], max_limit);
     }
 }
    return 0;
}
	.file	"tmpxft_00381170_00000000-6_naive_cpu.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB5175:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE5175:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation, @function
_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation:
.LFB5714:
	.cfi_startproc
	endbr64
	testl	%edx, %edx
	je	.L3
	decl	%edx
	jne	.L5
	movq	%rsi, (%rdi)
	jmp	.L5
.L3:
	leaq	_ZTIZ4mainEUlvE0_(%rip), %rax
	movq	%rax, (%rdi)
.L5:
	xorl	%eax, %eax
	ret
	.cfi_endproc
.LFE5714:
	.size	_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation, .-_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB5924:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE5924:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.type	_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data, @function
_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data:
.LFB5710:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	rand@PLT
	cvtsi2ssl	%eax, %xmm0
	divss	.LC1(%rip), %xmm0
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5710:
	.size	_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data, .-_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data
	.section	.text._ZNSt14_Function_baseD2Ev,"axG",@progbits,_ZNSt14_Function_baseD5Ev,comdat
	.align 2
	.weak	_ZNSt14_Function_baseD2Ev
	.type	_ZNSt14_Function_baseD2Ev, @function
_ZNSt14_Function_baseD2Ev:
.LFB3651:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3651
	endbr64
	movq	16(%rdi), %rax
	testq	%rax, %rax
	je	.L16
	movq	%rdi, %rsi
	pushq	%rdx
	.cfi_def_cfa_offset 16
	movl	$3, %edx
	call	*%rax
	popq	%rcx
	.cfi_def_cfa_offset 8
	ret
.L16:
	ret
	.cfi_endproc
.LFE3651:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table._ZNSt14_Function_baseD2Ev,"aG",@progbits,_ZNSt14_Function_baseD5Ev,comdat
.LLSDA3651:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3651-.LLSDACSB3651
.LLSDACSB3651:
.LLSDACSE3651:
	.section	.text._ZNSt14_Function_baseD2Ev,"axG",@progbits,_ZNSt14_Function_baseD5Ev,comdat
	.size	_ZNSt14_Function_baseD2Ev, .-_ZNSt14_Function_baseD2Ev
	.weak	_ZNSt14_Function_baseD1Ev
	.set	_ZNSt14_Function_baseD1Ev,_ZNSt14_Function_baseD2Ev
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB5198:
	.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
.LFE5198:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text._Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE,"axG",@progbits,_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE,comdat
	.weak	_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE
	.type	_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE, @function
_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE:
.LFB5542:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%r15d, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	xorl	%r14d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rdi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rcx, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
.L22:
	cmpl	%esi, %r15d
	jge	.L21
	movslq	%r14d, %rax
	xorl	%ebp, %ebp
	leaq	0(%r13,%rax,4), %rax
.L26:
	cmpl	%ebp, %r12d
	jle	.L28
	cmpq	$0, 16(%rbx)
	jne	.L23
	call	_ZSt25__throw_bad_function_callv@PLT
.L23:
	movl	%esi, 12(%rsp)
	movq	%rbx, %rdi
	movq	%rax, (%rsp)
	call	*24(%rbx)
	movq	(%rsp), %rax
	movl	12(%rsp), %esi
	movss	%xmm0, (%rax,%rbp,4)
	incq	%rbp
	jmp	.L26
.L28:
	incl	%r15d
	addl	%r12d, %r14d
	jmp	.L22
.L21:
	addq	$24, %rsp
	.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
.LFE5542:
	.size	_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE, .-_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE
	.section	.text._Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii,"axG",@progbits,_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii,comdat
	.weak	_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii
	.type	_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii, @function
_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii:
.LFB5543:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdx, %r10
	xorl	%r11d, %r11d
	movl	%ecx, %edx
	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
	movq	%rsi, %r12
	xorl	%esi, %esi
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movslq	%r9d, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	salq	$2, %rbp
	xorl	%ebx, %ebx
.L30:
	cmpl	%r8d, %ebx
	jge	.L29
	movslq	%r11d, %rax
	leaq	(%r10,%rax,4), %r15
	movslq	%esi, %rax
	leaq	(%rdi,%rax,4), %r14
	xorl	%eax, %eax
.L35:
	cmpl	%eax, %r9d
	jle	.L32
	leaq	(%r12,%rax,4), %r13
	xorl	%ecx, %ecx
	xorps	%xmm0, %xmm0
.L33:
	cmpl	%ecx, %edx
	jle	.L37
	movss	(%r14,%rcx,4), %xmm1
	mulss	0(%r13), %xmm1
	incq	%rcx
	addq	%rbp, %r13
	addss	%xmm1, %xmm0
	jmp	.L33
.L37:
	movss	%xmm0, (%r15,%rax,4)
	incq	%rax
	jmp	.L35
.L32:
	incl	%ebx
	addl	%r9d, %r11d
	addl	%edx, %esi
	jmp	.L30
.L29:
	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
.LFE5543:
	.size	_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii, .-_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii
	.section	.text._Z7maxDiffIfET_PS0_S1_ii,"axG",@progbits,_Z7maxDiffIfET_PS0_S1_ii,comdat
	.weak	_Z7maxDiffIfET_PS0_S1_ii
	.type	_Z7maxDiffIfET_PS0_S1_ii, @function
_Z7maxDiffIfET_PS0_S1_ii:
.LFB5544:
	.cfi_startproc
	endbr64
	movss	(%rdi), %xmm0
	subss	(%rsi), %xmm0
	xorl	%r9d, %r9d
	xorl	%r10d, %r10d
	movss	.LC3(%rip), %xmm2
.L39:
	cmpl	%edx, %r10d
	jge	.L38
	movslq	%r9d, %rax
	xorl	%r8d, %r8d
	salq	$2, %rax
.L43:
	cmpl	%ecx, %r8d
	jge	.L44
	movss	(%rdi,%rax), %xmm1
	subss	(%rsi,%rax), %xmm1
	incl	%r8d
	addq	$4, %rax
	andps	%xmm2, %xmm1
	maxss	%xmm0, %xmm1
	movaps	%xmm1, %xmm0
	jmp	.L43
.L44:
	incl	%r10d
	addl	%ecx, %r9d
	jmp	.L39
.L38:
	ret
	.cfi_endproc
.LFE5544:
	.size	_Z7maxDiffIfET_PS0_S1_ii, .-_Z7maxDiffIfET_PS0_S1_ii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC4:
	.string	"w+"
.LC5:
	.string	"Naive_Multiplication_CPU_Mohit.csv"
.LC6:
	.string	"Algorithm_Name,Input_Dimensions,Execution_Time(ms)"
.LC7:
	.string	"%d"
.LC8:
	.string	"Matrix Size : "
.LC10:
	.string	"PASS"
.LC11:
	.string	"FAIL"
.LC14:
	.string	"Using Milliseconds Clock: (AVG)"
.LC15:
	.string	" CPU time taken to execute for strassen matrices of size - "
.LC16:
	.string	" : "
.LC17:
	.string	" ms"
.LC18:
	.string	"\n%s,%d,%lf"
.LC19:
	.string	"Using Milliseconds Clock:(else above 1000) "
.LC20:
	.string	"MAX LIMIT"
.LC21:
	.string	"\n%s,%d,%d"
	.section	.rodata
	.align 32
.LC0:
	.long	16
	.long	32
	.long	48
	.long	64
	.long	80
	.long	96
	.long	112
	.long	128
	.long	144
	.long	160
	.long	176
	.long	192
	.long	208
	.long	224
	.long	240
	.long	256
	.long	272
	.long	288
	.long	304
	.long	320
	.long	336
	.long	352
	.long	368
	.long	384
	.long	400
	.long	416
	.long	432
	.long	448
	.long	464
	.long	480
	.long	496
	.long	512
	.long	528
	.long	544
	.long	560
	.long	576
	.long	592
	.long	608
	.long	624
	.long	640
	.long	656
	.long	672
	.long	688
	.long	704
	.long	720
	.long	736
	.long	752
	.long	768
	.long	784
	.long	800
	.long	816
	.long	832
	.long	848
	.long	864
	.long	880
	.long	896
	.long	912
	.long	928
	.long	944
	.long	960
	.long	976
	.long	992
	.long	1008
	.long	1040
	.long	1072
	.long	1104
	.long	1136
	.long	1168
	.long	1200
	.long	1232
	.long	1264
	.long	1296
	.long	1328
	.long	1360
	.long	1392
	.long	1424
	.long	1456
	.long	1488
	.long	1520
	.long	1552
	.long	1584
	.long	1616
	.long	1648
	.long	1680
	.long	1712
	.long	1744
	.long	1776
	.long	1808
	.long	1840
	.long	1872
	.long	1904
	.long	1936
	.long	1968
	.long	2000
	.long	2032
	.long	2064
	.long	2096
	.long	2128
	.long	2160
	.long	2192
	.long	2224
	.long	2256
	.long	2288
	.long	2320
	.long	2352
	.long	2384
	.long	2416
	.long	2448
	.long	2480
	.long	2512
	.long	2544
	.long	2576
	.long	2608
	.long	2640
	.long	2672
	.long	2704
	.long	2736
	.long	2768
	.long	2800
	.long	2832
	.long	2864
	.long	2896
	.long	2928
	.long	2960
	.long	2992
	.long	3024
	.long	3056
	.long	3088
	.long	3120
	.long	3152
	.long	3184
	.long	3216
	.long	3248
	.long	3280
	.long	3312
	.long	3344
	.long	3376
	.long	3408
	.long	3440
	.long	3472
	.long	3504
	.long	3536
	.long	3568
	.long	3600
	.long	3632
	.long	3664
	.long	3696
	.long	3728
	.long	3760
	.long	3792
	.long	3824
	.long	3856
	.long	3888
	.long	3920
	.long	3952
	.long	3984
	.long	4016
	.long	4080
	.long	4144
	.long	4208
	.long	4272
	.long	4336
	.long	4400
	.long	4464
	.long	4528
	.long	4592
	.long	4656
	.long	4720
	.long	4784
	.long	4848
	.long	4912
	.long	4976
	.long	5040
	.long	5104
	.long	5168
	.long	5232
	.long	5296
	.long	5360
	.long	5424
	.long	5488
	.long	5552
	.long	5616
	.long	5680
	.long	5744
	.long	5808
	.long	5872
	.long	5936
	.long	6000
	.long	6064
	.long	6128
	.long	6192
	.long	6256
	.long	6320
	.long	6384
	.long	6448
	.long	6512
	.long	6576
	.long	6640
	.long	6704
	.long	6768
	.long	6832
	.long	6896
	.long	6960
	.long	7024
	.long	7088
	.long	7152
	.long	7216
	.long	7280
	.long	7344
	.long	7408
	.long	7472
	.long	7536
	.long	7600
	.long	7664
	.long	7728
	.long	7792
	.long	7856
	.long	7920
	.long	7984
	.long	8048
	.long	8112
	.long	8176
	.long	8240
	.long	8304
	.long	8368
	.long	8432
	.long	8496
	.long	8560
	.long	8624
	.long	8688
	.long	8752
	.long	8816
	.long	8880
	.long	8944
	.long	9008
	.long	9072
	.long	9136
	.long	9200
	.long	9264
	.long	9328
	.long	9392
	.long	9456
	.long	9520
	.long	9584
	.long	9648
	.long	9712
	.long	9776
	.long	9840
	.long	9904
	.long	9968
	.long	10032
	.long	10096
	.long	10160
	.long	10224
	.long	10288
	.long	10352
	.long	10416
	.long	10480
	.long	10544
	.long	10608
	.long	10672
	.long	10736
	.long	10800
	.long	10864
	.long	10928
	.long	10992
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB5156:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA5156
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC4(%rip), %rsi
	leaq	.LC5(%rip), %rdi
	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
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$1272, %rsp
	.cfi_def_cfa_offset 1328
	movq	%fs:40, %rax
	movq	%rax, 1256(%rsp)
	xorl	%eax, %eax
	leaq	92(%rsp), %r15
.LEHB0:
	call	fopen@PLT
	leaq	.LC6(%rip), %rdx
	movl	$2, %esi
	movq	%rax, %rdi
	movq	%rax, (%rsp)
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movl	$266, %edx
	leaq	92(%rsp), %rdi
	xorl	%eax, %eax
	leaq	.LC0(%rip), %rsi
	movl	$266, %ecx
	rep movsl
	movl	$2, %edi
	leaq	.LC7(%rip), %rsi
	call	__printf_chk@PLT
	leaq	1172(%rsp), %rdi
	movl	$21, %ecx
	movabsq	$7020372283523424622, %rax
	movq	%rax, 1156(%rsp)
	xorl	%eax, %eax
	movabsq	$8462372961522381428, %rdx
	rep stosl
	movq	%rdx, 1164(%rsp)
.L57:
	movl	(%r15,%r12,4), %ebx
	cmpl	$1360, %ebx
	jg	.L46
	movl	%ebx, %ebp
	leaq	.LC8(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	imull	%ebx, %ebp
	movl	%ebx, %esi
	leaq	_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data(%rip), %r14
	movq	%rax, %rdi
	leaq	_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation(%rip), %r13
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movslq	%ebp, %rbp
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	salq	$2, %rbp
	leaq	16(%rsp), %rdi
	movl	$1, %edx
	movq	%rbp, %rsi
	call	cudaMallocManaged@PLT
	movq	%rbp, %rsi
	leaq	24(%rsp), %rdi
	movl	$1, %edx
	call	cudaMallocManaged@PLT
	movq	%rbp, %rsi
	leaq	32(%rsp), %rdi
	movl	$1, %edx
	call	cudaMallocManaged@PLT
	movq	%rbp, %rsi
	movl	$1, %edx
	leaq	40(%rsp), %rdi
	call	cudaMallocManaged@PLT
.LEHE0:
	xorl	%edi, %edi
	leaq	48(%rsp), %rbp
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	movq	16(%rsp), %rdi
	xorps	%xmm0, %xmm0
	movq	%rbp, %rcx
	movl	%ebx, %edx
	movl	%ebx, %esi
	movq	%r14, 72(%rsp)
	movq	%r13, 64(%rsp)
	movups	%xmm0, 48(%rsp)
.LEHB1:
	call	_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE
.LEHE1:
	movq	%rbp, %rdi
	call	_ZNSt14_Function_baseD2Ev
	movq	24(%rsp), %rdi
	xorps	%xmm0, %xmm0
	movq	%rbp, %rcx
	movl	%ebx, %edx
	movl	%ebx, %esi
	movq	%r14, 72(%rsp)
	movq	%r13, 64(%rsp)
	movups	%xmm0, 48(%rsp)
.LEHB2:
	call	_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE
.LEHE2:
	movq	%rbp, %rdi
	call	_ZNSt14_Function_baseD2Ev
	cmpl	$768, %ebx
	jg	.L47
	xorl	%r13d, %r13d
.L51:
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movq	40(%rsp), %rdx
	movl	%ebx, %r9d
	movl	%ebx, %r8d
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	%ebx, %ecx
	movq	%rax, %r14
	call	_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movq	40(%rsp), %rsi
	movl	%ebx, %ecx
	movl	%ebx, %edx
	movq	32(%rsp), %rdi
	movq	%rax, 8(%rsp)
	call	_Z7maxDiffIfET_PS0_S1_ii
	movss	.LC9(%rip), %xmm2
	andps	.LC3(%rip), %xmm0
	leaq	.LC10(%rip), %rsi
	comiss	%xmm0, %xmm2
	jnb	.L73
	leaq	.LC11(%rip), %rsi
.L73:
	leaq	_ZSt4cout(%rip), %rdi
.LEHB3:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	8(%rsp), %rax
	movl	$1000, %ecx
	subq	%r14, %rax
	cqto
	idivq	%rcx
	cvtsi2ssq	%rax, %xmm0
	divss	.LC12(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movsd	%xmm0, 0(%rbp,%r13,8)
	incq	%r13
	cmpq	$5, %r13
	jne	.L51
	leaq	.LC14(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movsd	48(%rsp), %xmm0
	addsd	56(%rsp), %xmm0
	addsd	64(%rsp), %xmm0
	addsd	72(%rsp), %xmm0
	addsd	80(%rsp), %xmm0
	divsd	.LC13(%rip), %xmm0
	movq	%xmm0, %rbp
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC15(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC16(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rbp, %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC17(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	1156(%rsp), %rcx
	movq	%rbp, %xmm0
	jmp	.L75
.L47:
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movq	40(%rsp), %rdx
	movl	%ebx, %r9d
	movl	%ebx, %r8d
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdi
	movl	%ebx, %ecx
	movq	%rax, %r13
	call	_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movq	40(%rsp), %rsi
	movl	%ebx, %ecx
	movl	%ebx, %edx
	movq	32(%rsp), %rdi
	movq	%rax, %rbp
	call	_Z7maxDiffIfET_PS0_S1_ii
	movss	.LC9(%rip), %xmm1
	andps	.LC3(%rip), %xmm0
	leaq	.LC10(%rip), %rsi
	comiss	%xmm0, %xmm1
	jnb	.L74
	leaq	.LC11(%rip), %rsi
.L74:
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	%rbp, %rax
	movl	$1000, %ecx
	leaq	.LC19(%rip), %rsi
	subq	%r13, %rax
	leaq	_ZSt4cout(%rip), %rdi
	cqto
	idivq	%rcx
	cvtsi2ssq	%rax, %xmm0
	divss	.LC12(%rip), %xmm0
	movss	%xmm0, 8(%rsp)
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC15(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC16(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movss	8(%rsp), %xmm0
	movq	%rax, %rdi
	cvtss2sd	%xmm0, %xmm0
	movsd	%xmm0, 8(%rsp)
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	leaq	.LC17(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movsd	8(%rsp), %xmm0
	leaq	1156(%rsp), %rcx
.L75:
	movq	(%rsp), %rdi
	movl	%ebx, %r8d
	leaq	.LC18(%rip), %rdx
	movb	$1, %al
	movl	$2, %esi
	call	__fprintf_chk@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	jmp	.L56
.L46:
	leaq	.LC20(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	(%rsp), %rdi
	movl	%ebx, %r8d
	xorl	%eax, %eax
	leaq	1156(%rsp), %rcx
	movl	$800000, %r9d
	movl	$2, %esi
	leaq	.LC21(%rip), %rdx
	call	__fprintf_chk@PLT
.L56:
	incq	%r12
	cmpq	$266, %r12
	jne	.L57
	movq	1256(%rsp), %rax
	subq	%fs:40, %rax
	je	.L62
	jmp	.L71
.L63:
	endbr64
	jmp	.L77
.L64:
	endbr64
.L77:
	movq	%rax, %rbx
.L60:
	movq	%rbp, %rdi
	call	_ZNSt14_Function_baseD2Ev
	movq	1256(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L71
	movq	%rbx, %rdi
	call	_Unwind_Resume@PLT
.LEHE3:
.L71:
	call	__stack_chk_fail@PLT
.L62:
	addq	$1272, %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
.LFE5156:
	.section	.gcc_except_table,"a",@progbits
.LLSDA5156:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE5156-.LLSDACSB5156
.LLSDACSB5156:
	.uleb128 .LEHB0-.LFB5156
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB5156
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L63-.LFB5156
	.uleb128 0
	.uleb128 .LEHB2-.LFB5156
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L64-.LFB5156
	.uleb128 0
	.uleb128 .LEHB3-.LFB5156
	.uleb128 .LEHE3-.LEHB3
	.uleb128 0
	.uleb128 0
.LLSDACSE5156:
	.section	.text.startup
	.size	main, .-main
	.section	.data.rel.ro,"aw"
	.align 8
	.type	_ZTIZ4mainEUlvE0_, @object
	.size	_ZTIZ4mainEUlvE0_, 16
_ZTIZ4mainEUlvE0_:
	.quad	_ZTVN10__cxxabiv117__class_type_infoE+16
	.quad	_ZTSZ4mainEUlvE0_
	.section	.rodata
	.align 8
	.type	_ZTSZ4mainEUlvE0_, @object
	.size	_ZTSZ4mainEUlvE0_, 15
_ZTSZ4mainEUlvE0_:
	.string	"*Z4mainEUlvE0_"
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC1:
	.long	1241715308
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC3:
	.long	2147483647
	.long	0
	.long	0
	.long	0
	.section	.rodata.cst4
	.align 4
.LC9:
	.long	897988541
	.align 4
.LC12:
	.long	1148846080
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC13:
	.long	0
	.long	1075052544
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.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 <iostream>
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
#include <math.h>
#include <functional>
#include <stdlib.h>     /* srand, rand */
#include <time.h>       /* time */
#include <chrono>
#include <list>
#include <hip/hip_runtime.h>
#include <cstdlib>
#include <ctime>
#include <cmath>
using namespace std;
#define ROW_TILE_WIDTH  64
#define COL_TILE_WIDTH  64

#define EPSILON         (1e-6)

template<typename T>
__global__
void naive_matrix_multiply(T* A, T* B, T* C, int width, int cRows, int cCols)
{
    __shared__ T shATile[ROW_TILE_WIDTH][COL_TILE_WIDTH];
    __shared__ T shBTile[ROW_TILE_WIDTH][COL_TILE_WIDTH];

    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    T pValue = 0;

    // iterate for width/COL_TILE_WIDTH number of times
    // to compute the C tile
    for (int p = 0; p < width / COL_TILE_WIDTH; p++) {
        //load values to tiles from A and B
        shATile[threadIdx.y][threadIdx.x] = A[row * width + p * ROW_TILE_WIDTH + threadIdx.x];
        shBTile[threadIdx.y][threadIdx.x] = B[(p * COL_TILE_WIDTH + threadIdx.y) * cCols + col];

        // wait until all threads finish loading values
        __syncthreads();
        // update pValue for this thread
        for (int i = 0; i < COL_TILE_WIDTH; i++) pValue += shATile[threadIdx.y][i] * shBTile[i][threadIdx.x];
        // wait until all threads finish computing pValue before overwriting
        __syncthreads();

    }

    C[row * cCols + col] = pValue;

}

template<typename T>
void initialize_matrix(T* M, int rows, int cols, std::function<float()> F) {
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            M[i * cols + j] = F();
        }
    }
}

template<typename T>
void initialize_matrix(T* M, int rows, int cols, std::function<float(int, int)> F) {
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            M[i * cols + j] = F(i, j);
        }
    }
}

template<typename T>
void print_matrix(T* M, int rows, int cols) {
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            std::cout << M[i * cols + j] << " ";
        }
        std::cout << "\n";
    }
}



template<typename T>
void naive_matrix_multiply_cpu(T* A, T* B, T* C, int width, int C_rows, int C_cols) {

    for (int i = 0; i < C_rows; i++)
        for (int j = 0; j < C_cols; j++) {
            T value = 0.0f;
            for (int k = 0; k < width; k++) {
                value += A[i * width + k] * B[k * C_cols + j];
            }


            C[i * C_cols + j] = value;
        }
}

template<typename T>
T maxDiff(T* A1, T* A2, int rows, int cols) {
    T maxDiff = A1[0] - A2[0];
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            T diff = abs(A1[i * cols + j] - A2[i * cols + j]);
            if (diff > maxDiff) {
                maxDiff = diff;
            }
        }
    }


    return maxDiff;
}


int main(void)
{
    FILE* fp;
    fp = fopen("Naive_Multiplication_CPU_Mohit.csv", "w+");
    fprintf(fp, "Algorithm_Name,Input_Dimensions,Execution_Time(ms)");
    int max_limit = 800000;
    int my_list[] = {16, 32, 48, 64, 80, 96, 112, 128, 144, 160, 176, 192, 208, 224, 240, 256, 272, 288, 304, 320, 336, 352, 368, 384, 400, 416, 432, 448, 464, 480, 496, 512, 528, 544, 560, 576, 592, 608, 624, 640, 656, 672, 688, 704, 720, 736, 752, 768, 784, 800, 816, 832, 848, 864, 880, 896, 912, 928, 944, 960, 976, 992, 1008, 1040, 1072, 1104, 1136, 1168, 1200, 1232, 1264, 1296, 1328, 1360, 1392, 1424, 1456, 1488, 1520, 1552, 1584, 1616, 1648, 1680, 1712, 1744, 1776, 1808, 1840, 1872, 1904, 1936, 1968, 2000, 2032, 2064, 2096, 2128, 2160, 2192, 2224, 2256, 2288, 2320, 2352, 2384, 2416, 2448, 2480, 2512, 2544, 2576, 2608, 2640, 2672, 2704, 2736, 2768, 2800, 2832, 2864, 2896, 2928, 2960, 2992, 3024, 3056, 3088, 3120, 3152, 3184, 3216, 3248, 3280, 3312, 3344, 3376, 3408, 3440, 3472, 3504, 3536, 3568, 3600, 3632, 3664, 3696, 3728, 3760, 3792, 3824, 3856, 3888, 3920, 3952, 3984, 4016, 4080, 4144, 4208, 4272, 4336, 4400, 4464, 4528, 4592, 4656, 4720, 4784, 4848, 4912,4976, 5040, 5104, 5168, 5232, 5296, 5360, 5424, 5488, 5552, 5616, 5680, 5744, 5808, 5872, 5936, 6000, 6064, 6128, 6192, 6256, 6320, 6384, 6448, 6512, 6576, 6640, 6704, 6768, 6832, 6896, 6960, 7024, 7088, 7152, 7216, 7280, 7344, 7408, 7472, 7536, 7600, 7664, 7728, 7792, 7856, 7920, 7984, 8048, 8112, 8176, 8240, 8304, 8368, 8432, 8496, 8560, 8624, 8688, 8752, 8816, 8880, 8944, 9008, 9072, 9136, 9200, 9264, 9328, 9392, 9456, 9520, 9584, 9648, 9712, 9776, 9840, 9904, 9968, 10032, 10096, 10160, 10224, 10288, 10352, 10416, 10480, 10544, 10608, 10672, 10736, 10800, 10864, 10928, 10992};
    int length = sizeof(my_list) / sizeof(my_list[0]);
    printf("%d", length);
    char algoname[100] = "naive_matrix_cpu";
    for (int i = 0; i < length; i++){
        if(my_list[i]<=1360){
        double l[5];
        int matSize = my_list[i];
        cout<<"Matrix Size : "<<matSize<<endl;
        int A_rows = matSize;
        int A_cols = matSize;
        int B_rows = matSize;
        int B_cols = matSize;
        int C_rows = A_rows;
        int C_cols = B_cols;
        int A_size = A_rows * A_cols;
        int B_size = B_rows * B_cols;
        int C_size = C_rows * C_cols;
        float* A, * B, * C, * C_cpu;
        // // timing
        // cudaEvent_t start_gpu, stop_gpu;
        // float gpu_time_ms = 0;
        // cudaEventCreate(&start_gpu);
        // cudaEventCreate(&stop_gpu);

        // Allocate Unified Memory – accessible from CPU or GPU
        hipMallocManaged(&A, A_size * sizeof(float));
        hipMallocManaged(&B, B_size * sizeof(float));
        hipMallocManaged(&C, C_size * sizeof(float));
        hipMallocManaged(&C_cpu, C_size * sizeof(float));

        // initialize A and B matrices
        auto all_ones = []() -> float {
            return 1.0f;
        };

        srand(time(NULL));
        auto rand_numbers = []() -> float {
            return static_cast<float>(rand()) / (static_cast<float>(RAND_MAX / 1000));
        };
        
        auto index_based = [](int i, int j) -> float {
            return j;
        };

        initialize_matrix<float>(A, A_rows, A_cols, rand_numbers);
        initialize_matrix<float>(B, B_rows, B_cols, rand_numbers);


        // launch kernel

        dim3 dim_grid(C_cols / COL_TILE_WIDTH, C_rows / ROW_TILE_WIDTH, 1);
        dim3 dim_block(COL_TILE_WIDTH, ROW_TILE_WIDTH, 1);

        //cudaEventRecord(start_gpu);
        //naive_matrix_multiply<float> << <dim_grid, dim_block >> > (A, B, C, A_cols, C_rows, C_cols);
        //cudaEventRecord(stop_gpu);

        //// Wait for GPU to finish before accessing on host
        //cudaDeviceSynchronize();

        //cudaEventSynchronize(stop_gpu);
        //cudaEventElapsedTime(&gpu_time_ms, start_gpu, stop_gpu);

        if(my_list[i]<=768){
        
        for(int j=0;j<5;j++){
        // check results on CPU
        auto t1 = std::chrono::system_clock::now();
        naive_matrix_multiply_cpu<float>(A, B, C_cpu, A_cols, C_rows, C_cols);
        auto t2 = std::chrono::system_clock::now();

        if (fabsf(maxDiff<float>(C, C_cpu, C_rows, C_cols)) <= (float)EPSILON)
            std::cout << "PASS" << std::endl;
        else
            std::cout << "FAIL" << std::endl;

        auto cpu_time_ms = std::chrono::duration_cast<std::chrono::microseconds>(t2 - t1).count() / 1000.0f;
        // std::cout << "GPU time = " << gpu_time_ms << "ms" << std::endl;
        l[j]=cpu_time_ms;    
    }
    double avg;
	avg=(l[0]+l[1]+l[2]+l[3]+l[4])/5;
		cout << "Using Milliseconds Clock: (AVG)"<< endl;
		cout << " CPU time taken to execute for strassen matrices of size - " 
		<< matSize << " : " <<avg<<" ms"<< endl;
		cout << endl;
        fprintf(fp,"\n%s,%d,%lf",algoname,matSize,avg);
    }else{
        auto t1 = std::chrono::system_clock::now();
        naive_matrix_multiply_cpu<float>(A, B, C_cpu, A_cols, C_rows, C_cols);
        auto t2 = std::chrono::system_clock::now();
        if (fabsf(maxDiff<float>(C, C_cpu, C_rows, C_cols)) <= (float)EPSILON)
            std::cout << "PASS" << std::endl;
        else
            std::cout << "FAIL" << std::endl;
        auto cpu_time_ms = std::chrono::duration_cast<std::chrono::microseconds>(t2 - t1).count() / 1000.0f;
        cout << "Using Milliseconds Clock:(else above 1000) "<< endl;
		cout << " CPU time taken to execute for strassen matrices of size - " 
		<< matSize << " : " <<cpu_time_ms<<" ms"<< endl;
		cout << endl;
        fprintf(fp,"\n%s,%d,%lf",algoname,matSize,cpu_time_ms);
    }
       
        // std::cout << "Speedup = " << cpu_time_ms / gpu_time_ms << std::endl;

  
        // fprintf(fp, "\n%s,%d,%lf", algoname, matSize, cpu_time_ms);

        
        // Free memory
        hipFree(A);
        hipFree(B);
        hipFree(C);
    }
    else
      {
           cout<<"MAX LIMIT"<<endl;
           fprintf(fp, "\n%s,%d,%d", algoname, my_list[i], max_limit);
     }
 }
    return 0;
}
	.text
	.file	"naive_cpu.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI0_0:
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
.LCPI0_4:
	.zero	16
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0
.LCPI0_1:
	.long	0x358637bd                      # float 9.99999997E-7
.LCPI0_2:
	.long	0x447a0000                      # float 1000
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI0_3:
	.quad	0x4014000000000000              # double 5
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
.Lfunc_begin0:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception0
# %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	$280, %rsp                      # imm = 0x118
	.cfi_def_cfa_offset 336
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, %r12
	movl	$.L.str.2, %edi
	movl	$50, %esi
	movl	$1, %edx
	movq	%rax, %rcx
	callq	fwrite@PLT
	movl	$.L.str.3, %edi
	movl	$266, %esi                      # imm = 0x10A
	xorl	%eax, %eax
	callq	printf
	leaq	176(%rsp), %r14
	movl	$12, %ecx
	movl	$.L__const.main.algoname, %esi
	movq	%r14, %rdi
	rep;movsq (%rsi), %es:(%rdi)
	movl	$0, 96(%r14)
	movl	$_ZSt4cout, %r15d
	xorl	%ebx, %ebx
	movq	%r12, 120(%rsp)                 # 8-byte Spill
.LBB0_1:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_10 Depth 2
	cmpq	$73, %rbx
	ja	.LBB0_14
# %bb.2:                                #   in Loop: Header=BB0_1 Depth=1
	movl	.L__const.main.my_list(,%rbx,4), %r13d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$14, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%r13d, %esi
	callq	_ZNSolsEi
	movq	%rax, %r14
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	%r13d, %r14d
	imull	%r14d, %r14d
	shlq	$2, %r14
	leaq	32(%rsp), %rdi
	movq	%r14, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	24(%rsp), %rdi
	movq	%r14, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	40(%rsp), %rdi
	movq	%r14, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	16(%rsp), %rdi
	movq	%r14, %rsi
	movl	$1, %edx
	callq	hipMallocManaged
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	movq	32(%rsp), %rdi
	xorps	%xmm0, %xmm0
	movaps	%xmm0, 80(%rsp)
	movq	$_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data, 104(%rsp)
	movq	$_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation, 96(%rsp)
.Ltmp0:
	movl	%r13d, %esi
	movl	%r13d, %edx
	leaq	80(%rsp), %rcx
	callq	_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE
.Ltmp1:
# %bb.3:                                #   in Loop: Header=BB0_1 Depth=1
	movq	96(%rsp), %rax
	testq	%rax, %rax
	je	.LBB0_5
# %bb.4:                                #   in Loop: Header=BB0_1 Depth=1
.Ltmp6:
	leaq	80(%rsp), %rdi
	movq	%rdi, %rsi
	movl	$3, %edx
	callq	*%rax
.Ltmp7:
.LBB0_5:                                # %_ZNSt14_Function_baseD2Ev.exit
                                        #   in Loop: Header=BB0_1 Depth=1
	movq	24(%rsp), %rdi
	xorps	%xmm0, %xmm0
	movaps	%xmm0, 48(%rsp)
	movq	$_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data, 72(%rsp)
	movq	$_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation, 64(%rsp)
.Ltmp9:
	movl	%r13d, %esi
	movl	%r13d, %edx
	leaq	48(%rsp), %rcx
	callq	_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE
.Ltmp10:
# %bb.6:                                #   in Loop: Header=BB0_1 Depth=1
	movq	64(%rsp), %rax
	testq	%rax, %rax
	je	.LBB0_8
# %bb.7:                                #   in Loop: Header=BB0_1 Depth=1
.Ltmp15:
	leaq	48(%rsp), %rdi
	movq	%rdi, %rsi
	movl	$3, %edx
	callq	*%rax
.Ltmp16:
.LBB0_8:                                # %_ZNSt14_Function_baseD2Ev.exit66
                                        #   in Loop: Header=BB0_1 Depth=1
	cmpq	$47, %rbx
	ja	.LBB0_15
# %bb.9:                                # %.preheader.preheader
                                        #   in Loop: Header=BB0_1 Depth=1
	xorl	%r12d, %r12d
.LBB0_10:                               # %.preheader
                                        #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	%rax, %rbp
	movq	32(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdx
	movl	%r13d, %ecx
	movl	%r13d, %r8d
	movl	%r13d, %r9d
	callq	_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	%rax, %r14
	movq	40(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%r13d, %edx
	movl	%r13d, %ecx
	callq	_Z7maxDiffIfET_PS0_S1_ii
	andps	.LCPI0_0(%rip), %xmm0
	movss	.LCPI0_1(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	ucomiss	%xmm0, %xmm1
	movl	$.L.str.6, %esi
	jb	.LBB0_12
# %bb.11:                               # %.preheader
                                        #   in Loop: Header=BB0_10 Depth=2
	movl	$.L.str.5, %esi
.LBB0_12:                               # %.preheader
                                        #   in Loop: Header=BB0_10 Depth=2
	movl	$_ZSt4cout, %edi
	movl	$4, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	subq	%rbp, %r14
	movq	%r14, %rax
	movabsq	$2361183241434822607, %rcx      # imm = 0x20C49BA5E353F7CF
	imulq	%rcx
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$7, %rdx
	addq	%rax, %rdx
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rdx, %xmm0
	divss	.LCPI0_2(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movsd	%xmm0, 128(%rsp,%r12,8)
	incq	%r12
	cmpq	$5, %r12
	jne	.LBB0_10
# %bb.13:                               #   in Loop: Header=BB0_1 Depth=1
	movsd	128(%rsp), %xmm0                # xmm0 = mem[0],zero
	addsd	136(%rsp), %xmm0
	addsd	144(%rsp), %xmm0
	addsd	152(%rsp), %xmm0
	addsd	160(%rsp), %xmm0
	divsd	.LCPI0_3(%rip), %xmm0
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.7, %esi
	movl	$31, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.8, %esi
	movl	$59, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%r13d, %esi
	callq	_ZNSolsEi
	movq	%rax, %r14
	movl	$.L.str.9, %esi
	movl	$3, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
	movq	120(%rsp), %r12                 # 8-byte Reload
	jmp	.LBB0_16
.LBB0_14:                               #   in Loop: Header=BB0_1 Depth=1
	movl	$_ZSt4cout, %edi
	movl	$.L.str.13, %esi
	movl	$9, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	.L__const.main.my_list(,%rbx,4), %ecx
	movl	$.L.str.14, %esi
	movq	%r12, %rdi
	movq	%r14, %rdx
	movl	$800000, %r8d                   # imm = 0xC3500
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB0_17
.LBB0_15:                               #   in Loop: Header=BB0_1 Depth=1
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	%rax, %rbp
	movq	32(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdx
	movl	%r13d, %ecx
	movl	%r13d, %r8d
	movl	%r13d, %r9d
	callq	_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	%rax, %r14
	movq	40(%rsp), %rdi
	movq	16(%rsp), %rsi
	movl	%r13d, %edx
	movl	%r13d, %ecx
	callq	_Z7maxDiffIfET_PS0_S1_ii
	andps	.LCPI0_0(%rip), %xmm0
	movss	.LCPI0_1(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	ucomiss	%xmm0, %xmm1
	movl	$.L.str.5, %esi
	movl	$.L.str.6, %eax
	cmovbq	%rax, %rsi
	movl	$_ZSt4cout, %edi
	movl	$4, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	subq	%rbp, %r14
	movq	%r14, %rax
	movabsq	$2361183241434822607, %rcx      # imm = 0x20C49BA5E353F7CF
	imulq	%rcx
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$7, %rdx
	addq	%rax, %rdx
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rdx, %xmm0
	divss	.LCPI0_2(%rip), %xmm0
	movss	%xmm0, 8(%rsp)                  # 4-byte Spill
	movl	$_ZSt4cout, %edi
	movl	$.L.str.12, %esi
	movl	$43, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.8, %esi
	movl	$59, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%r13d, %esi
	callq	_ZNSolsEi
	movq	%rax, %r14
	movl	$.L.str.9, %esi
	movl	$3, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorps	%xmm0, %xmm0
	cvtss2sd	8(%rsp), %xmm0                  # 4-byte Folded Reload
	movq	%r14, %rdi
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r14
.LBB0_16:                               #   in Loop: Header=BB0_1 Depth=1
	movl	$.L.str.10, %esi
	movl	$3, %edx
	movq	%r14, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%r14), %rax
	movq	-24(%rax), %rdi
	addq	%r14, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r14, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$.L.str.11, %esi
	movq	%r12, %rdi
	leaq	176(%rsp), %r14
	movq	%r14, %rdx
	movl	%r13d, %ecx
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$1, %al
	callq	fprintf
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
.LBB0_17:                               #   in Loop: Header=BB0_1 Depth=1
	incq	%rbx
	cmpq	$266, %rbx                      # imm = 0x10A
	jne	.LBB0_1
# %bb.18:
	xorl	%eax, %eax
	addq	$280, %rsp                      # imm = 0x118
	.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
.LBB0_19:
	.cfi_def_cfa_offset 336
.Ltmp8:
	jmp	.LBB0_24
.LBB0_20:
.Ltmp17:
	jmp	.LBB0_24
.LBB0_21:
.Ltmp2:
	movq	%rax, %rbx
	movq	96(%rsp), %rax
	testq	%rax, %rax
	je	.LBB0_27
# %bb.22:
.Ltmp3:
	leaq	80(%rsp), %rdi
	movq	%rdi, %rsi
	movl	$3, %edx
	callq	*%rax
.Ltmp4:
	jmp	.LBB0_27
.LBB0_23:
.Ltmp5:
	jmp	.LBB0_24
.LBB0_25:
.Ltmp11:
	movq	%rax, %rbx
	movq	64(%rsp), %rax
	testq	%rax, %rax
	je	.LBB0_27
# %bb.26:
.Ltmp12:
	leaq	48(%rsp), %rdi
	movq	%rdi, %rsi
	movl	$3, %edx
	callq	*%rax
.Ltmp13:
.LBB0_27:                               # %_ZNSt14_Function_baseD2Ev.exit68
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.LBB0_28:
.Ltmp14:
.LBB0_24:
	movq	%rax, %rdi
	callq	__clang_call_terminate
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table0:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	3                               # @TType Encoding = udata4
	.uleb128 .Lttbase0-.Lttbaseref0
.Lttbaseref0:
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end0-.Lcst_begin0
.Lcst_begin0:
	.uleb128 .Lfunc_begin0-.Lfunc_begin0    # >> Call Site 1 <<
	.uleb128 .Ltmp0-.Lfunc_begin0           #   Call between .Lfunc_begin0 and .Ltmp0
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp0-.Lfunc_begin0           # >> Call Site 2 <<
	.uleb128 .Ltmp1-.Ltmp0                  #   Call between .Ltmp0 and .Ltmp1
	.uleb128 .Ltmp2-.Lfunc_begin0           #     jumps to .Ltmp2
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp6-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp7-.Ltmp6                  #   Call between .Ltmp6 and .Ltmp7
	.uleb128 .Ltmp8-.Lfunc_begin0           #     jumps to .Ltmp8
	.byte	1                               #   On action: 1
	.uleb128 .Ltmp9-.Lfunc_begin0           # >> Call Site 4 <<
	.uleb128 .Ltmp10-.Ltmp9                 #   Call between .Ltmp9 and .Ltmp10
	.uleb128 .Ltmp11-.Lfunc_begin0          #     jumps to .Ltmp11
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp15-.Lfunc_begin0          # >> Call Site 5 <<
	.uleb128 .Ltmp16-.Ltmp15                #   Call between .Ltmp15 and .Ltmp16
	.uleb128 .Ltmp17-.Lfunc_begin0          #     jumps to .Ltmp17
	.byte	1                               #   On action: 1
	.uleb128 .Ltmp16-.Lfunc_begin0          # >> Call Site 6 <<
	.uleb128 .Ltmp3-.Ltmp16                 #   Call between .Ltmp16 and .Ltmp3
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp3-.Lfunc_begin0           # >> Call Site 7 <<
	.uleb128 .Ltmp4-.Ltmp3                  #   Call between .Ltmp3 and .Ltmp4
	.uleb128 .Ltmp5-.Lfunc_begin0           #     jumps to .Ltmp5
	.byte	1                               #   On action: 1
	.uleb128 .Ltmp12-.Lfunc_begin0          # >> Call Site 8 <<
	.uleb128 .Ltmp13-.Ltmp12                #   Call between .Ltmp12 and .Ltmp13
	.uleb128 .Ltmp14-.Lfunc_begin0          #     jumps to .Ltmp14
	.byte	1                               #   On action: 1
	.uleb128 .Ltmp13-.Lfunc_begin0          # >> Call Site 9 <<
	.uleb128 .Lfunc_end0-.Ltmp13            #   Call between .Ltmp13 and .Lfunc_end0
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.byte	1                               # >> Action Record 1 <<
                                        #   Catch TypeInfo 1
	.byte	0                               #   No further actions
	.p2align	2, 0x0
                                        # >> Catch TypeInfos <<
	.long	0                               # TypeInfo 1
.Lttbase0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE,"axG",@progbits,_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE,comdat
	.weak	_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE # -- Begin function _Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE
	.type	_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE,@function
_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE: # @_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE
	.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	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, 8(%rsp)                   # 8-byte Spill
	testl	%esi, %esi
	jle	.LBB1_7
# %bb.1:                                # %.preheader.lr.ph
	movq	%rcx, %rbx
	movl	%esi, %eax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movl	%edx, %r12d
	xorl	%r13d, %r13d
	xorl	%r14d, %r14d
	movl	%edx, 4(%rsp)                   # 4-byte Spill
.LBB1_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_4 Depth 2
	testl	%edx, %edx
	jle	.LBB1_6
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB1_2 Depth=1
	movl	%r13d, %eax
	movq	8(%rsp), %rcx                   # 8-byte Reload
	leaq	(%rcx,%rax,4), %r15
	xorl	%ebp, %ebp
.LBB1_4:                                #   Parent Loop BB1_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	cmpq	$0, 16(%rbx)
	je	.LBB1_8
# %bb.5:                                # %_ZNKSt8functionIFfvEEclEv.exit
                                        #   in Loop: Header=BB1_4 Depth=2
	movq	%rbx, %rdi
	callq	*24(%rbx)
	movss	%xmm0, (%r15,%rbp,4)
	incq	%rbp
	cmpq	%rbp, %r12
	jne	.LBB1_4
.LBB1_6:                                # %._crit_edge
                                        #   in Loop: Header=BB1_2 Depth=1
	incq	%r14
	movl	4(%rsp), %edx                   # 4-byte Reload
	addl	%edx, %r13d
	cmpq	16(%rsp), %r14                  # 8-byte Folded Reload
	jne	.LBB1_2
.LBB1_7:                                # %._crit_edge13
	addq	$24, %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
.LBB1_8:
	.cfi_def_cfa_offset 80
	callq	_ZSt25__throw_bad_function_callv
.Lfunc_end1:
	.size	_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE, .Lfunc_end1-_Z17initialize_matrixIfEvPT_iiSt8functionIFfvEE
	.cfi_endproc
                                        # -- End function
	.section	.text._Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii,"axG",@progbits,_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii,comdat
	.weak	_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii # -- Begin function _Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii
	.type	_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii,@function
_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii: # @_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii
	.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
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdx, -8(%rsp)                  # 8-byte Spill
	movq	%rsi, -16(%rsp)                 # 8-byte Spill
	movq	%rdi, -24(%rsp)                 # 8-byte Spill
	testl	%r8d, %r8d
	jle	.LBB2_9
# %bb.1:                                # %.preheader27.lr.ph
	movslq	%r9d, %rdi
	movl	%r8d, %r8d
	movl	%edi, %r10d
	movl	%ecx, %r11d
	leaq	(,%rdi,4), %rbx
	xorl	%r14d, %r14d
	xorl	%r15d, %r15d
.LBB2_2:                                # %.preheader27
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_4 Depth 2
                                        #       Child Loop BB2_6 Depth 3
	testl	%r9d, %r9d
	jle	.LBB2_8
# %bb.3:                                # %.preheader.lr.ph
                                        #   in Loop: Header=BB2_2 Depth=1
	movl	%r14d, %eax
	movq	-24(%rsp), %rdx                 # 8-byte Reload
	leaq	(%rdx,%rax,4), %r12
	movq	%r15, %rax
	imulq	%rdi, %rax
	movq	-8(%rsp), %rdx                  # 8-byte Reload
	leaq	(%rdx,%rax,4), %r13
	movq	-16(%rsp), %rsi                 # 8-byte Reload
	xorl	%eax, %eax
.LBB2_4:                                # %.preheader
                                        #   Parent Loop BB2_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB2_6 Depth 3
	xorps	%xmm0, %xmm0
	testl	%ecx, %ecx
	jle	.LBB2_7
# %bb.5:                                # %.lr.ph.preheader
                                        #   in Loop: Header=BB2_4 Depth=2
	movq	%rsi, %rbp
	xorl	%edx, %edx
.LBB2_6:                                # %.lr.ph
                                        #   Parent Loop BB2_2 Depth=1
                                        #     Parent Loop BB2_4 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movss	(%r12,%rdx,4), %xmm1            # xmm1 = mem[0],zero,zero,zero
	mulss	(%rbp), %xmm1
	addss	%xmm1, %xmm0
	incq	%rdx
	addq	%rbx, %rbp
	cmpq	%rdx, %r11
	jne	.LBB2_6
.LBB2_7:                                # %._crit_edge
                                        #   in Loop: Header=BB2_4 Depth=2
	movss	%xmm0, (%r13,%rax,4)
	incq	%rax
	addq	$4, %rsi
	cmpq	%r10, %rax
	jne	.LBB2_4
.LBB2_8:                                # %._crit_edge31
                                        #   in Loop: Header=BB2_2 Depth=1
	incq	%r15
	addl	%ecx, %r14d
	cmpq	%r8, %r15
	jne	.LBB2_2
.LBB2_9:                                # %._crit_edge33
	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	_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii, .Lfunc_end2-_Z25naive_matrix_multiply_cpuIfEvPT_S1_S1_iii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _Z7maxDiffIfET_PS0_S1_ii
.LCPI3_0:
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.long	0x7fffffff                      # float NaN
	.section	.text._Z7maxDiffIfET_PS0_S1_ii,"axG",@progbits,_Z7maxDiffIfET_PS0_S1_ii,comdat
	.weak	_Z7maxDiffIfET_PS0_S1_ii
	.type	_Z7maxDiffIfET_PS0_S1_ii,@function
_Z7maxDiffIfET_PS0_S1_ii:               # @_Z7maxDiffIfET_PS0_S1_ii
	.cfi_startproc
# %bb.0:
	movss	(%rdi), %xmm2                   # xmm2 = mem[0],zero,zero,zero
	subss	(%rsi), %xmm2
	testl	%edx, %edx
	jle	.LBB3_1
# %bb.3:                                # %.preheader.lr.ph
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movl	%edx, %eax
	movl	%ecx, %edx
	xorl	%r8d, %r8d
	movaps	.LCPI3_0(%rip), %xmm1           # xmm1 = [NaN,NaN,NaN,NaN]
	xorl	%r9d, %r9d
.LBB3_4:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_9 Depth 2
	testl	%ecx, %ecx
	jle	.LBB3_5
# %bb.8:                                # %.lr.ph
                                        #   in Loop: Header=BB3_4 Depth=1
	movl	%r8d, %r11d
	leaq	(%rsi,%r11,4), %r10
	leaq	(%rdi,%r11,4), %r11
	xorl	%ebx, %ebx
.LBB3_9:                                #   Parent Loop BB3_4 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movss	(%r11,%rbx,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
	subss	(%r10,%rbx,4), %xmm0
	andps	%xmm1, %xmm0
	maxss	%xmm2, %xmm0
	incq	%rbx
	movaps	%xmm0, %xmm2
	cmpq	%rbx, %rdx
	jne	.LBB3_9
	jmp	.LBB3_6
.LBB3_5:                                #   in Loop: Header=BB3_4 Depth=1
	movaps	%xmm2, %xmm0
.LBB3_6:                                # %._crit_edge
                                        #   in Loop: Header=BB3_4 Depth=1
	incq	%r9
	addl	%ecx, %r8d
	movaps	%xmm0, %xmm2
	cmpq	%rax, %r9
	jne	.LBB3_4
# %bb.7:
	popq	%rbx
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	retq
.LBB3_1:
	movaps	%xmm2, %xmm0
	retq
.Lfunc_end3:
	.size	_Z7maxDiffIfET_PS0_S1_ii, .Lfunc_end3-_Z7maxDiffIfET_PS0_S1_ii
	.cfi_endproc
                                        # -- End function
	.section	.text.__clang_call_terminate,"axG",@progbits,__clang_call_terminate,comdat
	.hidden	__clang_call_terminate          # -- Begin function __clang_call_terminate
	.weak	__clang_call_terminate
	.type	__clang_call_terminate,@function
__clang_call_terminate:                 # @__clang_call_terminate
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__cxa_begin_catch
	callq	_ZSt9terminatev
.Lfunc_end4:
	.size	__clang_call_terminate, .Lfunc_end4-__clang_call_terminate
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data
.LCPI5_0:
	.long	0x4a03126c                      # float 2147483
	.text
	.p2align	1, 0x90
	.type	_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data,@function
_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data: # @_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	rand
	cvtsi2ss	%eax, %xmm0
	divss	.LCPI5_0(%rip), %xmm0
	popq	%rax
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data, .Lfunc_end5-_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data
	.cfi_endproc
                                        # -- End function
	.p2align	1, 0x90                         # -- Begin function _ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation
	.type	_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation,@function
_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation: # @_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation
	.cfi_startproc
# %bb.0:
	testl	%edx, %edx
	je	.LBB6_1
# %bb.2:
	cmpl	$1, %edx
	je	.LBB6_3
	jmp	.LBB6_4
.LBB6_1:
	movl	$_ZTIZ4mainEUlvE0_, %esi
.LBB6_3:                                # %_ZNSt14_Function_base13_Base_managerIZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation.exit.sink.split
	movq	%rsi, (%rdi)
.LBB6_4:                                # %_ZNSt14_Function_base13_Base_managerIZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation.exit
	xorl	%eax, %eax
	retq
.Lfunc_end6:
	.size	_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation, .Lfunc_end6-_ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Naive_Multiplication_CPU_Mohit.csv"
	.size	.L.str, 35

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"w+"
	.size	.L.str.1, 3

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Algorithm_Name,Input_Dimensions,Execution_Time(ms)"
	.size	.L.str.2, 51

	.type	.L__const.main.my_list,@object  # @__const.main.my_list
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
.L__const.main.my_list:
	.long	16                              # 0x10
	.long	32                              # 0x20
	.long	48                              # 0x30
	.long	64                              # 0x40
	.long	80                              # 0x50
	.long	96                              # 0x60
	.long	112                             # 0x70
	.long	128                             # 0x80
	.long	144                             # 0x90
	.long	160                             # 0xa0
	.long	176                             # 0xb0
	.long	192                             # 0xc0
	.long	208                             # 0xd0
	.long	224                             # 0xe0
	.long	240                             # 0xf0
	.long	256                             # 0x100
	.long	272                             # 0x110
	.long	288                             # 0x120
	.long	304                             # 0x130
	.long	320                             # 0x140
	.long	336                             # 0x150
	.long	352                             # 0x160
	.long	368                             # 0x170
	.long	384                             # 0x180
	.long	400                             # 0x190
	.long	416                             # 0x1a0
	.long	432                             # 0x1b0
	.long	448                             # 0x1c0
	.long	464                             # 0x1d0
	.long	480                             # 0x1e0
	.long	496                             # 0x1f0
	.long	512                             # 0x200
	.long	528                             # 0x210
	.long	544                             # 0x220
	.long	560                             # 0x230
	.long	576                             # 0x240
	.long	592                             # 0x250
	.long	608                             # 0x260
	.long	624                             # 0x270
	.long	640                             # 0x280
	.long	656                             # 0x290
	.long	672                             # 0x2a0
	.long	688                             # 0x2b0
	.long	704                             # 0x2c0
	.long	720                             # 0x2d0
	.long	736                             # 0x2e0
	.long	752                             # 0x2f0
	.long	768                             # 0x300
	.long	784                             # 0x310
	.long	800                             # 0x320
	.long	816                             # 0x330
	.long	832                             # 0x340
	.long	848                             # 0x350
	.long	864                             # 0x360
	.long	880                             # 0x370
	.long	896                             # 0x380
	.long	912                             # 0x390
	.long	928                             # 0x3a0
	.long	944                             # 0x3b0
	.long	960                             # 0x3c0
	.long	976                             # 0x3d0
	.long	992                             # 0x3e0
	.long	1008                            # 0x3f0
	.long	1040                            # 0x410
	.long	1072                            # 0x430
	.long	1104                            # 0x450
	.long	1136                            # 0x470
	.long	1168                            # 0x490
	.long	1200                            # 0x4b0
	.long	1232                            # 0x4d0
	.long	1264                            # 0x4f0
	.long	1296                            # 0x510
	.long	1328                            # 0x530
	.long	1360                            # 0x550
	.long	1392                            # 0x570
	.long	1424                            # 0x590
	.long	1456                            # 0x5b0
	.long	1488                            # 0x5d0
	.long	1520                            # 0x5f0
	.long	1552                            # 0x610
	.long	1584                            # 0x630
	.long	1616                            # 0x650
	.long	1648                            # 0x670
	.long	1680                            # 0x690
	.long	1712                            # 0x6b0
	.long	1744                            # 0x6d0
	.long	1776                            # 0x6f0
	.long	1808                            # 0x710
	.long	1840                            # 0x730
	.long	1872                            # 0x750
	.long	1904                            # 0x770
	.long	1936                            # 0x790
	.long	1968                            # 0x7b0
	.long	2000                            # 0x7d0
	.long	2032                            # 0x7f0
	.long	2064                            # 0x810
	.long	2096                            # 0x830
	.long	2128                            # 0x850
	.long	2160                            # 0x870
	.long	2192                            # 0x890
	.long	2224                            # 0x8b0
	.long	2256                            # 0x8d0
	.long	2288                            # 0x8f0
	.long	2320                            # 0x910
	.long	2352                            # 0x930
	.long	2384                            # 0x950
	.long	2416                            # 0x970
	.long	2448                            # 0x990
	.long	2480                            # 0x9b0
	.long	2512                            # 0x9d0
	.long	2544                            # 0x9f0
	.long	2576                            # 0xa10
	.long	2608                            # 0xa30
	.long	2640                            # 0xa50
	.long	2672                            # 0xa70
	.long	2704                            # 0xa90
	.long	2736                            # 0xab0
	.long	2768                            # 0xad0
	.long	2800                            # 0xaf0
	.long	2832                            # 0xb10
	.long	2864                            # 0xb30
	.long	2896                            # 0xb50
	.long	2928                            # 0xb70
	.long	2960                            # 0xb90
	.long	2992                            # 0xbb0
	.long	3024                            # 0xbd0
	.long	3056                            # 0xbf0
	.long	3088                            # 0xc10
	.long	3120                            # 0xc30
	.long	3152                            # 0xc50
	.long	3184                            # 0xc70
	.long	3216                            # 0xc90
	.long	3248                            # 0xcb0
	.long	3280                            # 0xcd0
	.long	3312                            # 0xcf0
	.long	3344                            # 0xd10
	.long	3376                            # 0xd30
	.long	3408                            # 0xd50
	.long	3440                            # 0xd70
	.long	3472                            # 0xd90
	.long	3504                            # 0xdb0
	.long	3536                            # 0xdd0
	.long	3568                            # 0xdf0
	.long	3600                            # 0xe10
	.long	3632                            # 0xe30
	.long	3664                            # 0xe50
	.long	3696                            # 0xe70
	.long	3728                            # 0xe90
	.long	3760                            # 0xeb0
	.long	3792                            # 0xed0
	.long	3824                            # 0xef0
	.long	3856                            # 0xf10
	.long	3888                            # 0xf30
	.long	3920                            # 0xf50
	.long	3952                            # 0xf70
	.long	3984                            # 0xf90
	.long	4016                            # 0xfb0
	.long	4080                            # 0xff0
	.long	4144                            # 0x1030
	.long	4208                            # 0x1070
	.long	4272                            # 0x10b0
	.long	4336                            # 0x10f0
	.long	4400                            # 0x1130
	.long	4464                            # 0x1170
	.long	4528                            # 0x11b0
	.long	4592                            # 0x11f0
	.long	4656                            # 0x1230
	.long	4720                            # 0x1270
	.long	4784                            # 0x12b0
	.long	4848                            # 0x12f0
	.long	4912                            # 0x1330
	.long	4976                            # 0x1370
	.long	5040                            # 0x13b0
	.long	5104                            # 0x13f0
	.long	5168                            # 0x1430
	.long	5232                            # 0x1470
	.long	5296                            # 0x14b0
	.long	5360                            # 0x14f0
	.long	5424                            # 0x1530
	.long	5488                            # 0x1570
	.long	5552                            # 0x15b0
	.long	5616                            # 0x15f0
	.long	5680                            # 0x1630
	.long	5744                            # 0x1670
	.long	5808                            # 0x16b0
	.long	5872                            # 0x16f0
	.long	5936                            # 0x1730
	.long	6000                            # 0x1770
	.long	6064                            # 0x17b0
	.long	6128                            # 0x17f0
	.long	6192                            # 0x1830
	.long	6256                            # 0x1870
	.long	6320                            # 0x18b0
	.long	6384                            # 0x18f0
	.long	6448                            # 0x1930
	.long	6512                            # 0x1970
	.long	6576                            # 0x19b0
	.long	6640                            # 0x19f0
	.long	6704                            # 0x1a30
	.long	6768                            # 0x1a70
	.long	6832                            # 0x1ab0
	.long	6896                            # 0x1af0
	.long	6960                            # 0x1b30
	.long	7024                            # 0x1b70
	.long	7088                            # 0x1bb0
	.long	7152                            # 0x1bf0
	.long	7216                            # 0x1c30
	.long	7280                            # 0x1c70
	.long	7344                            # 0x1cb0
	.long	7408                            # 0x1cf0
	.long	7472                            # 0x1d30
	.long	7536                            # 0x1d70
	.long	7600                            # 0x1db0
	.long	7664                            # 0x1df0
	.long	7728                            # 0x1e30
	.long	7792                            # 0x1e70
	.long	7856                            # 0x1eb0
	.long	7920                            # 0x1ef0
	.long	7984                            # 0x1f30
	.long	8048                            # 0x1f70
	.long	8112                            # 0x1fb0
	.long	8176                            # 0x1ff0
	.long	8240                            # 0x2030
	.long	8304                            # 0x2070
	.long	8368                            # 0x20b0
	.long	8432                            # 0x20f0
	.long	8496                            # 0x2130
	.long	8560                            # 0x2170
	.long	8624                            # 0x21b0
	.long	8688                            # 0x21f0
	.long	8752                            # 0x2230
	.long	8816                            # 0x2270
	.long	8880                            # 0x22b0
	.long	8944                            # 0x22f0
	.long	9008                            # 0x2330
	.long	9072                            # 0x2370
	.long	9136                            # 0x23b0
	.long	9200                            # 0x23f0
	.long	9264                            # 0x2430
	.long	9328                            # 0x2470
	.long	9392                            # 0x24b0
	.long	9456                            # 0x24f0
	.long	9520                            # 0x2530
	.long	9584                            # 0x2570
	.long	9648                            # 0x25b0
	.long	9712                            # 0x25f0
	.long	9776                            # 0x2630
	.long	9840                            # 0x2670
	.long	9904                            # 0x26b0
	.long	9968                            # 0x26f0
	.long	10032                           # 0x2730
	.long	10096                           # 0x2770
	.long	10160                           # 0x27b0
	.long	10224                           # 0x27f0
	.long	10288                           # 0x2830
	.long	10352                           # 0x2870
	.long	10416                           # 0x28b0
	.long	10480                           # 0x28f0
	.long	10544                           # 0x2930
	.long	10608                           # 0x2970
	.long	10672                           # 0x29b0
	.long	10736                           # 0x29f0
	.long	10800                           # 0x2a30
	.long	10864                           # 0x2a70
	.long	10928                           # 0x2ab0
	.long	10992                           # 0x2af0
	.size	.L__const.main.my_list, 1064

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

	.type	.L__const.main.algoname,@object # @__const.main.algoname
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
.L__const.main.algoname:
	.asciz	"naive_matrix_cpu\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000\000"
	.size	.L__const.main.algoname, 100

	.type	.L.str.4,@object                # @.str.4
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.4:
	.asciz	"Matrix Size : "
	.size	.L.str.4, 15

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"PASS"
	.size	.L.str.5, 5

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"FAIL"
	.size	.L.str.6, 5

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Using Milliseconds Clock: (AVG)"
	.size	.L.str.7, 32

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	" CPU time taken to execute for strassen matrices of size - "
	.size	.L.str.8, 60

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

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	" ms"
	.size	.L.str.10, 4

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"\n%s,%d,%lf"
	.size	.L.str.11, 11

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"Using Milliseconds Clock:(else above 1000) "
	.size	.L.str.12, 44

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"MAX LIMIT"
	.size	.L.str.13, 10

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"\n%s,%d,%d"
	.size	.L.str.14, 10

	.type	_ZTSZ4mainEUlvE0_,@object       # @_ZTSZ4mainEUlvE0_
	.section	.rodata,"a",@progbits
_ZTSZ4mainEUlvE0_:
	.asciz	"Z4mainEUlvE0_"
	.size	_ZTSZ4mainEUlvE0_, 14

	.type	_ZTIZ4mainEUlvE0_,@object       # @_ZTIZ4mainEUlvE0_
	.p2align	3, 0x0
_ZTIZ4mainEUlvE0_:
	.quad	_ZTVN10__cxxabiv117__class_type_infoE+16
	.quad	_ZTSZ4mainEUlvE0_
	.size	_ZTIZ4mainEUlvE0_, 16

	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym _ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E9_M_invokeERKSt9_Any_data
	.addrsig_sym _ZNSt17_Function_handlerIFfvEZ4mainEUlvE0_E10_M_managerERSt9_Any_dataRKS3_St18_Manager_operation
	.addrsig_sym _Unwind_Resume
	.addrsig_sym _ZSt4cout
	.addrsig_sym _ZTVN10__cxxabiv117__class_type_infoE
	.addrsig_sym _ZTSZ4mainEUlvE0_
	.addrsig_sym _ZTIZ4mainEUlvE0_
	.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
