/**
 * TSM2 and ISM2 Testbed and Evaluation Platform
 * by Cody Rivera, 2019-2020
 *
 * Usage - ./multiply [-d] [-i] matrixA matrixB matrixC
 * where -d signifies double precision, and -i signifies
 * ISM2
 */

#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <fstream>
#include <iostream>

#include "cublas_v2.h"
#include "cuda_runtime.h"
#include "cuda_error.cuh"
#include "kernels.cuh"
#include "multiply.cuh"
#include "launch_cublas.cuh"

// Testing Parameters -- Adjust as needed
#define EPS 1e-3
#define N_WARMUP 10
#define N_ROUNDS 100

/**
 * Testbed helper functions.
 */
// Based on
// https://randomascii.wordpress.com/2012/02/25/comparing-floating-point-numbers-2012-edition/
bool approxEqual(double A, double B, double maxRelDiff = EPS) {
    // Calculate the difference.
    double diff = fabs(A - B);
    A = fabs(A);
    B = fabs(B);
    // Find the largest
    double largest = (B > A) ? B : A;

    if (diff <= largest * maxRelDiff) return true;
    return false;
}

template <typename FloatType>
bool matrixCompare(const FloatType* A, const FloatType* B, unsigned int m,
                   unsigned int n, unsigned int& iFail, unsigned int& jFail) {
    FloatType aVal, bVal;
    bool b = true;
    // Cache-friendly comparison pattern
    for (unsigned int j = 0; j < n && b; j++) {
        for (unsigned int i = 0; i < m && b; i++) {
            aVal = A[i + (j * m)];
            bVal = B[i + (j * m)];
            if (!approxEqual(aVal, bVal, EPS)) {
                iFail = i;
                jFail = j;
                b = false;
            }
        }
    }
    return b;
}

template <typename FloatType>
void reportTestSuccess(const char* testName, double GFLOPs,
                       double totalGFLOPs) {
    printf("%s succeeded: %g GFLOPs, %g GFLOPs acc. for transfers\n", testName,
           GFLOPs, totalGFLOPs);
}

template <typename FloatType>
void reportTestFailure(const char* testName, const FloatType* orig,
                       const FloatType* cand, unsigned int leadDim,
                       unsigned int iFail, unsigned int jFail) {
    double oVal = (double)orig[iFail + (jFail * leadDim)];
    double cVal = (double)cand[iFail + (jFail * leadDim)];
    fprintf(stderr,
            "%s failed: Original[%u, %u] = %.6f != Candidate[%u, %u] = %.6f\n",
            testName, iFail, jFail, oVal, iFail, jFail, cVal);
}

template <typename FloatType>
double getGFLOPs(double time, unsigned int m, unsigned int n, unsigned int k) {
    double instCount = ((double)m * (double)n * (double)k) / 1e9;
    double timeSeconds = time / 1000;
    return instCount / timeSeconds;
}

/**
 * Kernel launch wrapper. Runs both CUBLAS and TSM2/ISM2, for evaluation
 * purposes.
 */
template <typename FloatType>
bool runKernels(const FloatType* A, const FloatType* B, FloatType* C,
                const unsigned int m, const unsigned int n,
                const unsigned int k, const bool runIsm2) {
    // Candidate for C -- Used by GPU kernels
    FloatType* candC;
    // Device memory
    FloatType *devA, *devB, *devC;
    // Events used for timing
    cudaEvent_t start, end, startTotal, endTotal;
    float time, timeTotal;

    printf("Multiplying matrix A[%u, %u] by matrix B[%u, %u]\n\n", m, k, k, n);

    // Change test name depending on runIsm2
    const char* testName = "TSM2 Kernel Test";
    if (runIsm2) {
        testName = "ISM2 Kernel Test";
    }

    // Allocates new memory
    candC = (FloatType*)malloc(m * n * sizeof(FloatType));
    if (candC == NULL) {
        fprintf(stderr, "Not enough memory\n");
        return false;
    }

    cudaErrchk(cudaMalloc((FloatType**)&devA, m * k * sizeof(FloatType)));
    cudaErrchk(cudaMalloc((FloatType**)&devB, k * n * sizeof(FloatType)));
    cudaErrchk(cudaMalloc((FloatType**)&devC, m * n * sizeof(FloatType)));

    // Inits CUDA events
    cudaErrchk(cudaEventCreate(&start));
    cudaErrchk(cudaEventCreate(&end));
    cudaErrchk(cudaEventCreate(&startTotal));
    cudaErrchk(cudaEventCreate(&endTotal));

    // Runs CUBLAS call
    cublasHandle_t handle;
    cublasErrchk(cublasCreate(&handle));

    FloatType one = 1;
    FloatType zero = 0;

    cudaErrchk(cudaEventRecord(startTotal));

    // Cuda Memory Copy
    cudaErrchk(
        cudaMemcpy(devA, A, m * k * sizeof(FloatType), cudaMemcpyHostToDevice));
    cudaErrchk(
        cudaMemcpy(devB, B, k * n * sizeof(FloatType), cudaMemcpyHostToDevice));

    for (int i = 0; i < N_WARMUP; ++i) {
        cublasErrchk(launchCublas<FloatType>(handle, one, zero,
            devA, devB, devC, m, n, k));
    }
    
    cudaErrchk(cudaEventRecord(start));
    for (int i = 0; i < N_ROUNDS; ++i) {
        cublasErrchk(launchCublas<FloatType>(handle, one, zero,
                                            devA, devB, devC, m, n, k));
    }
    cudaErrchk(cudaEventRecord(end));

    // Copies result back
    cudaErrchk(
        cudaMemcpy(C, devC, m * n * sizeof(FloatType), cudaMemcpyDeviceToHost));

    cudaErrchk(cudaEventRecord(endTotal));
    cudaErrchk(cudaDeviceSynchronize());
    cudaErrchk(cudaEventElapsedTime(&time, start, end));
    cudaErrchk(cudaEventElapsedTime(&timeTotal, startTotal, endTotal));
    time /= N_ROUNDS;
    timeTotal /= N_ROUNDS;

    reportTestSuccess<FloatType>("CUBLAS Test", getGFLOPs<FloatType>(time, m, n, k),
                                getGFLOPs<FloatType>(timeTotal, m, n, k));

    cublasErrchk(cublasDestroy(handle));

    // Runs kernel
    // Failure flag
    bool status;
    // Failure indices
    unsigned int iFail, jFail;

    // Clear result matrix
    cudaErrchk(cudaMemset(devC, 0, m * n * sizeof(FloatType)));
    cudaErrchk(cudaEventRecord(startTotal));

    // Cuda Memory Copy
    cudaErrchk(
        cudaMemcpy(devA, A, m * k * sizeof(FloatType), cudaMemcpyHostToDevice));
    cudaErrchk(
        cudaMemcpy(devB, B, k * n * sizeof(FloatType), cudaMemcpyHostToDevice));

    for (int i = 0; i < N_WARMUP; ++i) {
        cudaErrchk(cudaMemset(devC, 0, m * n * sizeof(FloatType)));
        if (runIsm2) {
            launchKernelIsm2(devA, devB, devC, m, n, k);
        } else {
            launchKernelTsm2(devA, devB, devC, m, n, k);
        }
    }
    
    cudaErrchk(cudaEventRecord(start));
    for (int i = 0; i < N_ROUNDS; ++i) {
        cudaErrchk(cudaMemset(devC, 0, m * n * sizeof(FloatType)));
        if (runIsm2) {
            launchKernelIsm2(devA, devB, devC, m, n, k);
        } else {
            launchKernelTsm2(devA, devB, devC, m, n, k);
        }
    }
    cudaErrchk(cudaGetLastError());
    cudaErrchk(cudaEventRecord(end));

    // Copies result back
    cudaErrchk(cudaMemcpy(candC, devC, m * n * sizeof(FloatType),
                          cudaMemcpyDeviceToHost));

    cudaErrchk(cudaEventRecord(endTotal));
    cudaErrchk(cudaDeviceSynchronize());
    cudaErrchk(cudaEventElapsedTime(&time, start, end));
    cudaErrchk(cudaEventElapsedTime(&timeTotal, startTotal, endTotal));
    time /= N_ROUNDS;
    timeTotal /= N_ROUNDS;
    
    status = matrixCompare<FloatType>(C, candC, m, n, iFail, jFail);
    if (status) {
        reportTestSuccess<FloatType>(testName,
                                  getGFLOPs<FloatType>(time, m, n, k),
                                  getGFLOPs<FloatType>(timeTotal, m, n, k));
    } else {
        reportTestFailure<FloatType>(testName, C, candC, m, iFail, jFail);
    }

    cudaErrchk(cudaEventDestroy(start));
    cudaErrchk(cudaEventDestroy(end));
    cudaErrchk(cudaEventDestroy(startTotal));
    cudaErrchk(cudaEventDestroy(endTotal));
    free(candC);
    cudaErrchk(cudaFree(devA));
    cudaErrchk(cudaFree(devB));
    cudaErrchk(cudaFree(devC));

    return true;
}

