#include <iostream>
#include <cuda_runtime.h>
#include <chrono>
#include <string>

#define N 13
//#define TRANSPOSED

typedef std::chrono::high_resolution_clock Clock;

int32_t matrixCGlobal[N][N];

int32_t matrixAGlobal[N * N] = {
        14, 39, 117, 89, 111, 73, 79, 102, 52, 81, 123, 70, 39,
        82, 29, 125, 85, 51, 60, 102, 39, 120, 106, 19, 15, 58,
        124, 31, 32, 23, 19, 69, 60, 61, 10, 33, 72, 1, 91,
        96, 112, 32, 111, 90, 12, 63, 77, 47, 105, 115, 38, 90,
        13, 35, 23, 78, 57, 109, 122, 89, 21, 116, 86, 123, 113,
        27, 14, 80, 69, 9, 23, 106, 26, 115, 31, 6, 73, 112,
        53, 70, 64, 118, 121, 17, 6, 113, 30, 8, 5, 116, 66,
        12, 113, 71, 94, 98, 116, 2, 95, 66, 107, 54, 11, 34,
        90, 36, 81, 124, 73, 41, 105, 14, 127, 109, 87, 29, 2,
        84, 77, 56, 81, 21, 81, 110, 110, 123, 104, 113, 39, 54,
        75, 102, 44, 79, 61, 55, 90, 125, 52, 45, 4, 120, 12,
        20, 20, 105, 41, 20, 44, 108, 74, 72, 62, 76, 34, 111,
        38, 97, 124, 5, 97, 87, 85, 106, 12, 31, 87, 6, 77
};

int32_t matrixBGlobal[N * N] = {
        69, 96, 71, 89, 127, 108, 96, 121, 64, 65, 62, 91, 73,
        9, 67, 113, 48, 47, 53, 96, 66, 7, 63, 17, 9, 8,
        107, 45, 112, 33, 114, 48, 102, 70, 52, 47, 34, 81, 17,
        38, 15, 61, 1, 104, 82, 68, 53, 69, 110, 12, 25, 46,
        111, 89, 54, 0, 107, 81, 127, 124, 36, 17, 99, 117, 75,
        125, 72, 48, 67, 31, 104, 64, 98, 94, 57, 81, 15, 16,
        111, 16, 127, 119, 88, 41, 75, 125, 22, 50, 120, 6, 81,
        75, 7, 78, 38, 35, 115, 114, 37, 66, 106, 64, 91, 97,
        75, 102, 84, 112, 65, 76, 87, 22, 45, 100, 19, 18, 89,
        27, 25, 109, 18, 116, 19, 116, 33, 103, 31, 29, 78, 8,
        24, 12, 86, 20, 32, 53, 31, 13, 51, 36, 100, 56, 44,
        13, 8, 54, 24, 101, 73, 115, 120, 56, 23, 63, 39, 93,
        77, 50, 108, 56, 106, 58, 121, 74, 70, 88, 19, 49, 83
};

int32_t matrixB_transposed[N * N];
// Calculates AB + A + B
//void naiveMatrixComputation() {
//    int8_t c, d, k;
//    int32_t sum;
//    for (c = 0; c < N; c++) {
//        for (d = 0; d < N; d++) {
//            sum = 0;
//            for (k = 0; k < N; ++k) {
//                sum += matrixAGlobal[c][k] * matrixBGlobal[k][d];
//            }
//            sum += matrixAGlobal[c][d] + matrixBGlobal[c][d];
//            matrixCGlobal[c][d] = sum;
//        }
//    }
//}

// Calculates AB + A + B
__global__ void
gpuMatrixComputation(const int32_t* matrixA, const int32_t* matrixB, int32_t* matrixC) {
    int8_t c, d, k;
    for (c = 0; c < N; ++c) {
        for (d = 0; d < N; ++d) {
            matrixC[c * N + d] = 0;
            for (k = 0; k < N; ++k) {
                matrixC[c * N + d] += matrixA[c * N + k] * matrixB[k * N + d];
            }
            matrixC[c * N + d] += matrixA[c * N + d] + matrixB[c * N + d];
        }
    }
}

// Calculates AB + A + B
__global__ void
gpuParallelMatrixComputation(const int32_t* matrixA, const int32_t* matrixB, int32_t* matrixC, const int size) {
    int x = blockIdx.x;
    int y = threadIdx.x;
    
    int sum = 0;
    
#pragma unroll
    for (int k = 0; k < size; ++k) {
#ifdef TRANSPOSED
        sum += matrixA[x * size + k] * matrixB[y * size + k];
#else
        sum += matrixA[x * size + k] * matrixB[k * size + y];
#endif
    }
    sum += matrixA[x * size + y] + matrixB[x * size + y];
    matrixC[x * size + y] = sum;
}

int main() {
    /*
     * Measure the time it takes to execute 1000 times
     */

    std::cout << "Malloc device mem" << std::endl;
    for (int c = 0; c < N; ++c)
        for (int d = 0; d < N; ++d)
            matrixB_transposed[c * N + d] = matrixBGlobal[d * N + c];


    int32_t* gpu_a, * gpu_b;
    int32_t* gpu_c;

    int32_t* c_out;
    c_out = (int32_t*)malloc(N * N * sizeof(int32_t));

    // We need variables accessible to the GPU,
    // so cudaMallocManaged provides these
    if (cudaMallocManaged(&gpu_a, N * N * sizeof(int32_t)) != 0) {
        std::cout << "malloc failed" << std::endl;
    }
    if (cudaMallocManaged(&gpu_b, N * N * sizeof(int32_t)) != 0) {
        std::cout << "malloc failed" << std::endl;
    }

    if (cudaMallocManaged(&gpu_c, N * N * sizeof(int32_t)) != 0) {
        std::cout << "malloc failed" << std::endl;
    }

    std::cout << "Move to device mem" << std::endl;

    if (cudaMemcpy(gpu_a, matrixAGlobal, (N * N * sizeof(int32_t)), cudaMemcpyHostToDevice) != 0) {
        std::cout << "memcpy failed" << std::endl;
    }

#ifdef TRANSPOSED
    if (cudaMemcpy(gpu_b, matrixB_transposed, (N * N * sizeof(int32_t)), cudaMemcpyHostToDevice) != 0) {
        std::cout << "memcpy failed" << std::endl;
    }
#else
    if (cudaMemcpy(gpu_b, matrixBGlobal, (N * N * sizeof(int32_t)), cudaMemcpyHostToDevice) != 0) {
        std::cout << "memcpy failed" << std::endl;
    }
#endif

    std::cout << "Computation" << std::endl;

    //Hier een goede keuze maken
    dim3 threadsPerBlock = 13;  // Should be a factor of 32
    dim3 numBlocks = 13;        // Very overkill but we now make use of all possible core

    unsigned long time_taken = 0;
    for (uint16_t i = 0; i < 1e4; i++) {
        auto cpu_start = Clock::now();
        //88593
       // gpuMatrixComputation<<<1, 1>>>(gpu_a, gpu_b, gpu_c);

        gpuParallelMatrixComputation <<<numBlocks, threadsPerBlock >>> (gpu_a, gpu_b, gpu_c, N);

        cudaDeviceSynchronize();
        auto cpu_end = Clock::now();

        time_taken += std::chrono::duration_cast<std::chrono::nanoseconds>(cpu_end - cpu_start).count();
    }
    /*
     * Print the resultating matrix
     */
    int8_t c, d;

    std::cout << "Copy from device to host" << std::endl;
    int cudaError = cudaMemcpy(c_out, gpu_c, (N * N * sizeof(int32_t)), cudaMemcpyDeviceToHost);
    if (cudaError != 0) {
        std::cout << "memcpy output failed: " << cudaError << std::endl;
    }

    std::cout << "Printing mem" << std::endl;

    for (c = 0; c < N; c++) {
        for (d = 0; d < N; d++) {
            std::cout << std::to_string((int32_t)c_out[c * N + d]) + "\t";
        }
        std::cout << "\n";
    }

    std::cout << "Freeing mem" << std::endl;

    cudaFree(gpu_a);
    cudaFree(gpu_b);
    cudaFree(gpu_c);

    std::cout << "Calculation took: " << std::to_string(time_taken / 1000) << " nanoseconds" << std::endl;

    return 0;
}
	.file	"tmpxft_0021874d_00000000-6_main.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3739:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3739:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4453:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4453:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.globl	_Z46__device_stub__Z20gpuMatrixComputationPKiS0_PiPKiS0_Pi
	.type	_Z46__device_stub__Z20gpuMatrixComputationPKiS0_PiPKiS0_Pi, @function