/**
 * Runs testbed on specified input files. Handles file IO.
 */
template <typename FloatType>
bool runMatmul(std::istream& fileA, std::istream& fileB, std::ostream& outFile,
               bool runIsm2) {
    FloatType *A, *B, *C;
    int m, n, k, kCand;

    // Reads Matrix Sizes
    fileA.read((char*)&m, sizeof(unsigned int));
    fileA.read((char*)&k, sizeof(unsigned int));
    fileB.read((char*)&kCand, sizeof(unsigned int));
    fileB.read((char*)&n, sizeof(unsigned int));

    if (k != kCand) {
        fprintf(stderr,
                "Matrix multiplication is undefined where A's"
                "column count is not equal\n to B's row count\n\n"
                "Matrix A (%u x %u) and Matrix B (%u x %u)\n",
                m, k, kCand, n);
        return false;
    }

    // Mallocs Matrices on CPU
    A = (FloatType*)malloc((size_t)m * k * sizeof(FloatType));
    B = (FloatType*)malloc((size_t)k * n * sizeof(FloatType));
    C = (FloatType*)malloc((size_t)m * n * sizeof(FloatType));

    if (A == NULL || B == NULL || C == NULL) {
        fprintf(stderr, "Not enough memory\n");
        return false;
    }

    // Loads Data to Matrix A and B
    fileA.read((char*)A, (size_t)m * k * sizeof(FloatType));
    fileB.read((char*)B, (size_t)k * n * sizeof(FloatType));

    // Calls CUDA
    bool status = runKernels<FloatType>(A, B, C, m, n, k, runIsm2);
    if (!status) {
        free(A);
        free(B);
        free(C);
        return false;
    }

    // Writes output matrix
    outFile.write((const char*)&m, sizeof(unsigned int));
    outFile.write((const char*)&n, sizeof(unsigned int));
    outFile.write((const char*)C, (size_t)m * n * sizeof(FloatType));

    free(A);
    free(B);
    free(C);
    return true;
}

/**
 * Entry point
 */
int main(int argc, char** argv) {
    int fileArg[3];
    int nFiles = 0;
    bool isDouble = false;
    bool runIsm2 = false;
    for (int i = 1; i < argc; i++) {
        if (strcmp(argv[i], "-d") == 0) {
            isDouble = true;
        } else if (strcmp(argv[i], "-i") == 0) {
            runIsm2 = true;
        } else {
            if (nFiles < 3) {
                fileArg[nFiles] = i;
            }
            nFiles++;
        }
    }
    if (nFiles != 3) {
        fprintf(stderr, "Usage: %s [-d] [-i] matrixA matrixB matrixC\n",
                argv[0]);
        return 1;
    }

    std::ifstream fileA(argv[fileArg[0]], std::ios::binary),
        fileB(argv[fileArg[1]], std::ios::binary);
    std::ofstream outFile(argv[fileArg[2]], std::ios::binary);
    if (!fileA) {
        fprintf(stderr, "Cannot open %s for reading\n", argv[fileArg[0]]);
        return 1;
    }
    if (!fileB) {
        fprintf(stderr, "Cannot open %s for reading\n", argv[fileArg[1]]);
        return 1;
    }
    if (!outFile) {
        fprintf(stderr, "Cannot open %s for writing\n", argv[fileArg[2]]);
        return 1;
    }
    // Runs matmul
    bool status = false;
    if (isDouble) {
        status = runMatmul<double>(fileA, fileB, outFile, runIsm2);
    } else {
        status = runMatmul<float>(fileA, fileB, outFile, runIsm2);
    }
    fileA.close();
    fileB.close();
    outFile.close();
    if (status) {
        return 0;
    } else {
        return 1;
    }
}

	.file	"tmpxft_002a7585_00000000-6_multiply.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4583:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4583:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata._Z10cudaAssert9cudaErrorPKcib.str1.1,"aMS",@progbits,1
.LC0:
	.string	"CUDAassert: %s %s %d\n"
	.section	.text._Z10cudaAssert9cudaErrorPKcib,"axG",@progbits,_Z10cudaAssert9cudaErrorPKcib,comdat
	.weak	_Z10cudaAssert9cudaErrorPKcib
	.type	_Z10cudaAssert9cudaErrorPKcib, @function
_Z10cudaAssert9cudaErrorPKcib:
.LFB4571:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L8
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%ecx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movl	%edx, 12(%rsp)
	movq	%rsi, (%rsp)
	call	cudaGetErrorString@PLT
	movl	12(%rsp), %r9d
	movq	(%rsp), %r8
	leaq	.LC0(%rip), %rdx
	movq	stderr(%rip), %rdi
	movq	%rax, %rcx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	testb	%bpl, %bpl
	je	.L2
	movl	%ebx, %edi
	call	exit@PLT
.L2:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
.L8:
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE4571:
	.size	_Z10cudaAssert9cudaErrorPKcib, .-_Z10cudaAssert9cudaErrorPKcib
	.section	.rodata._Z12cublasAssert14cublasStatus_tPKcib.str1.1,"aMS",@progbits,1
.LC1:
	.string	"CUBLASassert Failure Code: %d %s %d\n"
	.section	.text._Z12cublasAssert14cublasStatus_tPKcib,"axG",@progbits,_Z12cublasAssert14cublasStatus_tPKcib,comdat
	.weak	_Z12cublasAssert14cublasStatus_tPKcib
	.type	_Z12cublasAssert14cublasStatus_tPKcib, @function
_Z12cublasAssert14cublasStatus_tPKcib:
.LFB4572:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L18
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%ecx, %ebp
	movl	%edi, %ecx
	movl	%edx, %r9d
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	xorl	%eax, %eax
	leaq	.LC1(%rip), %rdx
	pushq	%r8
	.cfi_def_cfa_offset 32
	movq	stderr(%rip), %rdi
	movq	%rsi, %r8
	movl	$2, %esi
	call	__fprintf_chk@PLT
	testb	%bpl, %bpl
	je	.L12
	movl	%ebx, %edi
	call	exit@PLT
.L12:
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
.L18:
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE4572:
	.size	_Z12cublasAssert14cublasStatus_tPKcib, .-_Z12cublasAssert14cublasStatus_tPKcib
	.text
	.globl	_Z11approxEqualddd
	.type	_Z11approxEqualddd, @function
_Z11approxEqualddd:
.LFB4573:
	.cfi_startproc
	endbr64
	movq	.LC2(%rip), %xmm4
	movaps	%xmm0, %xmm3
	subsd	%xmm1, %xmm3
	andps	%xmm4, %xmm0
	andps	%xmm4, %xmm1
	maxsd	%xmm0, %xmm1
	andps	%xmm4, %xmm3
	movaps	%xmm1, %xmm0
	mulsd	%xmm2, %xmm0
	comisd	%xmm3, %xmm0
	setnb	%al
	ret
	.cfi_endproc
.LFE4573:
	.size	_Z11approxEqualddd, .-_Z11approxEqualddd
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4606:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE4606:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata._Z10runKernelsIdEbPKT_S2_PS0_jjjb.str1.1,"aMS",@progbits,1
.LC3:
	.string	"TSM2 Kernel Test"
.LC4:
	.string	"ISM2 Kernel Test"
.LC5:
	.string	"Multiplying matrix A[%u, %u] by matrix B[%u, %u]\n\n"
.LC6:
	.string	"Not enough memory\n"
.LC7:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/EddieBurning/tsm2x-imp/master/src/multiply.cu"
.LC13:
	.string	"CUBLAS Test"
.LC14:
	.string	"%s succeeded: %g GFLOPs, %g GFLOPs acc. for transfers\n"
.LC16:
	.string	"%s failed: Original[%u, %u] = %.6f != Candidate[%u, %u] = %.6f\n"
	.section	.text._Z10runKernelsIdEbPKT_S2_PS0_jjjb,"axG",@progbits,_Z10runKernelsIdEbPKT_S2_PS0_jjjb,comdat
	.weak	_Z10runKernelsIdEbPKT_S2_PS0_jjjb
	.type	_Z10runKernelsIdEbPKT_S2_PS0_jjjb, @function