_Z46__device_stub__Z20gpuMatrixComputationPKiS0_PiPKiS0_Pi:
.LFB3761:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	40(%rsp), %rcx
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 96(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L3
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z20gpuMatrixComputationPKiS0_Pi(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	112(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L3:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L5
	call	__stack_chk_fail@PLT
.L5:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3761:
	.size	_Z46__device_stub__Z20gpuMatrixComputationPKiS0_PiPKiS0_Pi, .-_Z46__device_stub__Z20gpuMatrixComputationPKiS0_PiPKiS0_Pi
	.globl	_Z20gpuMatrixComputationPKiS0_Pi
	.type	_Z20gpuMatrixComputationPKiS0_Pi, @function
_Z20gpuMatrixComputationPKiS0_Pi:
.LFB3762:
	.cfi_startproc
	endbr64
	jmp	_Z46__device_stub__Z20gpuMatrixComputationPKiS0_PiPKiS0_Pi
	.cfi_endproc
.LFE3762:
	.size	_Z20gpuMatrixComputationPKiS0_Pi, .-_Z20gpuMatrixComputationPKiS0_Pi
	.globl	_Z55__device_stub__Z28gpuParallelMatrixComputationPKiS0_PiiPKiS0_Pii
	.type	_Z55__device_stub__Z28gpuParallelMatrixComputationPKiS0_PiiPKiS0_Pii, @function
_Z55__device_stub__Z28gpuParallelMatrixComputationPKiS0_PiiPKiS0_Pii:
.LFB3763:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z28gpuParallelMatrixComputationPKiS0_Pii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L9:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3763:
	.size	_Z55__device_stub__Z28gpuParallelMatrixComputationPKiS0_PiiPKiS0_Pii, .-_Z55__device_stub__Z28gpuParallelMatrixComputationPKiS0_PiiPKiS0_Pii
	.globl	_Z28gpuParallelMatrixComputationPKiS0_Pii
	.type	_Z28gpuParallelMatrixComputationPKiS0_Pii, @function
_Z28gpuParallelMatrixComputationPKiS0_Pii:
.LFB3764:
	.cfi_startproc
	endbr64
	jmp	_Z55__device_stub__Z28gpuParallelMatrixComputationPKiS0_PiiPKiS0_Pii
	.cfi_endproc
.LFE3764:
	.size	_Z28gpuParallelMatrixComputationPKiS0_Pii, .-_Z28gpuParallelMatrixComputationPKiS0_Pii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Malloc device mem"
.LC1:
	.string	"malloc failed"
.LC2:
	.string	"Move to device mem"
.LC3:
	.string	"memcpy failed"
.LC4:
	.string	"Computation"
.LC5:
	.string	"Copy from device to host"
.LC6:
	.string	"memcpy output failed: "
.LC7:
	.string	"Printing mem"
.LC9:
	.string	"\t"
.LC10:
	.string	"\n"
.LC11:
	.string	"Freeing mem"
.LC12:
	.string	"Calculation took: "
.LC13:
	.string	" nanoseconds"
.LC8:
	.string	"00010203040506070809101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3734:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA3734
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC0(%rip), %rsi
	leaq	_ZSt4cout(%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
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$376, %rsp
	.cfi_def_cfa_offset 432
	movq	%fs:40, %rax
	movq	%rax, 360(%rsp)
	xorl	%eax, %eax
.LEHB0:
	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	matrixBGlobal(%rip), %rdx
	leaq	matrixB_transposed(%rip), %rdi
	xorl	%esi, %esi
	movq	%rdx, %rbx
.L15:
	movq	%rdx, %rcx
	xorl	%eax, %eax
.L16:
	movl	(%rcx), %r8d
	addq	$52, %rcx
	movl	%r8d, (%rdi,%rax,4)
	incq	%rax
	cmpq	$13, %rax
	jne	.L16
	incl	%esi
	addq	$4, %rdx
	addq	$52, %rdi
	cmpl	$13, %esi
	jne	.L15
	movl	$676, %edi
	call	malloc@PLT
	leaq	40(%rsp), %rdi
	movl	$1, %edx
	movl	$676, %esi
	movq	%rax, %r13
	call	cudaMallocManaged@PLT
	testl	%eax, %eax
	je	.L18
	leaq	.LC1(%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
.L18:
	leaq	48(%rsp), %rdi
	movl	$1, %edx
	movl	$676, %esi
	call	cudaMallocManaged@PLT
	testl	%eax, %eax
	je	.L19
	leaq	.LC1(%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
.L19:
	leaq	56(%rsp), %rdi
	movl	$1, %edx
	movl	$676, %esi
	call	cudaMallocManaged@PLT
	testl	%eax, %eax
	je	.L20
	leaq	.LC1(%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
.L20:
	leaq	.LC2(%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	40(%rsp), %rdi
	movl	$1, %ecx
	movl	$676, %edx
	leaq	matrixAGlobal(%rip), %rsi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L21
	leaq	.LC3(%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
.L21:
	movq	48(%rsp), %rdi
	movl	$1, %ecx
	movl	$676, %edx
	movq	%rbx, %rsi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L22
	leaq	.LC3(%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
.L22:
	leaq	.LC4(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movl	$10000, %ebp
	xorl	%ebx, %ebx
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$1, 72(%rsp)
	movabsq	$4294967309, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
.L24:
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movl	72(%rsp), %ecx
	movl	84(%rsp), %esi
	xorl	%r9d, %r9d
	movq	64(%rsp), %rdx
	movq	76(%rsp), %rdi
	xorl	%r8d, %r8d
	movq	%rax, %r12
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L23
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	movl	$13, %ecx
	movq	40(%rsp), %rdi
	call	_Z55__device_stub__Z28gpuParallelMatrixComputationPKiS0_PiiPKiS0_Pii
.L23:
	call	cudaDeviceSynchronize@PLT
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	subq	%r12, %rax
	addq	%rax, %rbx
	decw	%bp
	jne	.L24
	leaq	.LC5(%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	56(%rsp), %rsi
	movl	$2, %ecx
	movq	%r13, %rdi
	movl	$676, %edx
	call	cudaMemcpy@PLT
	movl	%eax, %ebp
	testl	%eax, %eax
	je	.L25
	leaq	.LC6(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebp, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L25:
	leaq	.LC7(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	xorl	%r14d, %r14d
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.LEHE0:
.L26:
	imulq	$52, %r14, %r15
	xorl	%r12d, %r12d
	addq	%r13, %r15
.L37:
	movl	(%r15,%r12,4), %eax
	movl	$1, %r9d
	movl	$10000, %ecx
	movl	%eax, %r8d
	movl	%eax, %ebp
	shrl	$31, %r8d
	negl	%ebp
	cmovs	%eax, %ebp
	movl	%ebp, %eax
.L32:
	cmpl	$9, %eax
	jbe	.L28
	cmpl	$99, %eax
	ja	.L29
	incl	%r9d
	jmp	.L28
.L29:
	cmpl	$999, %eax
	ja	.L30
	addl	$2, %r9d
	jmp	.L28
.L30:
	cmpl	$9999, %eax
	ja	.L31
	addl	$3, %r9d
	jmp	.L28
.L31:
	movl	%eax, %eax
	xorl	%edx, %edx
	addl	$4, %r9d
	divq	%rcx
	jmp	.L32
.L28:
	leaq	88(%rsp), %rax
	leal	(%r8,%r9), %esi
	movl	$45, %edx
	movl	%r9d, 28(%rsp)
	movq	%rax, 8(%rsp)
	movq	8(%rsp), %rdi
	leaq	104(%rsp), %rax
	movq	%rax, 88(%rsp)
	movl	%r8d, 16(%rsp)
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructEmc@PLT
	movl	16(%rsp), %r8d
	leaq	159(%rsp), %rax
	movl	28(%rsp), %r9d
	movq	%rax, 16(%rsp)
	leaq	159(%rsp), %rdi
	movl	$201, %ecx
	leaq	.LC8(%rip), %rsi
	rep movsb
	movzbl	%r8b, %r8d
	decl	%r9d
	addq	88(%rsp), %r8
	movl	$100, %ecx
.L33:
	cmpl	$99, %ebp
	jbe	.L82
	movl	%ebp, %eax
	xorl	%edx, %edx
	movl	%r9d, %esi
	divl	%ecx
	addl	%edx, %edx
	movl	%eax, %ebp
	leal	1(%rdx), %eax
	movb	159(%rsp,%rdx), %dl
	movb	159(%rsp,%rax), %al
	movb	%al, (%r8,%rsi)
	leal	-1(%r9), %eax
	subl	$2, %r9d
	movb	%dl, (%r8,%rax)
	jmp	.L33
.L82:
	leal	48(%rbp), %eax
	cmpl	$9, %ebp
	jbe	.L36
	addl	%ebp, %ebp
	leal	1(%rbp), %eax
	movb	159(%rsp,%rax), %al
	movb	%al, 1(%r8)
	movb	159(%rsp,%rbp), %al
.L36:
	movb	%al, (%r8)
	movq	8(%rsp), %rdi
	leaq	.LC9(%rip), %rsi
.LEHB1:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKc@PLT
.LEHE1:
	leaq	120(%rsp), %rbp
	movq	%rax, %rsi
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC1EOS4_@PLT
	movq	%rbp, %rsi
	leaq	_ZSt4cout(%rip), %rdi
.LEHB2:
	call	_ZStlsIcSt11char_traitsIcESaIcEERSt13basic_ostreamIT_T0_ES7_RKNSt7__cxx1112basic_stringIS4_S5_T1_EE@PLT
.LEHE2:
	movq	%rbp, %rdi
	incq	%r12
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	8(%rsp), %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	cmpq	$13, %r12
	jne	.L37
	leaq	.LC10(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	incq	%r14
.LEHB3:
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$13, %r14
	jne	.L26
	leaq	.LC11(%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	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	leaq	.LC12(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$1000, %ecx
	xorl	%edx, %edx
	movl	$1, %esi
	movq	%rax, %r12
	movq	%rbx, %rax
	divq	%rcx
	movl	$10000, %ecx
	movq	%rax, %rbx
.L43:
	cmpq	$9, %rax
	jbe	.L39
	cmpq	$99, %rax
	ja	.L40
	incl	%esi
	jmp	.L39
.L40:
	cmpq	$999, %rax
	ja	.L41
	addl	$2, %esi
	jmp	.L39
.L41:
	cmpq	$9999, %rax
	ja	.L42
	addl	$3, %esi
	jmp	.L39
.L42:
	xorl	%edx, %edx
	addl	$4, %esi
	divq	%rcx
	jmp	.L43
.L39:
	leaq	136(%rsp), %rax
	movq	%rbp, %rdi
	xorl	%edx, %edx
	movq	%rax, 120(%rsp)
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructEmc@PLT
.LEHE3:
	movq	120(%rsp), %r8
	movq	128(%rsp), %rax
	leaq	.LC8(%rip), %rsi
	movq	16(%rsp), %rdi
	movl	$201, %ecx
	rep movsb
	leal	-1(%rax), %ecx
	movl	$100, %esi
.L44:
	cmpq	$99, %rbx
	jbe	.L83
	movq	%rbx, %rax
	xorl	%edx, %edx
	divq	%rsi
	addq	%rdx, %rdx
	movq	%rax, %rbx
	movl	%ecx, %eax
	movb	160(%rsp,%rdx), %dil
	movb	159(%rsp,%rdx), %dl
	movb	%dil, (%r8,%rax)
	leal	-1(%rcx), %eax
	subl	$2, %ecx
	movb	%dl, (%r8,%rax)
	jmp	.L44
.L83:
	leal	48(%rbx), %eax
	cmpq	$9, %rbx
	jbe	.L47
	addq	%rbx, %rbx
	movb	160(%rsp,%rbx), %al
	movb	%al, 1(%r8)
	movb	159(%rsp,%rbx), %al
.L47:
	movb	%al, (%r8)
	movq	%rbp, %rsi
	movq	%r12, %rdi
.LEHB4:
	call	_ZStlsIcSt11char_traitsIcESaIcEERSt13basic_ostreamIT_T0_ES7_RKNSt7__cxx1112basic_stringIS4_S5_T1_EE@PLT
	movq	%rax, %rdi
	leaq	.LC13(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.LEHE4:
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	360(%rsp), %rax
	subq	%fs:40, %rax
	je	.L53
	jmp	.L78
.L55:
	endbr64
	movq	%rax, %rbx
.L48:
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	jmp	.L49
.L54:
	endbr64
	movq	%rax, %rbx
.L49:
	movq	8(%rsp), %rdi
	jmp	.L81
.L56:
	endbr64
	movq	%rax, %rbx
.L51:
	movq	%rbp, %rdi
.L81:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	360(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L78
	movq	%rbx, %rdi
.LEHB5:
	call	_Unwind_Resume@PLT
.LEHE5:
.L78:
	call	__stack_chk_fail@PLT
.L53:
	addq	$376, %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
.LFE3734:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA3734:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE3734-.LLSDACSB3734
.LLSDACSB3734:
	.uleb128 .LEHB0-.LFB3734
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB3734
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L54-.LFB3734
	.uleb128 0
	.uleb128 .LEHB2-.LFB3734
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L55-.LFB3734
	.uleb128 0
	.uleb128 .LEHB3-.LFB3734
	.uleb128 .LEHE3-.LEHB3
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB4-.LFB3734
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L56-.LFB3734
	.uleb128 0
	.uleb128 .LEHB5-.LFB3734
	.uleb128 .LEHE5-.LEHB5
	.uleb128 0
	.uleb128 0
.LLSDACSE3734:
	.section	.text.startup
	.size	main, .-main
	.section	.rodata.str1.1
.LC14:
	.string	"_Z28gpuParallelMatrixComputationPKiS0_Pii"
.LC15:
	.string	"_Z20gpuMatrixComputationPKiS0_Pi"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3766:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC14(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z28gpuParallelMatrixComputationPKiS0_Pii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC15(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z20gpuMatrixComputationPKiS0_Pi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3766:
	.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	matrixB_transposed
	.bss
	.align 32
	.type	matrixB_transposed, @object
	.size	matrixB_transposed, 676
matrixB_transposed:
	.zero	676
	.globl	matrixBGlobal
	.data
	.align 32
	.type	matrixBGlobal, @object
	.size	matrixBGlobal, 676
matrixBGlobal:
	.long	69
	.long	96
	.long	71
	.long	89
	.long	127
	.long	108
	.long	96
	.long	121
	.long	64
	.long	65
	.long	62
	.long	91
	.long	73
	.long	9
	.long	67
	.long	113
	.long	48
	.long	47
	.long	53
	.long	96
	.long	66
	.long	7
	.long	63
	.long	17
	.long	9
	.long	8
	.long	107
	.long	45
	.long	112
	.long	33
	.long	114
	.long	48
	.long	102
	.long	70
	.long	52
	.long	47
	.long	34
	.long	81
	.long	17
	.long	38
	.long	15
	.long	61
	.long	1
	.long	104
	.long	82
	.long	68
	.long	53
	.long	69
	.long	110
	.long	12
	.long	25
	.long	46
	.long	111
	.long	89
	.long	54
	.long	0
	.long	107
	.long	81
	.long	127
	.long	124
	.long	36
	.long	17
	.long	99
	.long	117
	.long	75
	.long	125
	.long	72
	.long	48
	.long	67
	.long	31
	.long	104
	.long	64
	.long	98
	.long	94
	.long	57
	.long	81
	.long	15
	.long	16
	.long	111
	.long	16
	.long	127
	.long	119
	.long	88
	.long	41
	.long	75
	.long	125
	.long	22
	.long	50
	.long	120
	.long	6
	.long	81
	.long	75
	.long	7
	.long	78
	.long	38
	.long	35
	.long	115
	.long	114
	.long	37
	.long	66
	.long	106
	.long	64
	.long	91
	.long	97
	.long	75
	.long	102
	.long	84
	.long	112
	.long	65
	.long	76
	.long	87
	.long	22
	.long	45
	.long	100
	.long	19
	.long	18
	.long	89
	.long	27
	.long	25
	.long	109
	.long	18
	.long	116
	.long	19
	.long	116
	.long	33
	.long	103
	.long	31
	.long	29
	.long	78
	.long	8
	.long	24
	.long	12
	.long	86
	.long	20
	.long	32
	.long	53
	.long	31
	.long	13
	.long	51
	.long	36
	.long	100
	.long	56
	.long	44
	.long	13
	.long	8
	.long	54
	.long	24
	.long	101
	.long	73
	.long	115
	.long	120
	.long	56
	.long	23
	.long	63
	.long	39
	.long	93
	.long	77
	.long	50
	.long	108
	.long	56
	.long	106
	.long	58
	.long	121
	.long	74
	.long	70
	.long	88
	.long	19
	.long	49
	.long	83
	.globl	matrixAGlobal
	.align 32
	.type	matrixAGlobal, @object
	.size	matrixAGlobal, 676
matrixAGlobal:
	.long	14
	.long	39
	.long	117
	.long	89
	.long	111
	.long	73
	.long	79
	.long	102
	.long	52
	.long	81
	.long	123
	.long	70
	.long	39
	.long	82
	.long	29
	.long	125
	.long	85
	.long	51
	.long	60
	.long	102
	.long	39
	.long	120
	.long	106
	.long	19
	.long	15
	.long	58
	.long	124
	.long	31
	.long	32
	.long	23
	.long	19
	.long	69
	.long	60
	.long	61
	.long	10
	.long	33
	.long	72
	.long	1
	.long	91
	.long	96
	.long	112
	.long	32
	.long	111
	.long	90
	.long	12
	.long	63
	.long	77
	.long	47
	.long	105
	.long	115
	.long	38
	.long	90
	.long	13
	.long	35
	.long	23
	.long	78
	.long	57
	.long	109
	.long	122
	.long	89
	.long	21
	.long	116
	.long	86
	.long	123
	.long	113
	.long	27
	.long	14
	.long	80
	.long	69
	.long	9
	.long	23
	.long	106
	.long	26
	.long	115
	.long	31
	.long	6
	.long	73
	.long	112
	.long	53
	.long	70
	.long	64
	.long	118
	.long	121
	.long	17
	.long	6
	.long	113
	.long	30
	.long	8
	.long	5
	.long	116
	.long	66
	.long	12
	.long	113
	.long	71
	.long	94
	.long	98
	.long	116
	.long	2
	.long	95
	.long	66
	.long	107
	.long	54
	.long	11
	.long	34
	.long	90
	.long	36
	.long	81
	.long	124
	.long	73
	.long	41
	.long	105
	.long	14
	.long	127
	.long	109
	.long	87
	.long	29
	.long	2
	.long	84
	.long	77
	.long	56
	.long	81
	.long	21
	.long	81
	.long	110
	.long	110
	.long	123
	.long	104
	.long	113
	.long	39
	.long	54
	.long	75
	.long	102
	.long	44
	.long	79
	.long	61
	.long	55
	.long	90
	.long	125
	.long	52
	.long	45
	.long	4
	.long	120
	.long	12
	.long	20
	.long	20
	.long	105
	.long	41
	.long	20
	.long	44
	.long	108
	.long	74
	.long	72
	.long	62
	.long	76
	.long	34
	.long	111
	.long	38
	.long	97
	.long	124
	.long	5
	.long	97
	.long	87
	.long	85
	.long	106
	.long	12
	.long	31
	.long	87
	.long	6
	.long	77
	.globl	matrixCGlobal
	.bss
	.align 32
	.type	matrixCGlobal, @object
	.size	matrixCGlobal, 676
matrixCGlobal:
	.zero	676
	.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 <hip/hip_runtime.h>
#include <chrono>
#include <string>

#define N 13
//#define TRANSPOSED

typedef std::chrono::high_resolution_clock Clock;

int32_t matrixCGlobal[N][N];

int32_t matrixAGlobal[N * N] = {
        14, 39, 117, 89, 111, 73, 79, 102, 52, 81, 123, 70, 39,
        82, 29, 125, 85, 51, 60, 102, 39, 120, 106, 19, 15, 58,
        124, 31, 32, 23, 19, 69, 60, 61, 10, 33, 72, 1, 91,
        96, 112, 32, 111, 90, 12, 63, 77, 47, 105, 115, 38, 90,
        13, 35, 23, 78, 57, 109, 122, 89, 21, 116, 86, 123, 113,
        27, 14, 80, 69, 9, 23, 106, 26, 115, 31, 6, 73, 112,
        53, 70, 64, 118, 121, 17, 6, 113, 30, 8, 5, 116, 66,
        12, 113, 71, 94, 98, 116, 2, 95, 66, 107, 54, 11, 34,
        90, 36, 81, 124, 73, 41, 105, 14, 127, 109, 87, 29, 2,
        84, 77, 56, 81, 21, 81, 110, 110, 123, 104, 113, 39, 54,
        75, 102, 44, 79, 61, 55, 90, 125, 52, 45, 4, 120, 12,
        20, 20, 105, 41, 20, 44, 108, 74, 72, 62, 76, 34, 111,
        38, 97, 124, 5, 97, 87, 85, 106, 12, 31, 87, 6, 77
};

int32_t matrixBGlobal[N * N] = {
        69, 96, 71, 89, 127, 108, 96, 121, 64, 65, 62, 91, 73,
        9, 67, 113, 48, 47, 53, 96, 66, 7, 63, 17, 9, 8,
        107, 45, 112, 33, 114, 48, 102, 70, 52, 47, 34, 81, 17,
        38, 15, 61, 1, 104, 82, 68, 53, 69, 110, 12, 25, 46,
        111, 89, 54, 0, 107, 81, 127, 124, 36, 17, 99, 117, 75,
        125, 72, 48, 67, 31, 104, 64, 98, 94, 57, 81, 15, 16,
        111, 16, 127, 119, 88, 41, 75, 125, 22, 50, 120, 6, 81,
        75, 7, 78, 38, 35, 115, 114, 37, 66, 106, 64, 91, 97,
        75, 102, 84, 112, 65, 76, 87, 22, 45, 100, 19, 18, 89,
        27, 25, 109, 18, 116, 19, 116, 33, 103, 31, 29, 78, 8,
        24, 12, 86, 20, 32, 53, 31, 13, 51, 36, 100, 56, 44,
        13, 8, 54, 24, 101, 73, 115, 120, 56, 23, 63, 39, 93,
        77, 50, 108, 56, 106, 58, 121, 74, 70, 88, 19, 49, 83
};

int32_t matrixB_transposed[N * N];
// Calculates AB + A + B
//void naiveMatrixComputation() {
//    int8_t c, d, k;
//    int32_t sum;
//    for (c = 0; c < N; c++) {
//        for (d = 0; d < N; d++) {
//            sum = 0;
//            for (k = 0; k < N; ++k) {
//                sum += matrixAGlobal[c][k] * matrixBGlobal[k][d];
//            }
//            sum += matrixAGlobal[c][d] + matrixBGlobal[c][d];
//            matrixCGlobal[c][d] = sum;
//        }
//    }
//}

// Calculates AB + A + B
__global__ void
gpuMatrixComputation(const int32_t* matrixA, const int32_t* matrixB, int32_t* matrixC) {
    int8_t c, d, k;
    for (c = 0; c < N; ++c) {
        for (d = 0; d < N; ++d) {
            matrixC[c * N + d] = 0;
            for (k = 0; k < N; ++k) {
                matrixC[c * N + d] += matrixA[c * N + k] * matrixB[k * N + d];
            }
            matrixC[c * N + d] += matrixA[c * N + d] + matrixB[c * N + d];
        }
    }
}

// Calculates AB + A + B
__global__ void
gpuParallelMatrixComputation(const int32_t* matrixA, const int32_t* matrixB, int32_t* matrixC, const int size) {
    int x = blockIdx.x;
    int y = threadIdx.x;
    
    int sum = 0;
    
#pragma unroll
    for (int k = 0; k < size; ++k) {
#ifdef TRANSPOSED
        sum += matrixA[x * size + k] * matrixB[y * size + k];
#else
        sum += matrixA[x * size + k] * matrixB[k * size + y];
#endif
    }
    sum += matrixA[x * size + y] + matrixB[x * size + y];
    matrixC[x * size + y] = sum;
}

int main() {
    /*
     * Measure the time it takes to execute 1000 times
     */

    std::cout << "Malloc device mem" << std::endl;
    for (int c = 0; c < N; ++c)
        for (int d = 0; d < N; ++d)
            matrixB_transposed[c * N + d] = matrixBGlobal[d * N + c];


    int32_t* gpu_a, * gpu_b;
    int32_t* gpu_c;

    int32_t* c_out;
    c_out = (int32_t*)malloc(N * N * sizeof(int32_t));

    // We need variables accessible to the GPU,
    // so cudaMallocManaged provides these
    if (hipMallocManaged(&gpu_a, N * N * sizeof(int32_t)) != 0) {
        std::cout << "malloc failed" << std::endl;
    }
    if (hipMallocManaged(&gpu_b, N * N * sizeof(int32_t)) != 0) {
        std::cout << "malloc failed" << std::endl;
    }

    if (hipMallocManaged(&gpu_c, N * N * sizeof(int32_t)) != 0) {
        std::cout << "malloc failed" << std::endl;
    }

    std::cout << "Move to device mem" << std::endl;

    if (hipMemcpy(gpu_a, matrixAGlobal, (N * N * sizeof(int32_t)), hipMemcpyHostToDevice) != 0) {
        std::cout << "memcpy failed" << std::endl;
    }

#ifdef TRANSPOSED
    if (hipMemcpy(gpu_b, matrixB_transposed, (N * N * sizeof(int32_t)), hipMemcpyHostToDevice) != 0) {
        std::cout << "memcpy failed" << std::endl;
    }
#else
    if (hipMemcpy(gpu_b, matrixBGlobal, (N * N * sizeof(int32_t)), hipMemcpyHostToDevice) != 0) {
        std::cout << "memcpy failed" << std::endl;
    }
#endif

    std::cout << "Computation" << std::endl;

    //Hier een goede keuze maken
    dim3 threadsPerBlock = 13;  // Should be a factor of 32
    dim3 numBlocks = 13;        // Very overkill but we now make use of all possible core

    unsigned long time_taken = 0;
    for (uint16_t i = 0; i < 1e4; i++) {
        auto cpu_start = Clock::now();
        //88593
       // gpuMatrixComputation<<<1, 1>>>(gpu_a, gpu_b, gpu_c);

        gpuParallelMatrixComputation <<<numBlocks, threadsPerBlock >>> (gpu_a, gpu_b, gpu_c, N);

        hipDeviceSynchronize();
        auto cpu_end = Clock::now();

        time_taken += std::chrono::duration_cast<std::chrono::nanoseconds>(cpu_end - cpu_start).count();
    }
    /*
     * Print the resultating matrix
     */
    int8_t c, d;

    std::cout << "Copy from device to host" << std::endl;
    int hipError_t = hipMemcpy(c_out, gpu_c, (N * N * sizeof(int32_t)), hipMemcpyDeviceToHost);
    if (hipError_t != 0) {
        std::cout << "memcpy output failed: " << hipError_t << std::endl;
    }

    std::cout << "Printing mem" << std::endl;

    for (c = 0; c < N; c++) {
        for (d = 0; d < N; d++) {
            std::cout << std::to_string((int32_t)c_out[c * N + d]) + "\t";
        }
        std::cout << "\n";
    }

    std::cout << "Freeing mem" << std::endl;

    hipFree(gpu_a);
    hipFree(gpu_b);
    hipFree(gpu_c);

    std::cout << "Calculation took: " << std::to_string(time_taken / 1000) << " nanoseconds" << std::endl;

    return 0;
}
	.text
	.file	"main.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z35__device_stub__gpuMatrixComputationPKiS0_Pi # -- Begin function _Z35__device_stub__gpuMatrixComputationPKiS0_Pi
	.type	_Z35__device_stub__gpuMatrixComputationPKiS0_Pi,@function
_Z35__device_stub__gpuMatrixComputationPKiS0_Pi: # @_Z35__device_stub__gpuMatrixComputationPKiS0_Pi
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	movq	%rsi, 16(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z20gpuMatrixComputationPKiS0_Pi, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z35__device_stub__gpuMatrixComputationPKiS0_Pi, .Lfunc_end0-_Z35__device_stub__gpuMatrixComputationPKiS0_Pi
	.cfi_endproc
                                        # -- End function
	.globl	_Z43__device_stub__gpuParallelMatrixComputationPKiS0_Pii # -- Begin function _Z43__device_stub__gpuParallelMatrixComputationPKiS0_Pii
	.type	_Z43__device_stub__gpuParallelMatrixComputationPKiS0_Pii,@function
_Z43__device_stub__gpuParallelMatrixComputationPKiS0_Pii: # @_Z43__device_stub__gpuParallelMatrixComputationPKiS0_Pii
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z28gpuParallelMatrixComputationPKiS0_Pii, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z43__device_stub__gpuParallelMatrixComputationPKiS0_Pii, .Lfunc_end1-_Z43__device_stub__gpuParallelMatrixComputationPKiS0_Pii
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function 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	$120, %rsp
	.cfi_def_cfa_offset 176
	.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	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$17, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$matrixB_transposed, %eax
	movl	$matrixBGlobal, %ecx
	xorl	%edx, %edx
.LBB2_1:                                # %.preheader81
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_2 Depth 2
	movq	%rcx, %rsi
	xorl	%edi, %edi
.LBB2_2:                                #   Parent Loop BB2_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%rsi), %r8d
	movl	%r8d, (%rax,%rdi,4)
	incq	%rdi
	addq	$52, %rsi
	cmpq	$13, %rdi
	jne	.LBB2_2
# %bb.3:                                #   in Loop: Header=BB2_1 Depth=1
	incq	%rdx
	addq	$52, %rax
	addq	$4, %rcx
	cmpq	$13, %rdx
	jne	.LBB2_1
# %bb.4:
	movl	$676, %edi                      # imm = 0x2A4
	callq	malloc
	movq	%rax, 64(%rsp)                  # 8-byte Spill
	leaq	56(%rsp), %rdi
	movl	$676, %esi                      # imm = 0x2A4
	movl	$1, %edx
	callq	hipMallocManaged
	testl	%eax, %eax
	je	.LBB2_6
# %bb.5:
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$13, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB2_6:
	leaq	48(%rsp), %rdi
	movl	$676, %esi                      # imm = 0x2A4
	movl	$1, %edx
	callq	hipMallocManaged
	testl	%eax, %eax
	je	.LBB2_8
# %bb.7:
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$13, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB2_8:
	leaq	40(%rsp), %rdi
	movl	$676, %esi                      # imm = 0x2A4
	movl	$1, %edx
	callq	hipMallocManaged
	testl	%eax, %eax
	je	.LBB2_10
# %bb.9:
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$13, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB2_10:
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.2, %esi
	movl	$18, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	56(%rsp), %rdi
	movl	$matrixAGlobal, %esi
	movl	$676, %edx                      # imm = 0x2A4
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB2_12
# %bb.11:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$13, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB2_12:
	movq	48(%rsp), %rdi
	movl	$matrixBGlobal, %esi
	movl	$676, %edx                      # imm = 0x2A4
	movl	$1, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB2_14
# %bb.13:
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$13, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.LBB2_14:
	movabsq	$4294967309, %r14               # imm = 0x10000000D
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$11, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movw	$10000, %bx                     # imm = 0x2710
	xorl	%eax, %eax
	movq	%rax, 32(%rsp)                  # 8-byte Spill
.LBB2_15:                               # =>This Inner Loop Header: Depth=1
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	%rax, %r15
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%r14, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_17
# %bb.16:                               #   in Loop: Header=BB2_15 Depth=1
	movq	56(%rsp), %rdi
	movq	48(%rsp), %rsi
	movq	40(%rsp), %rdx
	movl	$13, %ecx
	callq	_Z43__device_stub__gpuParallelMatrixComputationPKiS0_Pii
.LBB2_17:                               #   in Loop: Header=BB2_15 Depth=1
	callq	hipDeviceSynchronize
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	32(%rsp), %rcx                  # 8-byte Reload
	subq	%r15, %rcx
	addq	%rax, %rcx
	movq	%rcx, 32(%rsp)                  # 8-byte Spill
	decw	%bx
	jne	.LBB2_15
# %bb.18:
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.5, %esi
	movl	$24, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	40(%rsp), %rsi
	movl	$676, %edx                      # imm = 0x2A4
	movq	64(%rsp), %rdi                  # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB2_20
# %bb.19:
	movl	%eax, %ebp
	movl	$_ZSt4cout, %edi
	movl	$.L.str.6, %esi
	movl	$22, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%ebp, %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
.LBB2_20:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.7, %esi
	movl	$12, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	xorl	%eax, %eax
	leaq	88(%rsp), %rbx
.LBB2_21:                               # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_22 Depth 2
                                        #       Child Loop BB2_24 Depth 3
	movq	%rax, 72(%rsp)                  # 8-byte Spill
	imulq	$52, %rax, %rax
	addq	64(%rsp), %rax                  # 8-byte Folded Reload
	movq	%rax, 80(%rsp)                  # 8-byte Spill
	xorl	%r15d, %r15d
.LBB2_22:                               #   Parent Loop BB2_21 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB2_24 Depth 3
	movq	80(%rsp), %rax                  # 8-byte Reload
	movl	(%rax,%r15,4), %r12d
	movl	%r12d, %ebp
	negl	%ebp
	cmovsl	%r12d, %ebp
	movl	$1, %r13d
	cmpl	$10, %ebp
	movl	$3518437209, %edx               # imm = 0xD1B71759
	jb	.LBB2_31
# %bb.23:                               # %.lr.ph.i.i.preheader
                                        #   in Loop: Header=BB2_22 Depth=2
	movl	$4, %r13d
	movl	%ebp, %eax
.LBB2_24:                               # %.lr.ph.i.i
                                        #   Parent Loop BB2_21 Depth=1
                                        #     Parent Loop BB2_22 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	cmpl	$99, %eax
	jbe	.LBB2_25
# %bb.26:                               #   in Loop: Header=BB2_24 Depth=3
	cmpl	$999, %eax                      # imm = 0x3E7
	jbe	.LBB2_27
# %bb.28:                               #   in Loop: Header=BB2_24 Depth=3
	cmpl	$10000, %eax                    # imm = 0x2710
	jb	.LBB2_31
# %bb.29:                               #   in Loop: Header=BB2_24 Depth=3
	movl	%eax, %ecx
	imulq	%rdx, %rcx
	shrq	$45, %rcx
	addl	$4, %r13d
	cmpl	$99999, %eax                    # imm = 0x1869F
	movl	%ecx, %eax
	ja	.LBB2_24
# %bb.30:                               # %_ZNSt8__detail14__to_chars_lenIjEEjT_i.exit.i.loopexit
                                        #   in Loop: Header=BB2_22 Depth=2
	addl	$-3, %r13d
	jmp	.LBB2_31
.LBB2_25:                               #   in Loop: Header=BB2_22 Depth=2
	addl	$-2, %r13d
	jmp	.LBB2_31
.LBB2_27:                               #   in Loop: Header=BB2_22 Depth=2
	decl	%r13d
.LBB2_31:                               # %_ZNSt8__detail14__to_chars_lenIjEEjT_i.exit.i
                                        #   in Loop: Header=BB2_22 Depth=2
	shrl	$31, %r12d
	leal	(%r12,%r13), %esi
	leaq	104(%rsp), %rax
	movq	%rax, 88(%rsp)
.Ltmp0:
	movq	%rbx, %rdi
	movl	$45, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructEmc
.Ltmp1:
# %bb.32:                               # %_ZNSt7__cxx119to_stringEi.exit
                                        #   in Loop: Header=BB2_22 Depth=2
	addq	88(%rsp), %r12
	movq	%r12, %rdi
	movl	%r13d, %esi
	movl	%ebp, %edx
	callq	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
.Ltmp3:
	movl	$.L.str.8, %esi
	movq	%rbx, %r14
	movq	%rbx, %rdi
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE6appendEPKc
.Ltmp4:
# %bb.33:                               # %.noexc
                                        #   in Loop: Header=BB2_22 Depth=2
	movq	%rax, %r12
	leaq	16(%rsp), %rdi
	movq	%rdi, (%rsp)
	movq	(%rax), %rsi
	movq	%rax, %r13
	addq	$16, %r13
	cmpq	%r13, %rsi
	je	.LBB2_34
# %bb.35:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
                                        #   in Loop: Header=BB2_22 Depth=2
	movq	%rsi, (%rsp)
	movq	16(%r12), %rax
	movq	%rax, 16(%rsp)
	movq	8(%r12), %rbp
	jmp	.LBB2_36
.LBB2_34:                               #   in Loop: Header=BB2_22 Depth=2
	movq	8(%r12), %rbp
	leaq	1(%rbp), %rdx
	callq	memcpy@PLT
.LBB2_36:                               #   in Loop: Header=BB2_22 Depth=2
	movq	%rbp, 8(%rsp)
	movq	%r13, (%r12)
	movq	$0, 8(%r12)
	movb	$0, 16(%r12)
	movq	(%rsp), %rsi
	movq	8(%rsp), %rdx
.Ltmp6:
	movl	$_ZSt4cout, %edi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp7:
# %bb.37:                               # %_ZStlsIcSt11char_traitsIcESaIcEERSt13basic_ostreamIT_T0_ES7_RKNSt7__cxx1112basic_stringIS4_S5_T1_EE.exit
                                        #   in Loop: Header=BB2_22 Depth=2
	movq	(%rsp), %rdi
	leaq	16(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB2_39
# %bb.38:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i50
                                        #   in Loop: Header=BB2_22 Depth=2
	callq	_ZdlPv
.LBB2_39:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
                                        #   in Loop: Header=BB2_22 Depth=2
	movq	88(%rsp), %rdi
	leaq	104(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB2_41
# %bb.40:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i51
                                        #   in Loop: Header=BB2_22 Depth=2
	callq	_ZdlPv
.LBB2_41:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit53
                                        #   in Loop: Header=BB2_22 Depth=2
	incq	%r15
	cmpq	$13, %r15
	movq	%r14, %rbx
	jne	.LBB2_22
# %bb.42:                               #   in Loop: Header=BB2_21 Depth=1
	movl	$_ZSt4cout, %edi
	movl	$.L.str.9, %esi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rbx
	movq	72(%rsp), %rax                  # 8-byte Reload
	incq	%rax
	cmpq	$13, %rax
	jne	.LBB2_21
# %bb.43:
	movl	$_ZSt4cout, %ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.10, %esi
	movl	$11, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	56(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	movq	40(%rsp), %rdi
	callq	hipFree
	movl	$_ZSt4cout, %edi
	movl	$.L.str.11, %esi
	movl	$18, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	32(%rsp), %rdi                  # 8-byte Reload
	movq	%rdi, %rax
	shrq	$3, %rax
	movabsq	$2361183241434822607, %rcx      # imm = 0x20C49BA5E353F7CF
	mulq	%rcx
	movq	%rdx, %r12
	shrq	$4, %r12
	movl	$1, %esi
	cmpq	$10000, %rdi                    # imm = 0x2710
	leaq	16(%rsp), %rax
	jb	.LBB2_56
# %bb.44:                               # %.lr.ph.i.i60.preheader
	movl	$4, %esi
	movabsq	$3777893186295716171, %rdi      # imm = 0x346DC5D63886594B
	movq	%r12, %rcx
.LBB2_45:                               # %.lr.ph.i.i60
                                        # =>This Inner Loop Header: Depth=1
	cmpq	$99, %rcx
	jbe	.LBB2_46
# %bb.51:                               #   in Loop: Header=BB2_45 Depth=1
	cmpq	$999, %rcx                      # imm = 0x3E7
	jbe	.LBB2_52
# %bb.53:                               #   in Loop: Header=BB2_45 Depth=1
	cmpq	$10000, %rcx                    # imm = 0x2710
	jb	.LBB2_56
# %bb.54:                               #   in Loop: Header=BB2_45 Depth=1
	movq	%rcx, %rax
	mulq	%rdi
	leaq	16(%rsp), %rax
	shrq	$11, %rdx
	addl	$4, %esi
	cmpq	$99999, %rcx                    # imm = 0x1869F
	movq	%rdx, %rcx
	ja	.LBB2_45
# %bb.55:                               # %_ZNSt7__cxx119to_stringEm.exit.loopexit
	addl	$-3, %esi
	jmp	.LBB2_56
.LBB2_46:
	addl	$-2, %esi
	jmp	.LBB2_56
.LBB2_52:
	decl	%esi
.LBB2_56:                               # %_ZNSt7__cxx119to_stringEm.exit
	movl	%esi, %esi
	movq	%rsp, %r15
	movq	%rax, (%r15)
	movq	%r15, %rdi
	xorl	%edx, %edx
	callq	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE12_M_constructEmc
	movq	(%r15), %rdi
	movl	8(%r15), %esi
	movq	%r12, %rdx
	callq	_ZNSt8__detail18__to_chars_10_implImEEvPcjT_
	movq	(%r15), %rsi
	movq	8(%r15), %rdx
.Ltmp9:
	movl	$_ZSt4cout, %edi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp10:
# %bb.57:                               # %_ZStlsIcSt11char_traitsIcESaIcEERSt13basic_ostreamIT_T0_ES7_RKNSt7__cxx1112basic_stringIS4_S5_T1_EE.exit63
.Ltmp11:
	movq	%rax, %rbx
	movl	$.L.str.12, %esi
	movl	$12, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
.Ltmp12:
# %bb.58:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc.exit
	movq	(%rbx), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
.Ltmp13:
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
.Ltmp14:
# %bb.59:                               # %.noexc72
.Ltmp15:
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
.Ltmp16:
# %bb.60:                               # %.noexc73
.Ltmp17:
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
.Ltmp18:
# %bb.61:                               # %_ZNSolsEPFRSoS_E.exit
	movq	(%rsp), %rdi
	leaq	16(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB2_63
# %bb.62:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i66
	callq	_ZdlPv
.LBB2_63:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit68
	xorl	%eax, %eax
	addq	$120, %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
.LBB2_64:
	.cfi_def_cfa_offset 176
.Ltmp19:
	movq	%rax, %r15
	movq	(%rsp), %rdi
	leaq	16(%rsp), %rax
	jmp	.LBB2_65
.LBB2_48:
.Ltmp8:
	movq	%rax, %r15
	movq	(%rsp), %rdi
	leaq	16(%rsp), %rax
	cmpq	%rax, %rdi
	je	.LBB2_50
# %bb.49:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i54
	callq	_ZdlPv
	jmp	.LBB2_50
.LBB2_47:
.Ltmp5:
	movq	%rax, %r15
.LBB2_50:                               # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit56
	movq	88(%rsp), %rdi
	leaq	104(%rsp), %rax
.LBB2_65:
	cmpq	%rax, %rdi
	je	.LBB2_67
# %bb.66:                               # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i69
	callq	_ZdlPv
.LBB2_67:
	movq	%r15, %rdi
	callq	_Unwind_Resume@PLT
.LBB2_68:
.Ltmp2:
	movq	%rax, %rdi
	callq	__clang_call_terminate
.Lfunc_end2:
	.size	main, .Lfunc_end2-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table2:
.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	1                               #   On action: 1
	.uleb128 .Ltmp3-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp4-.Ltmp3                  #   Call between .Ltmp3 and .Ltmp4
	.uleb128 .Ltmp5-.Lfunc_begin0           #     jumps to .Ltmp5
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp4-.Lfunc_begin0           # >> Call Site 4 <<
	.uleb128 .Ltmp6-.Ltmp4                  #   Call between .Ltmp4 and .Ltmp6
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp6-.Lfunc_begin0           # >> Call Site 5 <<
	.uleb128 .Ltmp7-.Ltmp6                  #   Call between .Ltmp6 and .Ltmp7
	.uleb128 .Ltmp8-.Lfunc_begin0           #     jumps to .Ltmp8
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp7-.Lfunc_begin0           # >> Call Site 6 <<
	.uleb128 .Ltmp9-.Ltmp7                  #   Call between .Ltmp7 and .Ltmp9
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp9-.Lfunc_begin0           # >> Call Site 7 <<
	.uleb128 .Ltmp18-.Ltmp9                 #   Call between .Ltmp9 and .Ltmp18
	.uleb128 .Ltmp19-.Lfunc_begin0          #     jumps to .Ltmp19
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp18-.Lfunc_begin0          # >> Call Site 8 <<
	.uleb128 .Lfunc_end2-.Ltmp18            #   Call between .Ltmp18 and .Lfunc_end2
	.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.__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_end3:
	.size	__clang_call_terminate, .Lfunc_end3-__clang_call_terminate
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt8__detail18__to_chars_10_implIjEEvPcjT_,"axG",@progbits,_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_,comdat
	.weak	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_ # -- Begin function _ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
	.type	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_,@function
_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_: # @_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
	.cfi_startproc
# %bb.0:
                                        # kill: def $esi killed $esi def $rsi
	cmpl	$100, %edx
	jb	.LBB4_1
# %bb.2:                                # %.lr.ph.preheader
	decl	%esi
.LBB4_3:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	%edx, %eax
	imulq	$1374389535, %rax, %rax         # imm = 0x51EB851F
	shrq	$37, %rax
	imull	$100, %eax, %ecx
	movl	%edx, %r8d
	subl	%ecx, %r8d
	movb	.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits+1(%r8,%r8), %cl
	movl	%esi, %r9d
	movb	%cl, (%rdi,%r9)
	movb	.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits(%r8,%r8), %cl
	leal	-1(%rsi), %r8d
	movb	%cl, (%rdi,%r8)
	addl	$-2, %esi
	cmpl	$9999, %edx                     # imm = 0x270F
	movl	%eax, %edx
	ja	.LBB4_3
	jmp	.LBB4_4
.LBB4_1:
	movl	%edx, %eax
.LBB4_4:                                # %._crit_edge
	cmpl	$10, %eax
	jb	.LBB4_6
# %bb.5:
	movl	%eax, %eax
	leaq	(%rax,%rax), %rcx
	movl	%ecx, %ecx
	movb	.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits+1(%rcx), %cl
	movb	%cl, 1(%rdi)
	movb	.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits(%rax,%rax), %al
	jmp	.LBB4_7
.LBB4_6:
	orb	$48, %al
.LBB4_7:
	movb	%al, (%rdi)
	retq
.Lfunc_end4:
	.size	_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_, .Lfunc_end4-_ZNSt8__detail18__to_chars_10_implIjEEvPcjT_
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt8__detail18__to_chars_10_implImEEvPcjT_,"axG",@progbits,_ZNSt8__detail18__to_chars_10_implImEEvPcjT_,comdat
	.weak	_ZNSt8__detail18__to_chars_10_implImEEvPcjT_ # -- Begin function _ZNSt8__detail18__to_chars_10_implImEEvPcjT_
	.type	_ZNSt8__detail18__to_chars_10_implImEEvPcjT_,@function
_ZNSt8__detail18__to_chars_10_implImEEvPcjT_: # @_ZNSt8__detail18__to_chars_10_implImEEvPcjT_
	.cfi_startproc
# %bb.0:
	movq	%rdx, %rcx
                                        # kill: def $esi killed $esi def $rsi
	cmpq	$100, %rdx
	jb	.LBB5_1
# %bb.2:                                # %.lr.ph.preheader
	decl	%esi
	movabsq	$2951479051793528259, %r8       # imm = 0x28F5C28F5C28F5C3
.LBB5_3:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	%rcx, %rax
	shrq	$2, %rax
	mulq	%r8
	shrq	$2, %rdx
	imull	$100, %edx, %eax
	movl	%ecx, %r9d
	subl	%eax, %r9d
	movb	.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits+1(%r9,%r9), %al
	movl	%esi, %r10d
	movb	%al, (%rdi,%r10)
	movb	.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits(%r9,%r9), %al
	leal	-1(%rsi), %r9d
	movb	%al, (%rdi,%r9)
	addl	$-2, %esi
	cmpq	$9999, %rcx                     # imm = 0x270F
	movq	%rdx, %rcx
	ja	.LBB5_3
	jmp	.LBB5_4
.LBB5_1:
	movq	%rcx, %rdx
.LBB5_4:                                # %._crit_edge
	cmpq	$10, %rdx
	jb	.LBB5_6
# %bb.5:
	movb	.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits+1(%rdx,%rdx), %al
	movb	%al, 1(%rdi)
	movb	.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits(%rdx,%rdx), %dl
	jmp	.LBB5_7
.LBB5_6:
	orb	$48, %dl
.LBB5_7:
	movb	%dl, (%rdi)
	retq
.Lfunc_end5:
	.size	_ZNSt8__detail18__to_chars_10_implImEEvPcjT_, .Lfunc_end5-_ZNSt8__detail18__to_chars_10_implImEEvPcjT_
	.cfi_endproc
                                        # -- End function
	.text
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z20gpuMatrixComputationPKiS0_Pi, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z28gpuParallelMatrixComputationPKiS0_Pii, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB7_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	matrixCGlobal,@object           # @matrixCGlobal
	.bss
	.globl	matrixCGlobal
	.p2align	4, 0x0
matrixCGlobal:
	.zero	676
	.size	matrixCGlobal, 676

	.type	matrixAGlobal,@object           # @matrixAGlobal
	.data
	.globl	matrixAGlobal
	.p2align	4, 0x0
matrixAGlobal:
	.long	14                              # 0xe
	.long	39                              # 0x27
	.long	117                             # 0x75
	.long	89                              # 0x59
	.long	111                             # 0x6f
	.long	73                              # 0x49
	.long	79                              # 0x4f
	.long	102                             # 0x66
	.long	52                              # 0x34
	.long	81                              # 0x51
	.long	123                             # 0x7b
	.long	70                              # 0x46
	.long	39                              # 0x27
	.long	82                              # 0x52
	.long	29                              # 0x1d
	.long	125                             # 0x7d
	.long	85                              # 0x55
	.long	51                              # 0x33
	.long	60                              # 0x3c
	.long	102                             # 0x66
	.long	39                              # 0x27
	.long	120                             # 0x78
	.long	106                             # 0x6a
	.long	19                              # 0x13
	.long	15                              # 0xf
	.long	58                              # 0x3a
	.long	124                             # 0x7c
	.long	31                              # 0x1f
	.long	32                              # 0x20
	.long	23                              # 0x17
	.long	19                              # 0x13
	.long	69                              # 0x45
	.long	60                              # 0x3c
	.long	61                              # 0x3d
	.long	10                              # 0xa
	.long	33                              # 0x21
	.long	72                              # 0x48
	.long	1                               # 0x1
	.long	91                              # 0x5b
	.long	96                              # 0x60
	.long	112                             # 0x70
	.long	32                              # 0x20
	.long	111                             # 0x6f
	.long	90                              # 0x5a
	.long	12                              # 0xc
	.long	63                              # 0x3f
	.long	77                              # 0x4d
	.long	47                              # 0x2f
	.long	105                             # 0x69
	.long	115                             # 0x73
	.long	38                              # 0x26
	.long	90                              # 0x5a
	.long	13                              # 0xd
	.long	35                              # 0x23
	.long	23                              # 0x17
	.long	78                              # 0x4e
	.long	57                              # 0x39
	.long	109                             # 0x6d
	.long	122                             # 0x7a
	.long	89                              # 0x59
	.long	21                              # 0x15
	.long	116                             # 0x74
	.long	86                              # 0x56
	.long	123                             # 0x7b
	.long	113                             # 0x71
	.long	27                              # 0x1b
	.long	14                              # 0xe
	.long	80                              # 0x50
	.long	69                              # 0x45
	.long	9                               # 0x9
	.long	23                              # 0x17
	.long	106                             # 0x6a
	.long	26                              # 0x1a
	.long	115                             # 0x73
	.long	31                              # 0x1f
	.long	6                               # 0x6
	.long	73                              # 0x49
	.long	112                             # 0x70
	.long	53                              # 0x35
	.long	70                              # 0x46
	.long	64                              # 0x40
	.long	118                             # 0x76
	.long	121                             # 0x79
	.long	17                              # 0x11
	.long	6                               # 0x6
	.long	113                             # 0x71
	.long	30                              # 0x1e
	.long	8                               # 0x8
	.long	5                               # 0x5
	.long	116                             # 0x74
	.long	66                              # 0x42
	.long	12                              # 0xc
	.long	113                             # 0x71
	.long	71                              # 0x47
	.long	94                              # 0x5e
	.long	98                              # 0x62
	.long	116                             # 0x74
	.long	2                               # 0x2
	.long	95                              # 0x5f
	.long	66                              # 0x42
	.long	107                             # 0x6b
	.long	54                              # 0x36
	.long	11                              # 0xb
	.long	34                              # 0x22
	.long	90                              # 0x5a
	.long	36                              # 0x24
	.long	81                              # 0x51
	.long	124                             # 0x7c
	.long	73                              # 0x49
	.long	41                              # 0x29
	.long	105                             # 0x69
	.long	14                              # 0xe
	.long	127                             # 0x7f
	.long	109                             # 0x6d
	.long	87                              # 0x57
	.long	29                              # 0x1d
	.long	2                               # 0x2
	.long	84                              # 0x54
	.long	77                              # 0x4d
	.long	56                              # 0x38
	.long	81                              # 0x51
	.long	21                              # 0x15
	.long	81                              # 0x51
	.long	110                             # 0x6e
	.long	110                             # 0x6e
	.long	123                             # 0x7b
	.long	104                             # 0x68
	.long	113                             # 0x71
	.long	39                              # 0x27
	.long	54                              # 0x36
	.long	75                              # 0x4b
	.long	102                             # 0x66
	.long	44                              # 0x2c
	.long	79                              # 0x4f
	.long	61                              # 0x3d
	.long	55                              # 0x37
	.long	90                              # 0x5a
	.long	125                             # 0x7d
	.long	52                              # 0x34
	.long	45                              # 0x2d
	.long	4                               # 0x4
	.long	120                             # 0x78
	.long	12                              # 0xc
	.long	20                              # 0x14
	.long	20                              # 0x14
	.long	105                             # 0x69
	.long	41                              # 0x29
	.long	20                              # 0x14
	.long	44                              # 0x2c
	.long	108                             # 0x6c
	.long	74                              # 0x4a
	.long	72                              # 0x48
	.long	62                              # 0x3e
	.long	76                              # 0x4c
	.long	34                              # 0x22
	.long	111                             # 0x6f
	.long	38                              # 0x26
	.long	97                              # 0x61
	.long	124                             # 0x7c
	.long	5                               # 0x5
	.long	97                              # 0x61
	.long	87                              # 0x57
	.long	85                              # 0x55
	.long	106                             # 0x6a
	.long	12                              # 0xc
	.long	31                              # 0x1f
	.long	87                              # 0x57
	.long	6                               # 0x6
	.long	77                              # 0x4d
	.size	matrixAGlobal, 676

	.type	matrixBGlobal,@object           # @matrixBGlobal
	.globl	matrixBGlobal
	.p2align	4, 0x0
matrixBGlobal:
	.long	69                              # 0x45
	.long	96                              # 0x60
	.long	71                              # 0x47
	.long	89                              # 0x59
	.long	127                             # 0x7f
	.long	108                             # 0x6c
	.long	96                              # 0x60
	.long	121                             # 0x79
	.long	64                              # 0x40
	.long	65                              # 0x41
	.long	62                              # 0x3e
	.long	91                              # 0x5b
	.long	73                              # 0x49
	.long	9                               # 0x9
	.long	67                              # 0x43
	.long	113                             # 0x71
	.long	48                              # 0x30
	.long	47                              # 0x2f
	.long	53                              # 0x35
	.long	96                              # 0x60
	.long	66                              # 0x42
	.long	7                               # 0x7
	.long	63                              # 0x3f
	.long	17                              # 0x11
	.long	9                               # 0x9
	.long	8                               # 0x8
	.long	107                             # 0x6b
	.long	45                              # 0x2d
	.long	112                             # 0x70
	.long	33                              # 0x21
	.long	114                             # 0x72
	.long	48                              # 0x30
	.long	102                             # 0x66
	.long	70                              # 0x46
	.long	52                              # 0x34
	.long	47                              # 0x2f
	.long	34                              # 0x22
	.long	81                              # 0x51
	.long	17                              # 0x11
	.long	38                              # 0x26
	.long	15                              # 0xf
	.long	61                              # 0x3d
	.long	1                               # 0x1
	.long	104                             # 0x68
	.long	82                              # 0x52
	.long	68                              # 0x44
	.long	53                              # 0x35
	.long	69                              # 0x45
	.long	110                             # 0x6e
	.long	12                              # 0xc
	.long	25                              # 0x19
	.long	46                              # 0x2e
	.long	111                             # 0x6f
	.long	89                              # 0x59
	.long	54                              # 0x36
	.long	0                               # 0x0
	.long	107                             # 0x6b
	.long	81                              # 0x51
	.long	127                             # 0x7f
	.long	124                             # 0x7c
	.long	36                              # 0x24
	.long	17                              # 0x11
	.long	99                              # 0x63
	.long	117                             # 0x75
	.long	75                              # 0x4b
	.long	125                             # 0x7d
	.long	72                              # 0x48
	.long	48                              # 0x30
	.long	67                              # 0x43
	.long	31                              # 0x1f
	.long	104                             # 0x68
	.long	64                              # 0x40
	.long	98                              # 0x62
	.long	94                              # 0x5e
	.long	57                              # 0x39
	.long	81                              # 0x51
	.long	15                              # 0xf
	.long	16                              # 0x10
	.long	111                             # 0x6f
	.long	16                              # 0x10
	.long	127                             # 0x7f
	.long	119                             # 0x77
	.long	88                              # 0x58
	.long	41                              # 0x29
	.long	75                              # 0x4b
	.long	125                             # 0x7d
	.long	22                              # 0x16
	.long	50                              # 0x32
	.long	120                             # 0x78
	.long	6                               # 0x6
	.long	81                              # 0x51
	.long	75                              # 0x4b
	.long	7                               # 0x7
	.long	78                              # 0x4e
	.long	38                              # 0x26
	.long	35                              # 0x23
	.long	115                             # 0x73
	.long	114                             # 0x72
	.long	37                              # 0x25
	.long	66                              # 0x42
	.long	106                             # 0x6a
	.long	64                              # 0x40
	.long	91                              # 0x5b
	.long	97                              # 0x61
	.long	75                              # 0x4b
	.long	102                             # 0x66
	.long	84                              # 0x54
	.long	112                             # 0x70
	.long	65                              # 0x41
	.long	76                              # 0x4c
	.long	87                              # 0x57
	.long	22                              # 0x16
	.long	45                              # 0x2d
	.long	100                             # 0x64
	.long	19                              # 0x13
	.long	18                              # 0x12
	.long	89                              # 0x59
	.long	27                              # 0x1b
	.long	25                              # 0x19
	.long	109                             # 0x6d
	.long	18                              # 0x12
	.long	116                             # 0x74
	.long	19                              # 0x13
	.long	116                             # 0x74
	.long	33                              # 0x21
	.long	103                             # 0x67
	.long	31                              # 0x1f
	.long	29                              # 0x1d
	.long	78                              # 0x4e
	.long	8                               # 0x8
	.long	24                              # 0x18
	.long	12                              # 0xc
	.long	86                              # 0x56
	.long	20                              # 0x14
	.long	32                              # 0x20
	.long	53                              # 0x35
	.long	31                              # 0x1f
	.long	13                              # 0xd
	.long	51                              # 0x33
	.long	36                              # 0x24
	.long	100                             # 0x64
	.long	56                              # 0x38
	.long	44                              # 0x2c
	.long	13                              # 0xd
	.long	8                               # 0x8
	.long	54                              # 0x36
	.long	24                              # 0x18
	.long	101                             # 0x65
	.long	73                              # 0x49
	.long	115                             # 0x73
	.long	120                             # 0x78
	.long	56                              # 0x38
	.long	23                              # 0x17
	.long	63                              # 0x3f
	.long	39                              # 0x27
	.long	93                              # 0x5d
	.long	77                              # 0x4d
	.long	50                              # 0x32
	.long	108                             # 0x6c
	.long	56                              # 0x38
	.long	106                             # 0x6a
	.long	58                              # 0x3a
	.long	121                             # 0x79
	.long	74                              # 0x4a
	.long	70                              # 0x46
	.long	88                              # 0x58
	.long	19                              # 0x13
	.long	49                              # 0x31
	.long	83                              # 0x53
	.size	matrixBGlobal, 676

	.type	matrixB_transposed,@object      # @matrixB_transposed
	.bss
	.globl	matrixB_transposed
	.p2align	4, 0x0
matrixB_transposed:
	.zero	676
	.size	matrixB_transposed, 676

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

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

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"Malloc device mem"
	.size	.L.str, 18

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"malloc failed"
	.size	.L.str.1, 14

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Move to device mem"
	.size	.L.str.2, 19

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"memcpy failed"
	.size	.L.str.3, 14

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

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"Copy from device to host"
	.size	.L.str.5, 25

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"memcpy output failed: "
	.size	.L.str.6, 23

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Printing mem"
	.size	.L.str.7, 13

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"\t"
	.size	.L.str.8, 2

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"\n"
	.size	.L.str.9, 2

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Freeing mem"
	.size	.L.str.10, 12

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"Calculation took: "
	.size	.L.str.11, 19

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	" nanoseconds"
	.size	.L.str.12, 13

	.type	.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits,@object # @__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits
	.section	.rodata.str1.16,"aMS",@progbits,1
	.p2align	4, 0x0
.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits:
	.asciz	"00010203040506070809101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899"
	.size	.L__const._ZNSt8__detail18__to_chars_10_implImEEvPcjT_.__digits, 201

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z20gpuMatrixComputationPKiS0_Pi"
	.size	.L__unnamed_1, 33

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z28gpuParallelMatrixComputationPKiS0_Pii"
	.size	.L__unnamed_2, 42

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.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 _Z35__device_stub__gpuMatrixComputationPKiS0_Pi
	.addrsig_sym _Z43__device_stub__gpuParallelMatrixComputationPKiS0_Pii
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym matrixAGlobal
	.addrsig_sym matrixBGlobal
	.addrsig_sym _Z20gpuMatrixComputationPKiS0_Pi
	.addrsig_sym _Z28gpuParallelMatrixComputationPKiS0_Pii
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z20gpuMatrixComputationPKiS0_Pi ; -- Begin function _Z20gpuMatrixComputationPKiS0_Pi
	.globl	_Z20gpuMatrixComputationPKiS0_Pi
	.p2align	8
	.type	_Z20gpuMatrixComputationPKiS0_Pi,@function
_Z20gpuMatrixComputationPKiS0_Pi:       ; @_Z20gpuMatrixComputationPKiS0_Pi
; %bb.0:
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_mov_b32_e32 v0, 0
	s_mov_b32 s3, 0
	s_mov_b32 s20, 0
	s_waitcnt lgkmcnt(0)
	s_mov_b64 s[8:9], s[4:5]
.LBB0_1:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_2 Depth 2
                                        ;       Child Loop BB0_3 Depth 3
	s_mul_i32 s21, s20, 13
	s_mov_b64 s[10:11], s[6:7]
	s_mov_b32 s22, s3
.LBB0_2:                                ;   Parent Loop BB0_1 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_3 Depth 3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_i32 s2, s22, s21
	v_mov_b32_e32 v1, 0
	s_lshl_b64 s[14:15], s[2:3], 2
	s_mov_b64 s[16:17], 0
	s_add_u32 s12, s0, s14
	s_addc_u32 s13, s1, s15
	s_mov_b64 s[18:19], s[10:11]
	global_store_b32 v0, v0, s[12:13]
.LBB0_3:                                ;   Parent Loop BB0_1 Depth=1
                                        ;     Parent Loop BB0_2 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_add_u32 s24, s8, s16
	s_addc_u32 s25, s9, s17
	s_add_u32 s16, s16, 4
	s_clause 0x1
	global_load_b32 v2, v0, s[24:25]
	global_load_b32 v3, v0, s[18:19]
	s_addc_u32 s17, s17, 0
	s_add_u32 s18, s18, 52
	s_addc_u32 s19, s19, 0
	s_cmp_lg_u32 s16, 52
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v2, v3, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v1, v1, v2
	global_store_b32 v0, v1, s[12:13]
	s_cbranch_scc1 .LBB0_3
; %bb.4:                                ;   in Loop: Header=BB0_2 Depth=2
	s_add_u32 s16, s4, s14
	s_addc_u32 s17, s5, s15
	s_add_u32 s14, s6, s14
	s_addc_u32 s15, s7, s15
	s_clause 0x1
	global_load_b32 v2, v0, s[16:17]
	global_load_b32 v3, v0, s[14:15]
	s_add_i32 s22, s22, 1
	s_add_u32 s10, s10, 4
	s_addc_u32 s11, s11, 0
	s_cmp_lg_u32 s22, 13
	s_waitcnt vmcnt(0)
	v_add_nc_u32_e32 v2, v3, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v1, v2, v1
	global_store_b32 v0, v1, s[12:13]
	s_cbranch_scc1 .LBB0_2
; %bb.5:                                ;   in Loop: Header=BB0_1 Depth=1
	s_add_i32 s20, s20, 1
	s_add_u32 s8, s8, 52
	s_addc_u32 s9, s9, 0
	s_cmp_lg_u32 s20, 13
	s_cbranch_scc1 .LBB0_1
; %bb.6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z20gpuMatrixComputationPKiS0_Pi
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 24
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 4
		.amdhsa_next_free_sgpr 26
		.amdhsa_reserve_vcc 0
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z20gpuMatrixComputationPKiS0_Pi, .Lfunc_end0-_Z20gpuMatrixComputationPKiS0_Pi
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 284
; NumSgprs: 26
; NumVgprs: 4
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 26
; NumVGPRsForWavesPerEU: 4
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z28gpuParallelMatrixComputationPKiS0_Pii ; -- Begin function _Z28gpuParallelMatrixComputationPKiS0_Pii
	.globl	_Z28gpuParallelMatrixComputationPKiS0_Pii
	.p2align	8
	.type	_Z28gpuParallelMatrixComputationPKiS0_Pii,@function
_Z28gpuParallelMatrixComputationPKiS0_Pii: ; @_Z28gpuParallelMatrixComputationPKiS0_Pii
; %bb.0:
	s_clause 0x2
	s_load_b32 s10, s[0:1], 0x18
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_waitcnt lgkmcnt(0)
	s_cmp_lt_i32 s10, 1
	s_mul_i32 s2, s15, s10
	s_cbranch_scc1 .LBB1_3
; %bb.1:                                ; %.lr.ph.preheader
	v_dual_mov_b32 v3, 0 :: v_dual_mov_b32 v2, v0
	s_ashr_i32 s3, s2, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_lshl_b64 s[8:9], s[2:3], 2
	v_mov_b32_e32 v1, v3
	s_add_u32 s8, s4, s8
	s_addc_u32 s9, s5, s9
	s_mov_b32 s3, s10
.LBB1_2:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_lshlrev_b64 v[4:5], 2, v[2:3]
	s_load_b32 s11, s[8:9], 0x0
	s_add_i32 s3, s3, -1
	s_add_u32 s8, s8, 4
	s_addc_u32 s9, s9, 0
	s_cmp_eq_u32 s3, 0
	v_add_co_u32 v4, vcc_lo, s6, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v5, vcc_lo
	global_load_b32 v6, v[4:5], off
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_mad_u64_u32 v[4:5], null, v6, s11, v[1:2]
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v1, v4 :: v_dual_add_nc_u32 v2, s10, v2
	s_cbranch_scc0 .LBB1_2
	s_branch .LBB1_4
.LBB1_3:
	v_mov_b32_e32 v1, 0
.LBB1_4:                                ; %Flow46
	v_add_nc_u32_e32 v2, s2, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v4, vcc_lo, s4, v2
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v3, vcc_lo
	v_add_co_u32 v6, vcc_lo, s6, v2
	v_add_co_ci_u32_e32 v7, vcc_lo, s7, v3, vcc_lo
	global_load_b32 v0, v[4:5], off
	global_load_b32 v4, v[6:7], off
	s_waitcnt vmcnt(0)
	v_add3_u32 v4, v0, v1, v4
	v_add_co_u32 v0, vcc_lo, s0, v2
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v3, vcc_lo
	global_store_b32 v[0:1], v4, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z28gpuParallelMatrixComputationPKiS0_Pii
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 28
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 8
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z28gpuParallelMatrixComputationPKiS0_Pii, .Lfunc_end1-_Z28gpuParallelMatrixComputationPKiS0_Pii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 280
; NumSgprs: 18
; NumVgprs: 8
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 8
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 24
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z20gpuMatrixComputationPKiS0_Pi
    .private_segment_fixed_size: 0
    .sgpr_count:     26
    .sgpr_spill_count: 0
    .symbol:         _Z20gpuMatrixComputationPKiS0_Pi.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     4
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 28
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z28gpuParallelMatrixComputationPKiS0_Pii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z28gpuParallelMatrixComputationPKiS0_Pii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     8
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