_Z10runKernelsIdEbPKT_S2_PS0_jjjb:
.LFB5110:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%r9d, %r12d
	movl	%r8d, %r9d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%r8d, %ebp
	movl	%r12d, %r8d
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%ecx, %ebx
	movl	%r12d, %ecx
	subq	$184, %rsp
	.cfi_def_cfa_offset 240
	movl	240(%rsp), %r13d
	movq	%rdi, 56(%rsp)
	movl	$2, %edi
	movq	%rsi, 64(%rsp)
	leaq	.LC5(%rip), %rsi
	movq	%rdx, 40(%rsp)
	movl	%ebx, %edx
	movb	%r13b, 79(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	testb	%r13b, %r13b
	leaq	.LC4(%rip), %rdx
	leaq	.LC3(%rip), %rax
	cmovne	%rdx, %rax
	movq	%rax, 48(%rsp)
	movl	%ebx, %eax
	imull	%ebp, %eax
	salq	$3, %rax
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	malloc@PLT
	testq	%rax, %rax
	jne	.L29
	movq	stderr(%rip), %rdi
	leaq	.LC6(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	xorl	%eax, %eax
	jmp	.L27
.L29:
	movq	%rax, %r14
	movl	%ebx, %eax
	leaq	88(%rsp), %rdi
	movl	$10, %r13d
	imull	%r12d, %eax
	salq	$3, %rax
	movq	%rax, %rsi
	movq	%rax, 24(%rsp)
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$121, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movl	%r12d, %eax
	leaq	96(%rsp), %rdi
	imull	%ebp, %eax
	salq	$3, %rax
	movq	%rax, %rsi
	movq	%rax, 32(%rsp)
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$122, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	8(%rsp), %rsi
	leaq	104(%rsp), %rdi
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$123, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	leaq	112(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$1, %ecx
	movl	$126, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	leaq	120(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$1, %ecx
	movl	$127, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	leaq	128(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$1, %ecx
	movl	$128, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	leaq	136(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$1, %ecx
	movl	$129, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	leaq	144(%rsp), %rdi
	call	cublasCreate_v2@PLT
	movl	$1, %ecx
	movl	$133, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z12cublasAssert14cublasStatus_tPKcib
	movq	.LC8(%rip), %rax
	movq	128(%rsp), %rdi
	xorl	%esi, %esi
	movq	$0x000000000, 160(%rsp)
	movq	%rax, 152(%rsp)
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$138, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	24(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	$1, %ecx
	movq	88(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$141, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	32(%rsp), %rdx
	movq	64(%rsp), %rsi
	movl	$1, %ecx
	movq	96(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$143, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
.L31:
	leaq	160(%rsp), %rdx
	leaq	152(%rsp), %rsi
	pushq	%rdi
	.cfi_def_cfa_offset 248
	pushq	%r12
	.cfi_def_cfa_offset 256
	pushq	%rbp
	.cfi_def_cfa_offset 264
	pushq	%rbx
	.cfi_def_cfa_offset 272
	movq	120(%rsp), %rcx
	movq	176(%rsp), %rdi
	movq	136(%rsp), %r9
	movq	128(%rsp), %r8
	call	_Z12launchCublasIdE14cublasStatus_tP13cublasContextRT_S4_PKS3_S6_PS3_jjj@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 240
	movl	$1, %ecx
	movl	$147, %edx
	movl	%eax, %edi
	leaq	.LC7(%rip), %rsi
	call	_Z12cublasAssert14cublasStatus_tPKcib
	decl	%r13d
	jne	.L31
	movq	112(%rsp), %rdi
	xorl	%esi, %esi
	movl	$100, %r13d
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$151, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
.L32:
	leaq	160(%rsp), %rdx
	leaq	152(%rsp), %rsi
	pushq	%rcx
	.cfi_def_cfa_offset 248
	pushq	%r12
	.cfi_def_cfa_offset 256
	pushq	%rbp
	.cfi_def_cfa_offset 264
	pushq	%rbx
	.cfi_def_cfa_offset 272
	movq	120(%rsp), %rcx
	movq	176(%rsp), %rdi
	movq	136(%rsp), %r9
	movq	128(%rsp), %r8
	call	_Z12launchCublasIdE14cublasStatus_tP13cublasContextRT_S4_PKS3_S6_PS3_jjj@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 240
	movl	$1, %ecx
	movl	$153, %edx
	movl	%eax, %edi
	leaq	.LC7(%rip), %rsi
	call	_Z12cublasAssert14cublasStatus_tPKcib
	decl	%r13d
	jne	.L32
	movq	120(%rsp), %rdi
	xorl	%esi, %esi
	movl	$10, %r13d
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$156, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	8(%rsp), %rdx
	movq	104(%rsp), %rsi
	movl	$2, %ecx
	movq	40(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$159, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	136(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$162, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	call	cudaDeviceSynchronize@PLT
	movl	$1, %ecx
	movl	$163, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	120(%rsp), %rdx
	movq	112(%rsp), %rsi
	leaq	80(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movl	$1, %ecx
	movl	$164, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	136(%rsp), %rdx
	leaq	84(%rsp), %rdi
	movq	128(%rsp), %rsi
	call	cudaEventElapsedTime@PLT
	movl	$1, %ecx
	movl	$165, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movl	%ebx, %eax
	movss	80(%rsp), %xmm1
	divss	.LC10(%rip), %xmm1
	cvtsi2sdq	%rax, %xmm2
	movl	%ebp, %eax
	movss	%xmm1, 80(%rsp)
	cvtss2sd	%xmm1, %xmm1
	cvtsi2sdq	%rax, %xmm3
	movl	%r12d, %eax
	divsd	.LC12(%rip), %xmm1
	movss	84(%rsp), %xmm0
	divss	.LC10(%rip), %xmm0
	movss	%xmm0, 84(%rsp)
	cvtss2sd	%xmm0, %xmm0
	divsd	.LC12(%rip), %xmm0
	leaq	.LC13(%rip), %rdx
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	mulsd	%xmm3, %xmm2
	cvtsi2sdq	%rax, %xmm3
	movb	$2, %al
	mulsd	%xmm3, %xmm2
	divsd	.LC11(%rip), %xmm2
	movsd	%xmm2, 16(%rsp)
	divsd	%xmm1, %xmm2
	movsd	16(%rsp), %xmm1
	divsd	%xmm0, %xmm1
	movaps	%xmm2, %xmm0
	call	__printf_chk@PLT
	movq	144(%rsp), %rdi
	call	cublasDestroy_v2@PLT
	movl	$1, %ecx
	movl	$172, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z12cublasAssert14cublasStatus_tPKcib
	movq	8(%rsp), %rdx
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movl	$1, %ecx
	movl	$181, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	128(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$182, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	24(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	$1, %ecx
	movq	88(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$185, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	32(%rsp), %rdx
	movq	64(%rsp), %rsi
	movl	$1, %ecx
	movq	96(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$187, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
.L41:
	movq	8(%rsp), %rdx
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movl	$1, %ecx
	movl	$191, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	cmpb	$0, 79(%rsp)
	movl	%r12d, %r9d
	movl	%ebp, %r8d
	movq	88(%rsp), %rdi
	movq	96(%rsp), %rsi
	movl	%ebx, %ecx
	movq	104(%rsp), %rdx
	je	.L39
	call	_Z16launchKernelIsm2IdEvPKT_S2_PS0_jjj@PLT
	jmp	.L40
.L39:
	call	_Z16launchKernelTsm2IdEvPKT_S2_PS0_jjj@PLT
.L40:
	decl	%r13d
	jne	.L41
	movq	112(%rsp), %rdi
	xorl	%esi, %esi
	movl	$100, %r13d
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$199, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
.L44:
	movq	8(%rsp), %rdx
	movq	104(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movl	$1, %ecx
	movl	$201, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	cmpb	$0, 79(%rsp)
	movl	%r12d, %r9d
	movl	%ebp, %r8d
	movq	88(%rsp), %rdi
	movq	96(%rsp), %rsi
	movl	%ebx, %ecx
	movq	104(%rsp), %rdx
	je	.L42
	call	_Z16launchKernelIsm2IdEvPKT_S2_PS0_jjj@PLT
	jmp	.L43
.L42:
	call	_Z16launchKernelTsm2IdEvPKT_S2_PS0_jjj@PLT
.L43:
	decl	%r13d
	jne	.L44
	call	cudaGetLastError@PLT
	movl	$1, %ecx
	movl	$208, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	120(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$209, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	8(%rsp), %rdx
	movq	104(%rsp), %rsi
	movq	%r14, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$212, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	136(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$215, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	call	cudaDeviceSynchronize@PLT
	movl	$1, %ecx
	movl	$216, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	120(%rsp), %rdx
	movq	112(%rsp), %rsi
	leaq	80(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movl	$1, %ecx
	movl	$217, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	136(%rsp), %rdx
	leaq	84(%rsp), %rdi
	movq	128(%rsp), %rsi
	call	cudaEventElapsedTime@PLT
	movl	$1, %ecx
	movl	$218, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movss	80(%rsp), %xmm6
	xorl	%esi, %esi
	xorl	%ecx, %ecx
	movss	84(%rsp), %xmm5
	divss	.LC10(%rip), %xmm6
	movb	$1, %dl
	divss	.LC10(%rip), %xmm5
	movss	%xmm6, 80(%rsp)
	movss	%xmm5, 84(%rsp)
.L45:
	cmpl	%ebp, %ecx
	setb	%al
	andb	%dl, %al
	je	.L65
	xorl	%edx, %edx
.L49:
	cmpl	%ebx, %edx
	jnb	.L56
	testb	%al, %al
	je	.L56
	movq	40(%rsp), %rdi
	leal	(%rdx,%rsi), %eax
	movsd	.LC15(%rip), %xmm2
	movsd	(%r14,%rax,8), %xmm1
	movsd	(%rdi,%rax,8), %xmm0
	call	_Z11approxEqualddd
	movl	72(%rsp), %edi
	testb	%al, %al
	cmove	%ecx, %edi
	cmove	%edx, %r15d
	incl	%edx
	movl	%edi, 72(%rsp)
	jmp	.L49
.L56:
	incl	%ecx
	addl	%ebx, %esi
	movl	%eax, %edx
	jmp	.L45
.L65:
	testb	%dl, %dl
	je	.L50
	movsd	16(%rsp), %xmm0
	movsd	16(%rsp), %xmm1
	cvtss2sd	%xmm5, %xmm5
	movb	$2, %al
	cvtss2sd	%xmm6, %xmm6
	divsd	.LC12(%rip), %xmm5
	divsd	.LC12(%rip), %xmm6
	movq	48(%rsp), %rdx
	divsd	%xmm6, %xmm0
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	divsd	%xmm5, %xmm1
	call	__printf_chk@PLT
	jmp	.L51
.L50:
	movl	72(%rsp), %eax
	movq	stderr(%rip), %rdi
	movl	%r15d, %r8d
	movl	$2, %esi
	leaq	.LC16(%rip), %rdx
	imull	%eax, %ebx
	leal	(%rbx,%r15), %eax
	movq	40(%rsp), %rbx
	movsd	(%r14,%rax,8), %xmm1
	movsd	(%rbx,%rax,8), %xmm0
	movl	72(%rsp), %ebx
	movb	$2, %al
	pushq	%rbx
	.cfi_def_cfa_offset 248
	pushq	%r15
	.cfi_def_cfa_offset 256
	movl	88(%rsp), %r9d
	movq	64(%rsp), %rcx
	call	__fprintf_chk@PLT
	popq	%rax
	.cfi_def_cfa_offset 248
	popq	%rdx
	.cfi_def_cfa_offset 240
.L51:
	movq	112(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$1, %ecx
	movl	$231, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	120(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$1, %ecx
	movl	$232, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	128(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$1, %ecx
	movl	$233, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	136(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$1, %ecx
	movl	$234, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	%r14, %rdi
	call	free@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$236, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$237, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$238, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movb	$1, %al
.L27:
	movq	168(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L52
	call	__stack_chk_fail@PLT
.L52:
	addq	$184, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5110:
	.size	_Z10runKernelsIdEbPKT_S2_PS0_jjjb, .-_Z10runKernelsIdEbPKT_S2_PS0_jjjb
	.section	.rodata._Z9runMatmulIdEbRSiS0_RSob.str1.1,"aMS",@progbits,1
.LC17:
	.string	"Matrix multiplication is undefined where A'scolumn count is not equal\n to B's row count\n\nMatrix A (%u x %u) and Matrix B (%u x %u)\n"
	.section	.text._Z9runMatmulIdEbRSiS0_RSob,"axG",@progbits,_Z9runMatmulIdEbRSiS0_RSob,comdat
	.weak	_Z9runMatmulIdEbRSiS0_RSob
	.type	_Z9runMatmulIdEbRSiS0_RSob, @function
_Z9runMatmulIdEbRSiS0_RSob:
.LFB4937:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%ecx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	%rsi, (%rsp)
	movq	%rdx, 8(%rsp)
	movl	$4, %edx
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movq	%rax, %rsi
	movq	%rax, 16(%rsp)
	call	_ZNSi4readEPcl@PLT
	leaq	64(%rsp), %rsi
	movl	$4, %edx
	movq	%r15, %rdi
	call	_ZNSi4readEPcl@PLT
	movq	(%rsp), %rdi
	leaq	68(%rsp), %rsi
	movl	$4, %edx
	call	_ZNSi4readEPcl@PLT
	leaq	60(%rsp), %rax
	movq	(%rsp), %rdi
	movl	$4, %edx
	movq	%rax, %rsi
	movq	%rax, 24(%rsp)
	call	_ZNSi4readEPcl@PLT
	movl	64(%rsp), %r8d
	movl	68(%rsp), %r9d
	movslq	56(%rsp), %rcx
	cmpl	%r9d, %r8d
	je	.L67
	pushq	%rsi
	.cfi_def_cfa_offset 152
	movl	68(%rsp), %eax
	leaq	.LC17(%rip), %rdx
	movl	$2, %esi
	movq	stderr(%rip), %rdi
	pushq	%rax
	.cfi_def_cfa_offset 160
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	popq	%rdi
	.cfi_def_cfa_offset 152
	popq	%r8
	.cfi_def_cfa_offset 144
	jmp	.L68
.L67:
	movslq	%r8d, %rbx
	movq	%rcx, %rdx
	movq	%rcx, 40(%rsp)
	imulq	%rbx, %rdx
	salq	$3, %rdx
	movq	%rdx, %rdi
	movq	%rdx, 32(%rsp)
	call	malloc@PLT
	movslq	60(%rsp), %r14
	movq	%rax, %r12
	imulq	%r14, %rbx
	leaq	0(,%rbx,8), %rdi
	call	malloc@PLT
	movq	40(%rsp), %rcx
	movq	%rax, %rbp
	imulq	%r14, %rcx
	leaq	0(,%rcx,8), %rdi
	call	malloc@PLT
	testq	%r12, %r12
	movq	%rax, %rbx
	sete	%al
	testq	%rbp, %rbp
	sete	%cl
	orb	%cl, %al
	jne	.L74
	testq	%rbx, %rbx
	movq	32(%rsp), %rdx
	jne	.L69
.L74:
	movq	stderr(%rip), %rdi
	leaq	.LC6(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L68:
	xorl	%r13d, %r13d
	jmp	.L66
.L69:
	movq	%r12, %rsi
	movq	%r15, %rdi
	movzbl	%r13b, %r13d
	call	_ZNSi4readEPcl@PLT
	movslq	64(%rsp), %rdx
	movslq	60(%rsp), %rax
	movq	%rbp, %rsi
	movq	(%rsp), %rdi
	imulq	%rax, %rdx
	salq	$3, %rdx
	call	_ZNSi4readEPcl@PLT
	movq	%rbx, %rdx
	movq	%rbp, %rsi
	movq	%r12, %rdi
	pushq	%rax
	.cfi_def_cfa_offset 152
	pushq	%r13
	.cfi_def_cfa_offset 160
	movl	80(%rsp), %r9d
	movl	76(%rsp), %r8d
	movl	72(%rsp), %ecx
	call	_Z10runKernelsIdEbPKT_S2_PS0_jjjb
	popq	%rdx
	.cfi_def_cfa_offset 152
	popq	%rcx
	.cfi_def_cfa_offset 144
	movl	%eax, %r13d
	testb	%al, %al
	jne	.L72
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	jmp	.L68
.L72:
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$4, %edx
	call	_ZNSo5writeEPKcl@PLT
	movq	24(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$4, %edx
	call	_ZNSo5writeEPKcl@PLT
	movslq	56(%rsp), %rdx
	movslq	60(%rsp), %rax
	movq	%rbx, %rsi
	movq	8(%rsp), %rdi
	imulq	%rax, %rdx
	salq	$3, %rdx
	call	_ZNSo5writeEPKcl@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
.L66:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L73
	call	__stack_chk_fail@PLT
.L73:
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	movl	%r13d, %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
.LFE4937:
	.size	_Z9runMatmulIdEbRSiS0_RSob, .-_Z9runMatmulIdEbRSiS0_RSob
	.section	.text._Z10runKernelsIfEbPKT_S2_PS0_jjjb,"axG",@progbits,_Z10runKernelsIfEbPKT_S2_PS0_jjjb,comdat
	.weak	_Z10runKernelsIfEbPKT_S2_PS0_jjjb
	.type	_Z10runKernelsIfEbPKT_S2_PS0_jjjb, @function
_Z10runKernelsIfEbPKT_S2_PS0_jjjb:
.LFB5111:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%r9d, %r12d
	movl	%r8d, %r9d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%r8d, %ebp
	movl	%r12d, %r8d
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%ecx, %ebx
	movl	%r12d, %ecx
	subq	$184, %rsp
	.cfi_def_cfa_offset 240
	movl	240(%rsp), %r13d
	movq	%rdi, 56(%rsp)
	movl	$2, %edi
	movq	%rsi, 64(%rsp)
	leaq	.LC5(%rip), %rsi
	movq	%rdx, 40(%rsp)
	movl	%ebx, %edx
	movb	%r13b, 79(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	testb	%r13b, %r13b
	leaq	.LC4(%rip), %rdx
	leaq	.LC3(%rip), %rax
	cmovne	%rdx, %rax
	movq	%rax, 48(%rsp)
	movl	%ebx, %eax
	imull	%ebp, %eax
	salq	$2, %rax
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	call	malloc@PLT
	testq	%rax, %rax
	jne	.L78
	movq	stderr(%rip), %rdi
	leaq	.LC6(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	xorl	%eax, %eax
	jmp	.L76
.L78:
	movq	%rax, %r14
	movl	%ebx, %eax
	leaq	104(%rsp), %rdi
	movl	$10, %r13d
	imull	%r12d, %eax
	salq	$2, %rax
	movq	%rax, %rsi
	movq	%rax, 24(%rsp)
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$121, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movl	%r12d, %eax
	leaq	112(%rsp), %rdi
	imull	%ebp, %eax
	salq	$2, %rax
	movq	%rax, %rsi
	movq	%rax, 32(%rsp)
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$122, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	8(%rsp), %rsi
	leaq	120(%rsp), %rdi
	call	cudaMalloc@PLT
	movl	$1, %ecx
	movl	$123, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	leaq	128(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$1, %ecx
	movl	$126, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	leaq	136(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$1, %ecx
	movl	$127, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	leaq	144(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$1, %ecx
	movl	$128, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	leaq	152(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$1, %ecx
	movl	$129, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	leaq	160(%rsp), %rdi
	call	cublasCreate_v2@PLT
	movl	$1, %ecx
	movl	$133, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z12cublasAssert14cublasStatus_tPKcib
	movq	144(%rsp), %rdi
	xorl	%esi, %esi
	movl	$0x3f800000, 96(%rsp)
	movl	$0x00000000, 100(%rsp)
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$138, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	24(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	$1, %ecx
	movq	104(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$141, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	32(%rsp), %rdx
	movq	64(%rsp), %rsi
	movl	$1, %ecx
	movq	112(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$143, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
.L80:
	leaq	100(%rsp), %rdx
	leaq	96(%rsp), %rsi
	pushq	%rdi
	.cfi_def_cfa_offset 248
	pushq	%r12
	.cfi_def_cfa_offset 256
	pushq	%rbp
	.cfi_def_cfa_offset 264
	pushq	%rbx
	.cfi_def_cfa_offset 272
	movq	136(%rsp), %rcx
	movq	192(%rsp), %rdi
	movq	152(%rsp), %r9
	movq	144(%rsp), %r8
	call	_Z12launchCublasIfE14cublasStatus_tP13cublasContextRT_S4_PKS3_S6_PS3_jjj@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 240
	movl	$1, %ecx
	movl	$147, %edx
	movl	%eax, %edi
	leaq	.LC7(%rip), %rsi
	call	_Z12cublasAssert14cublasStatus_tPKcib
	decl	%r13d
	jne	.L80
	movq	128(%rsp), %rdi
	xorl	%esi, %esi
	movl	$100, %r13d
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$151, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
.L81:
	leaq	100(%rsp), %rdx
	leaq	96(%rsp), %rsi
	pushq	%rcx
	.cfi_def_cfa_offset 248
	pushq	%r12
	.cfi_def_cfa_offset 256
	pushq	%rbp
	.cfi_def_cfa_offset 264
	pushq	%rbx
	.cfi_def_cfa_offset 272
	movq	136(%rsp), %rcx
	movq	192(%rsp), %rdi
	movq	152(%rsp), %r9
	movq	144(%rsp), %r8
	call	_Z12launchCublasIfE14cublasStatus_tP13cublasContextRT_S4_PKS3_S6_PS3_jjj@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 240
	movl	$1, %ecx
	movl	$153, %edx
	movl	%eax, %edi
	leaq	.LC7(%rip), %rsi
	call	_Z12cublasAssert14cublasStatus_tPKcib
	decl	%r13d
	jne	.L81
	movq	136(%rsp), %rdi
	xorl	%esi, %esi
	movl	$10, %r13d
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$156, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	8(%rsp), %rdx
	movq	120(%rsp), %rsi
	movl	$2, %ecx
	movq	40(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$159, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	152(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$162, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	call	cudaDeviceSynchronize@PLT
	movl	$1, %ecx
	movl	$163, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	136(%rsp), %rdx
	leaq	88(%rsp), %rdi
	movq	128(%rsp), %rsi
	call	cudaEventElapsedTime@PLT
	movl	$1, %ecx
	movl	$164, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	152(%rsp), %rdx
	leaq	92(%rsp), %rdi
	movq	144(%rsp), %rsi
	call	cudaEventElapsedTime@PLT
	movl	$1, %ecx
	movl	$165, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movl	%ebx, %eax
	movss	88(%rsp), %xmm1
	divss	.LC10(%rip), %xmm1
	cvtsi2sdq	%rax, %xmm2
	movl	%ebp, %eax
	movss	%xmm1, 88(%rsp)
	cvtss2sd	%xmm1, %xmm1
	cvtsi2sdq	%rax, %xmm3
	movl	%r12d, %eax
	divsd	.LC12(%rip), %xmm1
	movss	92(%rsp), %xmm0
	divss	.LC10(%rip), %xmm0
	movss	%xmm0, 92(%rsp)
	cvtss2sd	%xmm0, %xmm0
	divsd	.LC12(%rip), %xmm0
	leaq	.LC13(%rip), %rdx
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	mulsd	%xmm3, %xmm2
	cvtsi2sdq	%rax, %xmm3
	movb	$2, %al
	mulsd	%xmm3, %xmm2
	divsd	.LC11(%rip), %xmm2
	movsd	%xmm2, 16(%rsp)
	divsd	%xmm1, %xmm2
	movsd	16(%rsp), %xmm1
	divsd	%xmm0, %xmm1
	movaps	%xmm2, %xmm0
	call	__printf_chk@PLT
	movq	160(%rsp), %rdi
	call	cublasDestroy_v2@PLT
	movl	$1, %ecx
	movl	$172, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z12cublasAssert14cublasStatus_tPKcib
	movq	8(%rsp), %rdx
	movq	120(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movl	$1, %ecx
	movl	$181, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	144(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$182, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	24(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	$1, %ecx
	movq	104(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$185, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	32(%rsp), %rdx
	movq	64(%rsp), %rsi
	movl	$1, %ecx
	movq	112(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$187, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
.L90:
	movq	8(%rsp), %rdx
	movq	120(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movl	$1, %ecx
	movl	$191, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	cmpb	$0, 79(%rsp)
	movl	%r12d, %r9d
	movl	%ebp, %r8d
	movq	104(%rsp), %rdi
	movq	112(%rsp), %rsi
	movl	%ebx, %ecx
	movq	120(%rsp), %rdx
	je	.L88
	call	_Z16launchKernelIsm2IfEvPKT_S2_PS0_jjj@PLT
	jmp	.L89
.L88:
	call	_Z16launchKernelTsm2IfEvPKT_S2_PS0_jjj@PLT
.L89:
	decl	%r13d
	jne	.L90
	movq	128(%rsp), %rdi
	xorl	%esi, %esi
	movl	$100, %r13d
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$199, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
.L93:
	movq	8(%rsp), %rdx
	movq	120(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaMemset@PLT
	movl	$1, %ecx
	movl	$201, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	cmpb	$0, 79(%rsp)
	movl	%r12d, %r9d
	movl	%ebp, %r8d
	movq	104(%rsp), %rdi
	movq	112(%rsp), %rsi
	movl	%ebx, %ecx
	movq	120(%rsp), %rdx
	je	.L91
	call	_Z16launchKernelIsm2IfEvPKT_S2_PS0_jjj@PLT
	jmp	.L92
.L91:
	call	_Z16launchKernelTsm2IfEvPKT_S2_PS0_jjj@PLT
.L92:
	decl	%r13d
	jne	.L93
	call	cudaGetLastError@PLT
	movl	$1, %ecx
	movl	$208, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	136(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$209, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	8(%rsp), %rdx
	movq	120(%rsp), %rsi
	movq	%r14, %rdi
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movl	$1, %ecx
	movl	$212, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	152(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$1, %ecx
	movl	$215, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	call	cudaDeviceSynchronize@PLT
	movl	$1, %ecx
	movl	$216, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	136(%rsp), %rdx
	leaq	88(%rsp), %rdi
	movq	128(%rsp), %rsi
	call	cudaEventElapsedTime@PLT
	movl	$1, %ecx
	movl	$217, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	152(%rsp), %rdx
	leaq	92(%rsp), %rdi
	movq	144(%rsp), %rsi
	call	cudaEventElapsedTime@PLT
	movl	$1, %ecx
	movl	$218, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movss	88(%rsp), %xmm6
	xorl	%esi, %esi
	xorl	%ecx, %ecx
	movss	92(%rsp), %xmm5
	divss	.LC10(%rip), %xmm6
	movb	$1, %dl
	divss	.LC10(%rip), %xmm5
	movss	%xmm6, 88(%rsp)
	movss	%xmm5, 92(%rsp)
.L94:
	cmpl	%ebp, %ecx
	setb	%al
	andb	%dl, %al
	je	.L114
	xorl	%edx, %edx
.L98:
	cmpl	%ebx, %edx
	jnb	.L105
	testb	%al, %al
	je	.L105
	movq	40(%rsp), %rdi
	movsd	.LC15(%rip), %xmm2
	leal	(%rdx,%rsi), %eax
	cvtss2sd	(%r14,%rax,4), %xmm1
	cvtss2sd	(%rdi,%rax,4), %xmm0
	call	_Z11approxEqualddd
	movl	72(%rsp), %edi
	testb	%al, %al
	cmove	%ecx, %edi
	cmove	%edx, %r15d
	incl	%edx
	movl	%edi, 72(%rsp)
	jmp	.L98
.L105:
	incl	%ecx
	addl	%ebx, %esi
	movl	%eax, %edx
	jmp	.L94
.L114:
	testb	%dl, %dl
	je	.L99
	movsd	16(%rsp), %xmm0
	movsd	16(%rsp), %xmm1
	cvtss2sd	%xmm5, %xmm5
	movb	$2, %al
	cvtss2sd	%xmm6, %xmm6
	divsd	.LC12(%rip), %xmm5
	divsd	.LC12(%rip), %xmm6
	movq	48(%rsp), %rdx
	divsd	%xmm6, %xmm0
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	divsd	%xmm5, %xmm1
	call	__printf_chk@PLT
	jmp	.L100
.L99:
	movl	72(%rsp), %eax
	movq	stderr(%rip), %rdi
	movl	%r15d, %r8d
	movl	$2, %esi
	leaq	.LC16(%rip), %rdx
	imull	%eax, %ebx
	leal	(%rbx,%r15), %eax
	movq	40(%rsp), %rbx
	cvtss2sd	(%r14,%rax,4), %xmm1
	cvtss2sd	(%rbx,%rax,4), %xmm0
	movl	72(%rsp), %ebx
	movb	$2, %al
	pushq	%rbx
	.cfi_def_cfa_offset 248
	pushq	%r15
	.cfi_def_cfa_offset 256
	movl	88(%rsp), %r9d
	movq	64(%rsp), %rcx
	call	__fprintf_chk@PLT
	popq	%rax
	.cfi_def_cfa_offset 248
	popq	%rdx
	.cfi_def_cfa_offset 240
.L100:
	movq	128(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$1, %ecx
	movl	$231, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	136(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$1, %ecx
	movl	$232, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	144(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$1, %ecx
	movl	$233, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	152(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$1, %ecx
	movl	$234, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	%r14, %rdi
	call	free@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$236, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	112(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$237, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movq	120(%rsp), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$238, %edx
	leaq	.LC7(%rip), %rsi
	movl	%eax, %edi
	call	_Z10cudaAssert9cudaErrorPKcib
	movb	$1, %al
.L76:
	movq	168(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L101
	call	__stack_chk_fail@PLT
.L101:
	addq	$184, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5111:
	.size	_Z10runKernelsIfEbPKT_S2_PS0_jjjb, .-_Z10runKernelsIfEbPKT_S2_PS0_jjjb
	.section	.text._Z9runMatmulIfEbRSiS0_RSob,"axG",@progbits,_Z9runMatmulIfEbRSiS0_RSob,comdat
	.weak	_Z9runMatmulIfEbRSiS0_RSob
	.type	_Z9runMatmulIfEbRSiS0_RSob, @function
_Z9runMatmulIfEbRSiS0_RSob:
.LFB4938:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%ecx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	%rsi, (%rsp)
	movq	%rdx, 8(%rsp)
	movl	$4, %edx
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movq	%rax, %rsi
	movq	%rax, 16(%rsp)
	call	_ZNSi4readEPcl@PLT
	leaq	64(%rsp), %rsi
	movl	$4, %edx
	movq	%r15, %rdi
	call	_ZNSi4readEPcl@PLT
	movq	(%rsp), %rdi
	leaq	68(%rsp), %rsi
	movl	$4, %edx
	call	_ZNSi4readEPcl@PLT
	leaq	60(%rsp), %rax
	movq	(%rsp), %rdi
	movl	$4, %edx
	movq	%rax, %rsi
	movq	%rax, 24(%rsp)
	call	_ZNSi4readEPcl@PLT
	movl	64(%rsp), %r8d
	movl	68(%rsp), %r9d
	movslq	56(%rsp), %rcx
	cmpl	%r9d, %r8d
	je	.L116
	pushq	%rsi
	.cfi_def_cfa_offset 152
	movl	68(%rsp), %eax
	leaq	.LC17(%rip), %rdx
	movl	$2, %esi
	movq	stderr(%rip), %rdi
	pushq	%rax
	.cfi_def_cfa_offset 160
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	popq	%rdi
	.cfi_def_cfa_offset 152
	popq	%r8
	.cfi_def_cfa_offset 144
	jmp	.L117
.L116:
	movslq	%r8d, %rbx
	movq	%rcx, %rdx
	movq	%rcx, 40(%rsp)
	imulq	%rbx, %rdx
	salq	$2, %rdx
	movq	%rdx, %rdi
	movq	%rdx, 32(%rsp)
	call	malloc@PLT
	movslq	60(%rsp), %r14
	movq	%rax, %r12
	imulq	%r14, %rbx
	leaq	0(,%rbx,4), %rdi
	call	malloc@PLT
	movq	40(%rsp), %rcx
	movq	%rax, %rbp
	imulq	%r14, %rcx
	leaq	0(,%rcx,4), %rdi
	call	malloc@PLT
	testq	%r12, %r12
	movq	%rax, %rbx
	sete	%al
	testq	%rbp, %rbp
	sete	%cl
	orb	%cl, %al
	jne	.L123
	testq	%rbx, %rbx
	movq	32(%rsp), %rdx
	jne	.L118
.L123:
	movq	stderr(%rip), %rdi
	leaq	.LC6(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L117:
	xorl	%r13d, %r13d
	jmp	.L115
.L118:
	movq	%r12, %rsi
	movq	%r15, %rdi
	movzbl	%r13b, %r13d
	call	_ZNSi4readEPcl@PLT
	movslq	64(%rsp), %rdx
	movslq	60(%rsp), %rax
	movq	%rbp, %rsi
	movq	(%rsp), %rdi
	imulq	%rax, %rdx
	salq	$2, %rdx
	call	_ZNSi4readEPcl@PLT
	movq	%rbx, %rdx
	movq	%rbp, %rsi
	movq	%r12, %rdi
	pushq	%rax
	.cfi_def_cfa_offset 152
	pushq	%r13
	.cfi_def_cfa_offset 160
	movl	80(%rsp), %r9d
	movl	76(%rsp), %r8d
	movl	72(%rsp), %ecx
	call	_Z10runKernelsIfEbPKT_S2_PS0_jjjb
	popq	%rdx
	.cfi_def_cfa_offset 152
	popq	%rcx
	.cfi_def_cfa_offset 144
	movl	%eax, %r13d
	testb	%al, %al
	jne	.L121
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
	jmp	.L117
.L121:
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$4, %edx
	call	_ZNSo5writeEPKcl@PLT
	movq	24(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$4, %edx
	call	_ZNSo5writeEPKcl@PLT
	movslq	56(%rsp), %rdx
	movslq	60(%rsp), %rax
	movq	%rbx, %rsi
	movq	8(%rsp), %rdi
	imulq	%rax, %rdx
	salq	$2, %rdx
	call	_ZNSo5writeEPKcl@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	%rbx, %rdi
	call	free@PLT
.L115:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L122
	call	__stack_chk_fail@PLT
.L122:
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	movl	%r13d, %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
.LFE4938:
	.size	_Z9runMatmulIfEbRSiS0_RSob, .-_Z9runMatmulIfEbRSiS0_RSob
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC20:
	.string	"-d"
.LC21:
	.string	"-i"
.LC22:
	.string	"Usage: %s [-d] [-i] matrixA matrixB matrixC\n"
.LC23:
	.string	"Cannot open %s for reading\n"
.LC24:
	.string	"Cannot open %s for writing\n"
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB4580:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4580
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$1, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	xorl	%r14d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%edi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	subq	$1608, %rsp
	.cfi_def_cfa_offset 1664
	movq	%fs:40, %rax
	movq	%rax, 1592(%rsp)
	xorl	%eax, %eax
.L126:
	cmpl	%r15d, %r13d
	jle	.L156
	movq	(%rbx,%r15,8), %rdi
	leaq	.LC20(%rip), %rsi
	movq	%rdi, (%rsp)
	call	strcmp@PLT
	testl	%eax, %eax
	je	.L144
	movq	(%rsp), %rdi
	leaq	.LC21(%rip), %rsi
	call	strcmp@PLT
	testl	%eax, %eax
	je	.L145
	cmpl	$2, %r12d
	jg	.L128
	movslq	%r12d, %rax
	movl	%r15d, 28(%rsp,%rax,4)
.L128:
	incl	%r12d
	jmp	.L127
.L144:
	movb	$1, %r14b
	jmp	.L127
.L145:
	movb	$1, %bpl
.L127:
	incq	%r15
	jmp	.L126
.L156:
	cmpl	$3, %r12d
	je	.L130
	movq	(%rbx), %rcx
	movl	$2, %esi
	xorl	%eax, %eax
	movl	$1, %ebp
	movq	stderr(%rip), %rdi
	leaq	.LC22(%rip), %rdx
.LEHB0:
	call	__fprintf_chk@PLT
	jmp	.L125
.L130:
	movslq	28(%rsp), %rax
	movl	$4, %edx
	leaq	(%rbx,%rax,8), %r13
	leaq	552(%rsp), %rax
	movq	0(%r13), %rsi
	movq	%rax, %rdi
	movq	%rax, (%rsp)
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode@PLT
.LEHE0:
	movslq	32(%rsp), %rax
	movl	$4, %edx
	leaq	(%rbx,%rax,8), %r12
	leaq	1072(%rsp), %rax
	movq	(%r12), %rsi
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
.LEHB1:
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode@PLT
.LEHE1:
	movslq	36(%rsp), %rax
	movl	$4, %edx
	leaq	(%rbx,%rax,8), %r15
	leaq	40(%rsp), %rbx
	movq	(%r15), %rsi
	movq	%rbx, %rdi
.LEHB2:
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEEC1EPKcSt13_Ios_Openmode@PLT
.LEHE2:
	testb	$5, 840(%rsp)
	je	.L132
	movq	0(%r13), %rcx
.L151:
	leaq	.LC23(%rip), %rdx
.L153:
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	xorl	%eax, %eax
.LEHB3:
	call	__fprintf_chk@PLT
	movl	$1, %ebp
	jmp	.L133
.L132:
	testb	$5, 1360(%rsp)
	je	.L134
	movq	(%r12), %rcx
	jmp	.L151
.L134:
	testb	$5, 320(%rsp)
	je	.L136
	movq	(%r15), %rcx
	leaq	.LC24(%rip), %rdx
	jmp	.L153
.L136:
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	movzbl	%bpl, %ecx
	movq	%rbx, %rdx
	testb	%r14b, %r14b
	je	.L137
	call	_Z9runMatmulIdEbRSiS0_RSob
	jmp	.L154
.L137:
	call	_Z9runMatmulIfEbRSiS0_RSob
.L154:
	movq	(%rsp), %rdi
	movl	%eax, %ebp
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEE5closeEv@PLT
	movq	8(%rsp), %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEE5closeEv@PLT
	movq	%rbx, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEE5closeEv@PLT
.LEHE3:
	xorl	$1, %ebp
	movzbl	%bpl, %ebp
.L133:
	movq	%rbx, %rdi
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	8(%rsp), %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	(%rsp), %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	jmp	.L125
.L148:
	endbr64
	movq	%rax, %rbp
.L139:
	movq	%rbx, %rdi
	movq	%rbp, %rbx
	call	_ZNSt14basic_ofstreamIcSt11char_traitsIcEED1Ev@PLT
	jmp	.L140
.L147:
	endbr64
	movq	%rax, %rbx
.L140:
	movq	8(%rsp), %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	jmp	.L141
.L146:
	endbr64
	movq	%rax, %rbx
.L141:
	movq	(%rsp), %rdi
	call	_ZNSt14basic_ifstreamIcSt11char_traitsIcEED1Ev@PLT
	movq	1592(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L155
	movq	%rbx, %rdi
.LEHB4:
	call	_Unwind_Resume@PLT
.LEHE4:
.L125:
	movq	1592(%rsp), %rax
	subq	%fs:40, %rax
	je	.L143
.L155:
	call	__stack_chk_fail@PLT
.L143:
	addq	$1608, %rsp
	.cfi_def_cfa_offset 56
	movl	%ebp, %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
.LFE4580:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4580:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4580-.LLSDACSB4580
.LLSDACSB4580:
	.uleb128 .LEHB0-.LFB4580
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB4580
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L146-.LFB4580
	.uleb128 0
	.uleb128 .LEHB2-.LFB4580
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L147-.LFB4580
	.uleb128 0
	.uleb128 .LEHB3-.LFB4580
	.uleb128 .LEHE3-.LEHB3
	.uleb128 .L148-.LFB4580
	.uleb128 0
	.uleb128 .LEHB4-.LFB4580
	.uleb128 .LEHE4-.LEHB4
	.uleb128 0
	.uleb128 0
.LLSDACSE4580:
	.section	.text.startup
	.size	main, .-main
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC2:
	.long	-1
	.long	2147483647
	.long	0
	.long	0
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC8:
	.long	0
	.long	1072693248
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC10:
	.long	1120403456
	.section	.rodata.cst8
	.align 8
.LC11:
	.long	0
	.long	1104006501
	.align 8
.LC12:
	.long	0
	.long	1083129856
	.align 8
.LC15:
	.long	-755914244
	.long	1062232653
	.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:

/**
 * Testbed declarations
 * by Cody Rivera, 2019-2020
 */

#ifndef _MULTIPLY_CUH
#define _MULTIPLY_CUH

// Functions
template<typename FloatType>
bool runKernels(const FloatType* A, const FloatType* B, FloatType* C,
                const unsigned int m, const unsigned int n,
                const unsigned int k, const bool runIsm2);

#endif
	.text
	.file	"multiply.hip"
	.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 __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
