#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <string.h>
#include <vector>

#define CSC(call)                                   \
do {                                                 \
	cudaError_t res = call;                             \
	if (res != cudaSuccess) {                            \
		fprintf(stderr, "ERROR in %s:%d. Message: %s\n",    \
			__FILE__, __LINE__, cudaGetErrorString(res));      \
		exit(0);                                              \
	}                                                        \
} while(0)

typedef unsigned char uchar;

#define NC_MAX 32 // максимальное количество классов

double3 avg[NC_MAX]; // средние
double cov[3 * 3 * NC_MAX]; // ковариационные матрицы
__constant__ double3 AVG[NC_MAX];
__constant__ double COV[3 * 3 * NC_MAX];

__host__ __device__
double dot(const double *A, const double *a, const double *b, int n)
{
    /*
    * res = a ^ T * A * b
    */
    int i, j;
    double res = 0;
    for (i = 0; i < n; i++)
    {
        for (j = 0; j < n; j++)
        {
            res += A[i * n + j] * a[i] * b[j];
        }
    }
    return res;
}

__host__ __device__ // работает везде
void mat_mul(const double *A, const double *B, double *C, int m, int n, int l, double alpha=1, double beta=0)
{
    /*
    * C = alpha * A * B + beta * C
    *
    * m - число строк A
    * n - число столбцов A и число строк B
    * l - число столбцов B
    */
    int i, j, k;
    double dot;

    for (i = 0; i < m; i++)
    {
        for (j = 0; j < l; j++)
        {
            dot = 0;
            for (k = 0; k < n; k++)
            {
                dot += A[i * n + k] * B[k * l + j];
            }
            C[i * l + j] = alpha * dot + (beta == 0 ? 0 : beta * C[i * l + j]);
        }
    }
}

__host__ __device__
void mat_mul_C(const double *A, double *B, double c, int m, int n)
{
    /*
    * B = c * A
    */
    int i, j;
    for (i = 0; i < m; i++)
    {
        for (j = 0; j < n; j++)
        {
            B[i * n + j] = c * A[i * n + j];
        }
    }
}

__host__ __device__
void mat_sum(const double *A, const double *B, double *C, int m, int n, double alpha=1, double beta=1)
{
    /*
    * C = alpha * A + beta * B
    */
    int i, j;
    for (i = 0; i < m; i++)
    {
        for (j = 0; j < n; j++)
        {
            C[i * n + j] = alpha * A[i * n + j] + beta * B[i * n + j];
        }
    }
}

__host__ __device__
void mat_set(double *A, int n, double alpha=0, double beta=0)
{
    /*
    * Инициализация A значениями alpha на диагонали и
    * значениями beta вне диагонали
    */
    int i, j;
    for (i = 0; i < n; i++)
    {
        for (j = 0; j < n; j++)
        {
            if (i == j)
            {
                A[i * n + j] = alpha;
            }
            else
            {
                A[i * n + j] = beta;
            }
        }
    }
}

void mat_swap_rows(double *A, int n, int i1, int i2)
{
    int j;
    double tmp;
    for (j = 0; j < n; j++)
    {
        tmp = A[i1 * n + j];
        A[i1 * n + j] = A[i2 * n + j];
        A[i2 * n + j] = tmp;
    }
}

void mat_transpose(const double *A, double *A_t, int n)
{
    int i, j;
    for (i = 0; i < n; i++)
    {
        for (j = 0; j < n; j++)
        {
            A_t[i * n + j] = A[j * n + i];
        }
    }
}

bool LUP(double *A, int n, int *pi, double eps=1e-10)
{
    int i, j, k, k_;
    int tmp;

    for (i = 0; i < n; i++)
    {
        pi[i] = i;
    }

    for (k = 0; k < n; k++)
    {
        k_ = k;
        for (i = k + 1; i < n; i++)
        {
            if (fabs(A[i * n + k]) > fabs(A[k_ * n + k]))
            {
                k_ = i;
            }
        }

        if (fabs(A[k_ * n + k]) < eps)
        {
            return false;
        }

        if (k != k_)
        {
            tmp = pi[k];
            pi[k] = pi[k_];
            pi[k_] = tmp;
            mat_swap_rows(A, n, k, k_);
        }

        for (i = k + 1; i < n; i++)
        {
            A[i * n + k] /= A[k * n + k];
            for (j = k + 1; j < n; j++)
            {
                A[i * n + j] -= A[i * n + k] * A[k * n + j];
            }
        }
    }
    return true;
}

void LUP_solve(const double *LU, const int *pi, const double *b, int n, double *x, double *work)
{
    /*
    * work - вектор размерности n
    */
    int i, j;
    double sum;
    double *y = work;

    memset(y, 0, n * sizeof(double));

    for (i = 0; i < n; i++) // прямой ход
    {
        for (sum = 0, j = 0; j <= i - 1; sum += LU[i * n + j] * y[j], j++);
        y[i] = b[pi[i]] - sum;
    }

    for (i = n - 1; i >= 0; i--) // обратный ход
    {
        for (sum = 0, j = i + 1; j < n; sum += LU[i * n + j] * x[j], j++);
        x[i] = (y[i] - sum) / LU[i * n + i];
    }
}

void LUP_inv_mat(double *A, double *A_inv, int n, double *work, int *iwork, double eps=1e-10) // инвариантная матрица при помощи LUP разложения (украдено из Кормана)
{
    /*
    * work  - вектор длины n ^ 2 + 2 * n
    * iwork - вектор длины n
    */
    int i;
    double *X, *e, *space;
    int *pi;

    X = work;
    e = X + n * n;
    space = e + n;

    pi = iwork;

    memset(e, 0, n * sizeof(double));
    e[0] = 1;

    LUP(A, n, pi, eps);
    for (i = 0; i < n - 1; i++)
    {
        LUP_solve(A, pi, e, n, X + i * n, space);
        e[i] = 0;
        e[i + 1] = 1;
    }
    LUP_solve(A, pi, e, n, X + i * n, space);
    mat_transpose(X, A_inv, n);
}

__global__
void kernel(int nc, uchar4 *im, int w, int h)
{
    int i, j, jc, idx;
    int offset;
    double3 to_dot;
    double dist, dist_max;

    idx = blockDim.x * blockIdx.x + threadIdx.x; // абсолютный номер потока
    offset = blockDim.x * gridDim.x; // общее кол-во потоков

    for (i = idx; i < w * h; i += offset) // идем по всем пикселям
    {
        dist_max = -INFINITY;
        jc = 0;
        for (j = 0; j < nc; j++) // цикл по числу классов
        {

            to_dot = make_double3((double)im[i].x - AVG[j].x,
								(double)im[i].y - AVG[j].y,
								(double)im[i].z - AVG[j].z);
            dist = -dot(COV + j * 3 * 3, (double *)&to_dot, (double *)&to_dot, 3); // произведение из основной формулы
            if (dist > dist_max) // для каждого класса вычисляем dist и выбираем максимум
            {
                jc = j;
                dist_max = dist;
            }
        }
        im[i].w = jc; // записываем номер класса с max dist
    }
}

int main()
{
    int nc, np, i, j, x, y, w, h;
    uchar4 *im = NULL, *im_dev = NULL;
    std::vector<double3> v(0);
    double mat[3 * 3],
           work[3 * 3 + 2 * 3];
    int iwork[3];

    FILE *fp;
    char name_src_im[256], name_dst_im[256];

    dim3 blocks(256), threads(256);

    scanf("%s\n%s\n%d", name_src_im, name_dst_im, &nc);

    fp = fopen(name_src_im, "rb");
    if (fp == NULL)
    {
        fprintf(stderr, "Error: can't open %s\n", name_src_im);
        return 0;
    }
    fread(&w, sizeof(int), 1, fp);
    fread(&h, sizeof(int), 1, fp);

    im = (uchar4 *)malloc(w * h * sizeof(uchar4));
    if (im == NULL)
    {
        fprintf(stderr, "Error: not enough memory in CPU\n");
        goto FREE;
    }
    CSC(cudaMalloc(&im_dev, w * h * sizeof(uchar4)));

    fread(im, sizeof(uchar4), w * h, fp); // считывание пикселей
    fclose(fp);

    for (j = 0; j < nc; j++) // цикл по числу классов
    {
        scanf("%d", &np); // количество пикселей в классе
        if(v.size() < np) v.resize(np);

        avg[j] = make_double3(0, 0, 0);
        for (i = 0; i < np; i++) // просчёт среднего по j-му классу
        {
            scanf("%d %d", &x, &y);
            v[i] = make_double3((double)im[y * w + x].x,
								(double)im[y * w + x].y,
								(double)im[y * w + x].z);

            avg[j].x += v[i].x;
            avg[j].y += v[i].y;
            avg[j].z += v[i].z;
        }
        avg[j].x /= np;
        avg[j].y /= np;
        avg[j].z /= np;

        if (np > 1)
        {
            mat_set(mat, 3, 0, 0);          // инициализация нулями
            for (i = 0; i < np; i++)        // просчёт ковариации по i-му классу
            {
                v[i].x = v[i].x - avg[j].x;
                v[i].y = v[i].y - avg[j].y;
                v[i].z = v[i].z - avg[j].z;
								// double3 приводим к вектору double
                mat_mul((double *)(v.data() + i), (double *)(v.data() + i), mat, 3, 1, 3, 1, 1); // умножение векторов как матриц 3х1 1х3 получаем 3х3
            }
            mat_mul_C(mat, mat, 1. / (np - 1), 3, 3); // делим матрицу на коэффициент
            LUP_inv_mat(mat, cov + j * 3 * 3, 3, work, iwork); // получаем матрицу ковариации для j-того класса (work, iwork -- место, чтобы внутри функции не делать malloc)
        }
        else
        {
            mat_set(cov + j * 3 * 3, 3, 1, 0); // единичная матрица
        }
    }

    /* Копирование в константную память */
    CSC(cudaMemcpyToSymbol(AVG, avg, nc * sizeof(double3), 0, cudaMemcpyHostToDevice));
    CSC(cudaMemcpyToSymbol(COV, cov, nc * 3 * 3 * sizeof(double), 0, cudaMemcpyHostToDevice));

    /* Копирование изображения */
    cudaMemcpy(im_dev, im, w * h * sizeof(uchar4), cudaMemcpyHostToDevice); // копирование изображения

    kernel<<<blocks, threads>>>(nc, im_dev, w, h);
    CSC(cudaGetLastError());
    CSC(cudaMemcpy(im, im_dev, w * h * sizeof(uchar4), cudaMemcpyDeviceToHost));

    fp = fopen(name_dst_im, "wb");
    if (fp == NULL)
    {
        fprintf(stderr, "Error: can't open %s\n", name_dst_im);
        goto FREE;
    }
    fwrite(&w, sizeof(int), 1, fp);
    fwrite(&h, sizeof(int), 1, fp);
    fwrite(im, sizeof(uchar4), w * h, fp);
    fclose(fp);

FREE:
    free(im);
    cudaFree(im_dev);
    return 0;
}

	.file	"tmpxft_0033e1aa_00000000-6_lab3.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2909:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2909:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z3dotPKdS0_S0_i
	.type	_Z3dotPKdS0_S0_i, @function
_Z3dotPKdS0_S0_i:
.LFB2896:
	.cfi_startproc
	endbr64
	movslq	%ecx, %r9
	xorl	%r8d, %r8d
	xorps	%xmm0, %xmm0
	salq	$3, %r9
.L3:
	cmpl	%r8d, %ecx
	jle	.L2
	movsd	(%rsi,%r8,8), %xmm2
	xorl	%eax, %eax
.L4:
	movsd	(%rdi,%rax,8), %xmm1
	mulsd	%xmm2, %xmm1
	mulsd	(%rdx,%rax,8), %xmm1
	incq	%rax
	addsd	%xmm1, %xmm0
	cmpl	%eax, %ecx
	jg	.L4
	incq	%r8
	addq	%r9, %rdi
	jmp	.L3
.L2:
	ret
	.cfi_endproc
.LFE2896:
	.size	_Z3dotPKdS0_S0_i, .-_Z3dotPKdS0_S0_i
	.globl	_Z7mat_mulPKdS0_Pdiiidd
	.type	_Z7mat_mulPKdS0_Pdiiidd, @function
_Z7mat_mulPKdS0_Pdiiidd:
.LFB2897:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdx, %r11
	xorps	%xmm4, %xmm4
	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
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movslq	%r9d, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	salq	$3, %rbp
	xorl	%ebx, %ebx
	movq	%rsi, -8(%rsp)
	xorl	%esi, %esi
.L9:
	cmpl	%r13d, %r12d
	jge	.L8
	movslq	%ebx, %rax
	movslq	%esi, %rdx
	movq	-8(%rsp), %r10
	xorl	%ecx, %ecx
	leaq	(%r11,%rax,8), %rax
	leaq	(%rdi,%rdx,8), %r15
.L15:
	cmpl	%r9d, %ecx
	jge	.L20
	movq	%r10, %r14
	xorl	%edx, %edx
	xorps	%xmm2, %xmm2
.L13:
	cmpl	%edx, %r8d
	jle	.L21
	movsd	(%r15,%rdx,8), %xmm3
	mulsd	(%r14), %xmm3
	incq	%rdx
	addq	%rbp, %r14
	addsd	%xmm3, %xmm2
	jmp	.L13
.L21:
	mulsd	%xmm0, %xmm2
	ucomisd	%xmm4, %xmm1
	jp	.L18
	xorps	%xmm3, %xmm3
	je	.L11
.L18:
	movsd	(%rax), %xmm3
	mulsd	%xmm1, %xmm3
.L11:
	addsd	%xmm3, %xmm2
	incl	%ecx
	addq	$8, %rax
	addq	$8, %r10
	movsd	%xmm2, -8(%rax)
	jmp	.L15
.L20:
	incl	%r12d
	addl	%r9d, %ebx
	addl	%r8d, %esi
	jmp	.L9
.L8:
	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
.LFE2897:
	.size	_Z7mat_mulPKdS0_Pdiiidd, .-_Z7mat_mulPKdS0_Pdiiidd
	.globl	_Z9mat_mul_CPKdPddii
	.type	_Z9mat_mul_CPKdPddii, @function
_Z9mat_mul_CPKdPddii:
.LFB2898:
	.cfi_startproc
	endbr64
	movq	%rdi, %r9
	movq	%rsi, %r10
	xorl	%edi, %edi
	xorl	%r8d, %r8d
.L23:
	cmpl	%edx, %r8d
	jge	.L22
	movslq	%edi, %rax
	xorl	%esi, %esi
	salq	$3, %rax
.L26:
	cmpl	%ecx, %esi
	jge	.L27
	movsd	(%r9,%rax), %xmm1
	incl	%esi
	mulsd	%xmm0, %xmm1
	movsd	%xmm1, (%r10,%rax)
	addq	$8, %rax
	jmp	.L26
.L27:
	incl	%r8d
	addl	%ecx, %edi
	jmp	.L23
.L22:
	ret
	.cfi_endproc
.LFE2898:
	.size	_Z9mat_mul_CPKdPddii, .-_Z9mat_mul_CPKdPddii
	.globl	_Z7mat_sumPKdS0_Pdiidd
	.type	_Z7mat_sumPKdS0_Pdiidd, @function
_Z7mat_sumPKdS0_Pdiidd:
.LFB2899:
	.cfi_startproc
	endbr64
	movq	%rsi, %r9
	movl	%ecx, %r11d
	movq	%rdx, %r10
	xorl	%ecx, %ecx
	xorl	%esi, %esi
.L29:
	cmpl	%r11d, %esi
	jge	.L28
	movslq	%ecx, %rax
	xorl	%edx, %edx
	salq	$3, %rax
.L32:
	cmpl	%r8d, %edx
	jge	.L33
	movsd	(%rdi,%rax), %xmm2
	movsd	(%r9,%rax), %xmm3
	incl	%edx
	mulsd	%xmm0, %xmm2
	mulsd	%xmm1, %xmm3
	addsd	%xmm3, %xmm2
	movsd	%xmm2, (%r10,%rax)
	addq	$8, %rax
	jmp	.L32
.L33:
	incl	%esi
	addl	%r8d, %ecx
	jmp	.L29
.L28:
	ret
	.cfi_endproc
.LFE2899:
	.size	_Z7mat_sumPKdS0_Pdiidd, .-_Z7mat_sumPKdS0_Pdiidd
	.globl	_Z7mat_setPdidd
	.type	_Z7mat_setPdidd, @function
_Z7mat_setPdidd:
.LFB2900:
	.cfi_startproc
	endbr64
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.L35:
	cmpl	%esi, %edx
	jge	.L41
	movslq	%ecx, %rax
	leaq	(%rdi,%rax,8), %r8
	xorl	%eax, %eax
.L37:
	movaps	%xmm0, %xmm2
	cmpl	%eax, %edx
	je	.L36
	movaps	%xmm1, %xmm2
.L36:
	movsd	%xmm2, (%r8,%rax,8)
	incq	%rax
	cmpl	%eax, %esi
	jg	.L37
	incl	%edx
	addl	%esi, %ecx
	jmp	.L35
.L41:
	ret
	.cfi_endproc
.LFE2900:
	.size	_Z7mat_setPdidd, .-_Z7mat_setPdidd
	.globl	_Z13mat_swap_rowsPdiii
	.type	_Z13mat_swap_rowsPdiii, @function
_Z13mat_swap_rowsPdiii:
.LFB2901:
	.cfi_startproc
	endbr64
	imull	%esi, %edx
	imull	%esi, %ecx
	movslq	%edx, %rdx
	movslq	%ecx, %rcx
	leaq	(%rdi,%rdx,8), %rax
	leaq	(%rdi,%rcx,8), %rdx
	xorl	%ecx, %ecx
.L43:
	cmpl	%esi, %ecx
	jge	.L45
	movsd	(%rax), %xmm0
	movsd	(%rdx), %xmm1
	incl	%ecx
	addq	$8, %rax
	addq	$8, %rdx
	movsd	%xmm1, -8(%rax)
	movsd	%xmm0, -8(%rdx)
	jmp	.L43
.L45:
	ret
	.cfi_endproc
.LFE2901:
	.size	_Z13mat_swap_rowsPdiii, .-_Z13mat_swap_rowsPdiii
	.globl	_Z13mat_transposePKdPdi
	.type	_Z13mat_transposePKdPdi, @function
_Z13mat_transposePKdPdi:
.LFB2902:
	.cfi_startproc
	endbr64
	movslq	%edx, %rax
	xorl	%r8d, %r8d
	salq	$3, %rax
.L47:
	cmpl	%edx, %r8d
	jge	.L51
	movq	%rdi, %r9
	xorl	%ecx, %ecx
.L48:
	movsd	(%r9), %xmm0
	addq	%rax, %r9
	movsd	%xmm0, (%rsi,%rcx,8)
	incq	%rcx
	cmpl	%ecx, %edx
	jg	.L48
	incl	%r8d
	addq	$8, %rdi
	addq	%rax, %rsi
	jmp	.L47
.L51:
	ret
	.cfi_endproc
.LFE2902:
	.size	_Z13mat_transposePKdPdi, .-_Z13mat_transposePKdPdi
	.globl	_Z3LUPPdiPid
	.type	_Z3LUPPdiPid, @function
_Z3LUPPdiPid:
.LFB2903:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdx, %r10
	movaps	%xmm0, %xmm2
	xorl	%eax, %eax
	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	$16, %rsp
	.cfi_def_cfa_offset 72
.L53:
	cmpl	%eax, %esi
	jle	.L69
	movl	%eax, (%r10,%rax,4)
	incq	%rax
	jmp	.L53
.L69:
	movslq	%esi, %r9
	leal	1(%rsi), %eax
	movq	.LC1(%rip), %xmm3
	xorl	%ebp, %ebp
	salq	$3, %r9
	movl	%eax, 8(%rsp)
	xorl	%r14d, %r14d
	xorl	%r8d, %r8d
	leaq	8(%r9), %rax
	leaq	(%rdi,%r9), %r11
	movq	%rax, (%rsp)
	movq	%r11, %r12
.L55:
	movl	%r8d, 12(%rsp)
	movl	%r8d, %r15d
	cmpl	%r8d, %esi
	jle	.L70
	leal	1(%r8), %ebx
	movq	%r12, %r13
	movl	%r8d, %ecx
	movl	%ebx, %edx
.L56:
	movl	%ecx, %eax
	imull	%esi, %eax
	addl	%r15d, %eax
	cltq
	movsd	(%rdi,%rax,8), %xmm1
	andps	%xmm3, %xmm1
	cmpl	%esi, %edx
	je	.L71
	movsd	0(%r13), %xmm0
	andps	%xmm3, %xmm0
	ucomisd	%xmm1, %xmm0
	cmova	%edx, %ecx
	incl	%edx
	addq	%r9, %r13
	jmp	.L56
.L71:
	comisd	%xmm1, %xmm2
	ja	.L67
	movl	12(%rsp), %eax
	cmpl	%eax, %ecx
	je	.L60
	movslq	%ecx, %rax
	movl	(%r10,%r8,4), %edx
	leaq	(%r10,%rax,4), %rax
	movl	(%rax), %r13d
	movl	%r13d, (%r10,%r8,4)
	movl	%edx, (%rax)
	movl	%r15d, %edx
	call	_Z13mat_swap_rowsPdiii
.L60:
	movslq	%ebp, %rax
	movslq	%r14d, %rdx
	addl	%esi, %ebp
	leaq	(%rdi,%rax,8), %rcx
	leaq	(%rdi,%rdx,8), %r13
	movq	%r11, %rax
.L61:
	cmpl	%esi, %ebx
	je	.L72
	movsd	(%rax,%r8,8), %xmm0
	leaq	1(%r8), %rdx
	divsd	0(%r13), %xmm0
	movsd	%xmm0, (%rax,%r8,8)
.L62:
	cmpl	%edx, %esi
	jle	.L73
	movsd	(%rax,%r8,8), %xmm1
	mulsd	(%rcx,%rdx,8), %xmm1
	movsd	(%rax,%rdx,8), %xmm0
	subsd	%xmm1, %xmm0
	movsd	%xmm0, (%rax,%rdx,8)
	incq	%rdx
	jmp	.L62
.L73:
	incl	%ebx
	addq	%r9, %rax
	jmp	.L61
.L72:
	movl	8(%rsp), %eax
	incq	%r8
	addq	%r9, %r11
	addl	%eax, %r14d
	movq	(%rsp), %rax
	addq	%rax, %r12
	jmp	.L55
.L70:
	movb	$1, %al
	jmp	.L52
.L67:
	xorl	%eax, %eax
.L52:
	addq	$16, %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
.LFE2903:
	.size	_Z3LUPPdiPid, .-_Z3LUPPdiPid
	.globl	_Z9LUP_solvePKdPKiS0_iPdS3_
	.type	_Z9LUP_solvePKdPKiS0_iPdS3_, @function
_Z9LUP_solvePKdPKiS0_iPdS3_:
.LFB2904:
	.cfi_startproc
	endbr64
	movslq	%ecx, %rcx
	xorl	%eax, %eax
	movq	%rdi, %r10
	movq	%rdx, %r11
	movq	%r9, %rdi
	movq	%rcx, %rdx
	salq	$3, %rcx
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	rep stosb
	xorl	%edi, %edi
	xorl	%eax, %eax
.L75:
	movl	%eax, %ebp
	cmpl	%eax, %edx
	jle	.L77
	movslq	%edi, %rcx
	xorps	%xmm0, %xmm0
	leaq	(%r10,%rcx,8), %rbx
	xorl	%ecx, %ecx
.L78:
	cmpl	%ecx, %ebp
	jle	.L84
	movsd	(%rbx,%rcx,8), %xmm1
	mulsd	(%r9,%rcx,8), %xmm1
	incq	%rcx
	addsd	%xmm1, %xmm0
	jmp	.L78
.L84:
	movslq	(%rsi,%rax,4), %rcx
	addl	%edx, %edi
	movsd	(%r11,%rcx,8), %xmm1
	subsd	%xmm0, %xmm1
	movsd	%xmm1, (%r9,%rax,8)
	incq	%rax
	jmp	.L75
.L77:
	leal	-1(%rdx), %eax
	movl	%edx, %esi
	movslq	%eax, %rcx
	imull	%edx, %eax
	notl	%esi
	movslq	%esi, %rsi
	salq	$3, %rsi
	movslq	%eax, %rdi
	addq	%rcx, %rdi
	leaq	(%r10,%rdi,8), %r11
.L79:
	testl	%ecx, %ecx
	js	.L85
	movslq	%eax, %rbx
	leaq	1(%rcx), %rdi
	xorps	%xmm1, %xmm1
	leaq	(%r10,%rbx,8), %rbx
.L80:
	cmpl	%edi, %edx
	jle	.L86
	movsd	(%rbx,%rdi,8), %xmm0
	mulsd	(%r8,%rdi,8), %xmm0
	incq	%rdi
	addsd	%xmm0, %xmm1
	jmp	.L80
.L86:
	movsd	(%r9,%rcx,8), %xmm0
	subl	%edx, %eax
	subsd	%xmm1, %xmm0
	divsd	(%r11), %xmm0
	addq	%rsi, %r11
	movsd	%xmm0, (%r8,%rcx,8)
	decq	%rcx
	jmp	.L79
.L85:
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2904:
	.size	_Z9LUP_solvePKdPKiS0_iPdS3_, .-_Z9LUP_solvePKdPKiS0_iPdS3_
	.globl	_Z11LUP_inv_matPdS_iS_Pid
	.type	_Z11LUP_inv_matPdS_iS_Pid, @function
_Z11LUP_inv_matPdS_iS_Pid:
.LFB2905:
	.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
	xorl	%r14d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rcx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movslq	%edx, %rbp
	movq	%r8, %rdx
	movq	%r12, %r15
	movl	%ebp, %eax
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rbp, %rbx
	imull	%ebp, %eax
	salq	$3, %rbp
	subq	$32, %rsp
	.cfi_def_cfa_offset 88
	movsd	.LC2(%rip), %xmm4
	cltq
	movq	%rdi, (%rsp)
	leaq	(%rcx,%rax,8), %r13
	xorl	%eax, %eax
	movq	%rbp, %rcx
	movq	%rsi, 16(%rsp)
	movq	%r13, %rdi
	leaq	0(%r13,%rbp), %r9
	movl	%ebx, %esi
	movq	%r8, 8(%rsp)
	rep stosb
	movq	(%rsp), %rdi
	movsd	%xmm4, 0(%r13)
	movq	%r9, 24(%rsp)
	call	_Z3LUPPdiPid
	movq	24(%rsp), %r9
.L88:
	leal	-1(%rbx), %edx
	movq	%r15, %r8
	addq	%rbp, %r15
	cmpl	%r14d, %edx
	jle	.L91
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdi
	movl	%ebx, %ecx
	movq	%r13, %rdx
	call	_Z9LUP_solvePKdPKiS0_iPdS3_
	movsd	%xmm4, 8(%r13,%r14,8)
	movq	$0x000000000, 0(%r13,%r14,8)
	incq	%r14
	jmp	.L88
.L91:
	movl	%edx, %eax
	xorl	%edx, %edx
	testl	%ebx, %ebx
	movq	8(%rsp), %rsi
	cmovle	%rdx, %rax
	movq	(%rsp), %rdi
	movq	%r13, %rdx
	movl	%ebx, %ecx
	imulq	%rbp, %rax
	leaq	(%r12,%rax), %r8
	call	_Z9LUP_solvePKdPKiS0_iPdS3_
	movq	16(%rsp), %rsi
	addq	$32, %rsp
	.cfi_def_cfa_offset 56
	movl	%ebx, %edx
	movq	%r12, %rdi
	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
	jmp	_Z13mat_transposePKdPdi
	.cfi_endproc
.LFE2905:
	.size	_Z11LUP_inv_matPdS_iS_Pid, .-_Z11LUP_inv_matPdS_iS_Pid
	.globl	_Z34__device_stub__Z6kerneliP6uchar4iiiP6uchar4ii
	.type	_Z34__device_stub__Z6kerneliP6uchar4iiiP6uchar4ii, @function
_Z34__device_stub__Z6kerneliP6uchar4iiiP6uchar4ii:
.LFB2931:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 24(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 12(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	12(%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	.L92
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z6kerneliP6uchar4ii(%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
.L92:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L94
	call	__stack_chk_fail@PLT
.L94:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2931:
	.size	_Z34__device_stub__Z6kerneliP6uchar4iiiP6uchar4ii, .-_Z34__device_stub__Z6kerneliP6uchar4iiiP6uchar4ii
	.globl	_Z6kerneliP6uchar4ii
	.type	_Z6kerneliP6uchar4ii, @function
_Z6kerneliP6uchar4ii:
.LFB2932:
	.cfi_startproc
	endbr64
	jmp	_Z34__device_stub__Z6kerneliP6uchar4iiiP6uchar4ii
	.cfi_endproc
.LFE2932:
	.size	_Z6kerneliP6uchar4ii, .-_Z6kerneliP6uchar4ii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"_Z6kerneliP6uchar4ii"
.LC4:
	.string	"AVG"
.LC5:
	.string	"COV"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2934:
	.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	.LC3(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z6kerneliP6uchar4ii(%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	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	movl	$768, %r9d
	leaq	_ZL3AVG(%rip), %rsi
	pushq	$1
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$1
	.cfi_def_cfa_offset 32
	leaq	.LC5(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$2304, %r9d
	leaq	_ZL3COV(%rip), %rsi
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.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
.LFE2934:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text._ZNSt6vectorI7double3SaIS0_EED2Ev,"axG",@progbits,_ZNSt6vectorI7double3SaIS0_EED5Ev,comdat
	.align 2
	.weak	_ZNSt6vectorI7double3SaIS0_EED2Ev
	.type	_ZNSt6vectorI7double3SaIS0_EED2Ev, @function
_ZNSt6vectorI7double3SaIS0_EED2Ev:
.LFB2992:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L99
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L99:
	ret
	.cfi_endproc
.LFE2992:
	.size	_ZNSt6vectorI7double3SaIS0_EED2Ev, .-_ZNSt6vectorI7double3SaIS0_EED2Ev
	.weak	_ZNSt6vectorI7double3SaIS0_EED1Ev
	.set	_ZNSt6vectorI7double3SaIS0_EED1Ev,_ZNSt6vectorI7double3SaIS0_EED2Ev
	.section	.text._ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIP7double3mEET_S4_T0_,"axG",@progbits,_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIP7double3mEET_S4_T0_,comdat
	.weak	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIP7double3mEET_S4_T0_
	.type	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIP7double3mEET_S4_T0_, @function
_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIP7double3mEET_S4_T0_:
.LFB3130:
	.cfi_startproc
	endbr64
	movq	%rdi, %rdx
	testq	%rsi, %rsi
	je	.L101
	movq	$0x000000000, (%rdi)
	leaq	24(%rdi), %rdx
	movq	$0x000000000, 8(%rdi)
	movq	$0x000000000, 16(%rdi)
	decq	%rsi
	je	.L101
	imulq	$24, %rsi, %rsi
	movq	%rdx, %rax
	addq	%rdx, %rsi
.L103:
	cmpq	%rax, %rsi
	je	.L110
	movups	(%rdi), %xmm0
	addq	$24, %rax
	movups	%xmm0, -24(%rax)
	movq	16(%rdi), %rcx
	movq	%rcx, -8(%rax)
	jmp	.L103
.L110:
	movabsq	$768614336404564651, %rcx
	subq	%rdi, %rsi
	leaq	-24(%rsi), %rax
	shrq	$3, %rax
	imulq	%rcx, %rax
	movabsq	$2305843009213693951, %rcx
	andq	%rcx, %rax
	imulq	$24, %rax, %rax
	addq	%rax, %rdx
.L101:
	movq	%rdx, %rax
	ret
	.cfi_endproc
.LFE3130:
	.size	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIP7double3mEET_S4_T0_, .-_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIP7double3mEET_S4_T0_
	.section	.rodata._ZNSt6vectorI7double3SaIS0_EE17_M_default_appendEm.str1.1,"aMS",@progbits,1
.LC6:
	.string	"vector::_M_default_append"
	.section	.text._ZNSt6vectorI7double3SaIS0_EE17_M_default_appendEm,"axG",@progbits,_ZNSt6vectorI7double3SaIS0_EE17_M_default_appendEm,comdat
	.align 2
	.weak	_ZNSt6vectorI7double3SaIS0_EE17_M_default_appendEm
	.type	_ZNSt6vectorI7double3SaIS0_EE17_M_default_appendEm, @function
_ZNSt6vectorI7double3SaIS0_EE17_M_default_appendEm:
.LFB3059:
	.cfi_startproc
	endbr64
	testq	%rsi, %rsi
	je	.L125
	movabsq	$-6148914691236517205, %rdx
	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
	movq	%rdi, %r13
	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
	movq	%rsi, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movq	8(%rdi), %rdi
	movq	16(%r13), %rax
	movq	0(%r13), %r14
	subq	%rdi, %rax
	sarq	$3, %rax
	imulq	%rdx, %rax
	cmpq	%rsi, %rax
	jb	.L113
	call	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIP7double3mEET_S4_T0_
	movq	%rax, 8(%r13)
	jmp	.L111
.L113:
	movabsq	$384307168202282325, %rax
	subq	%r14, %rdi
	movq	%rdi, %r15
	movq	%rdi, 8(%rsp)
	sarq	$3, %r15
	imulq	%rdx, %r15
	subq	%r15, %rax
	cmpq	%rsi, %rax
	jnb	.L114
	leaq	.LC6(%rip), %rdi
	call	_ZSt20__throw_length_errorPKc@PLT
.L114:
	movabsq	$384307168202282325, %rax
	cmpq	%r15, %rsi
	movq	%r15, %r12
	cmovnb	%rsi, %r12
	addq	%r15, %r12
	cmpq	%rax, %r12
	cmova	%rax, %r12
	imulq	$24, %r12, %r12
	movq	%r12, %rdi
	call	_Znwm@PLT
	movq	%rbx, %rsi
	movq	%rax, %rbp
	movq	8(%rsp), %rax
	leaq	0(%rbp,%rax), %rdi
	call	_ZNSt27__uninitialized_default_n_1ILb1EE18__uninit_default_nIP7double3mEET_S4_T0_
	cmpq	$0, 8(%rsp)
	je	.L115
	movq	8(%rsp), %rdx
	movq	%r14, %rsi
	movq	%rbp, %rdi
	call	memmove@PLT
	jmp	.L116
.L115:
	testq	%r14, %r14
	je	.L117
.L116:
	movq	16(%r13), %rsi
	movq	%r14, %rdi
	subq	%r14, %rsi
	call	_ZdlPvm@PLT
.L117:
	addq	%r15, %rbx
	movq	%rbp, 0(%r13)
	imulq	$24, %rbx, %rbx
	addq	%rbp, %rbx
	addq	%r12, %rbp
	movq	%rbx, 8(%r13)
	movq	%rbp, 16(%r13)
.L111:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
.L125:
	.cfi_restore 3
	.cfi_restore 6
	.cfi_restore 12
	.cfi_restore 13
	.cfi_restore 14
	.cfi_restore 15
	ret
	.cfi_endproc
.LFE3059:
	.size	_ZNSt6vectorI7double3SaIS0_EE17_M_default_appendEm, .-_ZNSt6vectorI7double3SaIS0_EE17_M_default_appendEm
	.section	.rodata.str1.1
.LC7:
	.string	"%s\n%s\n%d"
.LC8:
	.string	"rb"
.LC9:
	.string	"Error: can't open %s\n"
.LC10:
	.string	"Error: not enough memory in CPU\n"
.LC11:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/iDneprov/university/master/PGP-POD/3/lab3.cu"
.LC12:
	.string	"ERROR in %s:%d. Message: %s\n"
.LC13:
	.string	"%d"
.LC14:
	.string	"%d %d"
.LC16:
	.string	"wb"
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB2906:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA2906
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC7(%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	$840, %rsp
	.cfi_def_cfa_offset 896
	movq	%fs:40, %rax
	movq	%rax, 824(%rsp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	leaq	24(%rsp), %rcx
	leaq	312(%rsp), %rbx
	movq	%rax, 48(%rsp)
	movq	%rbx, %rsi
	movq	%rax, 80(%rsp)
	movq	%rax, 96(%rsp)
	movq	%rax, 88(%rsp)
	movl	$16777217, %eax
	salq	$8, %rax
	movl	$1, 64(%rsp)
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	leaq	568(%rsp), %rax
	movq	%rax, 8(%rsp)
	movq	%rax, %rdx
	xorl	%eax, %eax
	movl	$1, 76(%rsp)
.LEHB0:
	call	__isoc23_scanf@PLT
	leaq	.LC8(%rip), %rsi
	movq	%rbx, %rdi
	call	fopen@PLT
	movq	%rax, %r12
	testq	%rax, %rax
	jne	.L129
	movq	stderr(%rip), %rdi
	movq	%rbx, %rcx
	movl	$2, %esi
	xorl	%eax, %eax
	leaq	.LC9(%rip), %rdx
	call	__fprintf_chk@PLT
	jmp	.L130
.L129:
	leaq	40(%rsp), %rdi
	movq	%rax, %r8
	movl	$1, %ecx
	movl	$4, %edx
	movl	$4, %esi
	call	__fread_chk@PLT
	leaq	44(%rsp), %rdi
	movq	%r12, %r8
	movl	$1, %ecx
	movl	$4, %edx
	movl	$4, %esi
	call	__fread_chk@PLT
	movl	40(%rsp), %ebx
	imull	44(%rsp), %ebx
	movslq	%ebx, %rbx
	salq	$2, %rbx
	movq	%rbx, %rdi
	call	malloc@PLT
	movq	%rax, %rbp
	testq	%rax, %rax
	jne	.L131
	movq	stderr(%rip), %rdi
	leaq	.LC10(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	jmp	.L132
.L131:
	leaq	48(%rsp), %rdi
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L133
	call	cudaGetErrorString@PLT
	movq	%rax, %r9
	movl	$317, %r8d
.L157:
	movq	stderr(%rip), %rdi
	leaq	.LC11(%rip), %rcx
	leaq	.LC12(%rip), %rdx
	xorl	%eax, %eax
	movl	$2, %esi
	call	__fprintf_chk@PLT
	xorl	%edi, %edi
	call	exit@PLT
.L133:
	movl	40(%rsp), %ecx
	imull	44(%rsp), %ecx
	movq	%r12, %r8
	movq	%rbx, %rsi
	movl	$4, %edx
	movq	%rbp, %rdi
	movslq	%ecx, %rcx
	call	__fread_chk@PLT
	movq	%r12, %rdi
	call	fclose@PLT
	leaq	cov(%rip), %rax
	leaq	avg(%rip), %rbx
	xorl	%r14d, %r14d
	movq	%rax, (%rsp)
.L134:
	movslq	24(%rsp), %rdx
	cmpl	%r14d, %edx
	jle	.L161
	leaq	28(%rsp), %rsi
	leaq	.LC13(%rip), %rdi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movq	88(%rsp), %rax
	subq	80(%rsp), %rax
	movabsq	$-6148914691236517205, %rdx
	sarq	$3, %rax
	movslq	28(%rsp), %rsi
	imulq	%rdx, %rax
	cmpq	%rsi, %rax
	jnb	.L135
	subq	%rax, %rsi
	leaq	80(%rsp), %rdi
	call	_ZNSt6vectorI7double3SaIS0_EE17_M_default_appendEm
.L135:
	movq	$0x000000000, (%rbx)
	xorl	%r13d, %r13d
	xorps	%xmm0, %xmm0
	leaq	.LC14(%rip), %r15
	movq	$0x000000000, 8(%rbx)
.L136:
	movl	28(%rsp), %r12d
	movsd	%xmm0, 16(%rbx)
	cmpl	%r13d, %r12d
	jle	.L162
	leaq	36(%rsp), %rdx
	leaq	32(%rsp), %rsi
	movq	%r15, %rdi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movl	36(%rsp), %eax
	imull	40(%rsp), %eax
	addl	32(%rsp), %eax
	cltq
	leaq	0(%rbp,%rax,4), %rax
	movzbl	2(%rax), %edx
	cvtsi2sdl	%edx, %xmm0
	movzbl	1(%rax), %edx
	movzbl	(%rax), %eax
	cvtsi2sdl	%edx, %xmm1
	cvtsi2sdl	%eax, %xmm2
	imulq	$24, %r13, %rax
	addq	80(%rsp), %rax
	incq	%r13
	movsd	%xmm0, 16(%rax)
	movsd	%xmm1, 8(%rax)
	movsd	%xmm2, (%rax)
	addsd	(%rbx), %xmm2
	addsd	8(%rbx), %xmm1
	addsd	16(%rbx), %xmm0
	movsd	%xmm2, (%rbx)
	movsd	%xmm1, 8(%rbx)
	jmp	.L136
.L162:
	cvtsi2sdl	%r12d, %xmm1
	movsd	(%rbx), %xmm2
	divsd	%xmm1, %xmm2
	divsd	%xmm1, %xmm0
	movsd	%xmm2, (%rbx)
	movsd	8(%rbx), %xmm2
	divsd	%xmm1, %xmm2
	movsd	%xmm0, 16(%rbx)
	movsd	%xmm2, 8(%rbx)
	cmpl	$1, %r12d
	jle	.L138
	xorps	%xmm1, %xmm1
	leaq	120(%rsp), %rdi
	movl	$3, %esi
	xorl	%r15d, %r15d
	movaps	%xmm1, %xmm0
	call	_Z7mat_setPdidd
	movsd	.LC2(%rip), %xmm1
.L139:
	imulq	$24, %r15, %rdi
	addq	80(%rsp), %rdi
	incq	%r15
	movl	$3, %r9d
	movsd	(%rdi), %xmm0
	subsd	(%rbx), %xmm0
	leaq	120(%rsp), %r13
	movq	%rdi, %rsi
	movl	$1, %r8d
	movl	$3, %ecx
	movq	%r13, %rdx
	movsd	%xmm0, (%rdi)
	movsd	8(%rdi), %xmm0
	subsd	8(%rbx), %xmm0
	movsd	%xmm0, 8(%rdi)
	movsd	16(%rdi), %xmm0
	subsd	16(%rbx), %xmm0
	movsd	%xmm0, 16(%rdi)
	movaps	%xmm1, %xmm0
	call	_Z7mat_mulPKdS0_Pdiiidd
	cmpl	%r15d, %r12d
	jg	.L139
	leal	-1(%r12), %eax
	movq	%r13, %rsi
	movq	%r13, %rdi
	movl	$3, %ecx
	cvtsi2sdl	%eax, %xmm0
	movl	$3, %edx
	divsd	%xmm0, %xmm1
	movaps	%xmm1, %xmm0
	call	_Z9mat_mul_CPKdPddii
	movsd	.LC15(%rip), %xmm0
	movq	(%rsp), %rsi
	movq	%r13, %rdi
	leaq	192(%rsp), %rcx
	leaq	108(%rsp), %r8
	call	_Z11LUP_inv_matPdS_iS_Pid
	jmp	.L140
.L138:
	movsd	.LC2(%rip), %xmm0
	movq	(%rsp), %rdi
	xorps	%xmm1, %xmm1
	movl	$3, %esi
	call	_Z7mat_setPdidd
.L140:
	addq	$72, (%rsp)
	incl	%r14d
	addq	$24, %rbx
	jmp	.L134
.L161:
	imulq	$24, %rdx, %rdx
	movl	$1, %r8d
	xorl	%ecx, %ecx
	leaq	avg(%rip), %rsi
	leaq	_ZL3AVG(%rip), %rdi
	call	cudaMemcpyToSymbol@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L142
	call	cudaGetErrorString@PLT
	movq	%rax, %r9
	movl	$364, %r8d
	jmp	.L157
.L142:
	imull	$9, 24(%rsp), %edx
	movl	$1, %r8d
	xorl	%ecx, %ecx
	leaq	cov(%rip), %rsi
	leaq	_ZL3COV(%rip), %rdi
	movslq	%edx, %rdx
	salq	$3, %rdx
	call	cudaMemcpyToSymbol@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L144
	call	cudaGetErrorString@PLT
	movq	%rax, %r9
	movl	$365, %r8d
	jmp	.L157
.L144:
	movl	40(%rsp), %edx
	imull	44(%rsp), %edx
	movl	$1, %ecx
	movq	%rbp, %rsi
	movq	48(%rsp), %rdi
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	76(%rsp), %ecx
	movq	68(%rsp), %rdx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	56(%rsp), %rdi
	movl	64(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L145
	movl	44(%rsp), %ecx
	movl	40(%rsp), %edx
	movq	48(%rsp), %rsi
	movl	24(%rsp), %edi
	call	_Z34__device_stub__Z6kerneliP6uchar4iiiP6uchar4ii
.L145:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L146
	call	cudaGetErrorString@PLT
	movq	%rax, %r9
	movl	$371, %r8d
	jmp	.L157
.L146:
	movl	40(%rsp), %edx
	imull	44(%rsp), %edx
	movl	$2, %ecx
	movq	%rbp, %rdi
	movq	48(%rsp), %rsi
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L147
	call	cudaGetErrorString@PLT
	movq	%rax, %r9
	movl	$372, %r8d
	jmp	.L157
.L147:
	movq	8(%rsp), %rdi
	leaq	.LC16(%rip), %rsi
	call	fopen@PLT
	movq	%rax, %rbx
	testq	%rax, %rax
	jne	.L148
	movq	8(%rsp), %rcx
	movq	stderr(%rip), %rdi
	leaq	.LC9(%rip), %rdx
	xorl	%eax, %eax
	movl	$2, %esi
	call	__fprintf_chk@PLT
	jmp	.L132
.L148:
	leaq	40(%rsp), %rdi
	movq	%rax, %rcx
	movl	$1, %edx
	movl	$4, %esi
	call	fwrite@PLT
	leaq	44(%rsp), %rdi
	movq	%rbx, %rcx
	movl	$1, %edx
	movl	$4, %esi
	call	fwrite@PLT
	movl	40(%rsp), %edx
	imull	44(%rsp), %edx
	movq	%rbx, %rcx
	movq	%rbp, %rdi
	movl	$4, %esi
	movslq	%edx, %rdx
	call	fwrite@PLT
	movq	%rbx, %rdi
	call	fclose@PLT
.L132:
	movq	%rbp, %rdi
	call	free@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
.LEHE0:
.L130:
	leaq	80(%rsp), %rdi
	call	_ZNSt6vectorI7double3SaIS0_EED1Ev
	movq	824(%rsp), %rax
	subq	%fs:40, %rax
	je	.L151
	jmp	.L154
.L152:
	endbr64
	movq	%rax, %rbx
.L149:
	leaq	80(%rsp), %rdi
	call	_ZNSt6vectorI7double3SaIS0_EED1Ev
	movq	824(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L154
	movq	%rbx, %rdi
.LEHB1:
	call	_Unwind_Resume@PLT
.LEHE1:
.L154:
	call	__stack_chk_fail@PLT
.L151:
	addq	$840, %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
.LFE2906:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA2906:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE2906-.LLSDACSB2906
.LLSDACSB2906:
	.uleb128 .LEHB0-.LFB2906
	.uleb128 .LEHE0-.LEHB0
	.uleb128 .L152-.LFB2906
	.uleb128 0
	.uleb128 .LEHB1-.LFB2906
	.uleb128 .LEHE1-.LEHB1
	.uleb128 0
	.uleb128 0
.LLSDACSE2906:
	.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
	.local	_ZL3COV
	.comm	_ZL3COV,2304,32
	.local	_ZL3AVG
	.comm	_ZL3AVG,768,32
	.globl	cov
	.bss
	.align 32
	.type	cov, @object
	.size	cov, 2304
cov:
	.zero	2304
	.globl	avg
	.align 32
	.type	avg, @object
	.size	avg, 768
avg:
	.zero	768
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC1:
	.long	-1
	.long	2147483647
	.long	0
	.long	0
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC2:
	.long	0
	.long	1072693248
	.align 8
.LC15:
	.long	-640172613
	.long	1037794527
	.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 "hip/hip_runtime.h"


#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <string.h>
#include <vector>

#define CSC(call)                                   \
do {                                                 \
	hipError_t res = call;                             \
	if (res != hipSuccess) {                            \
		fprintf(stderr, "ERROR in %s:%d. Message: %s\n",    \
			__FILE__, __LINE__, hipGetErrorString(res));      \
		exit(0);                                              \
	}                                                        \
} while(0)

typedef unsigned char uchar;

#define NC_MAX 32 // максимальное количество классов

double3 avg[NC_MAX]; // средние
double cov[3 * 3 * NC_MAX]; // ковариационные матрицы
__constant__ double3 AVG[NC_MAX];
__constant__ double COV[3 * 3 * NC_MAX];

__host__ __device__
double dot(const double *A, const double *a, const double *b, int n)
{
    /*
    * res = a ^ T * A * b
    */
    int i, j;
    double res = 0;
    for (i = 0; i < n; i++)
    {
        for (j = 0; j < n; j++)
        {
            res += A[i * n + j] * a[i] * b[j];
        }
    }
    return res;
}

__host__ __device__ // работает везде
void mat_mul(const double *A, const double *B, double *C, int m, int n, int l, double alpha=1, double beta=0)
{
    /*
    * C = alpha * A * B + beta * C
    *
    * m - число строк A
    * n - число столбцов A и число строк B
    * l - число столбцов B
    */
    int i, j, k;
    double dot;

    for (i = 0; i < m; i++)
    {
        for (j = 0; j < l; j++)
        {
            dot = 0;
            for (k = 0; k < n; k++)
            {
                dot += A[i * n + k] * B[k * l + j];
            }
            C[i * l + j] = alpha * dot + (beta == 0 ? 0 : beta * C[i * l + j]);
        }
    }
}

__host__ __device__
void mat_mul_C(const double *A, double *B, double c, int m, int n)
{
    /*
    * B = c * A
    */
    int i, j;
    for (i = 0; i < m; i++)
    {
        for (j = 0; j < n; j++)
        {
            B[i * n + j] = c * A[i * n + j];
        }
    }
}

__host__ __device__
void mat_sum(const double *A, const double *B, double *C, int m, int n, double alpha=1, double beta=1)
{
    /*
    * C = alpha * A + beta * B
    */
    int i, j;
    for (i = 0; i < m; i++)
    {
        for (j = 0; j < n; j++)
        {
            C[i * n + j] = alpha * A[i * n + j] + beta * B[i * n + j];
        }
    }
}

__host__ __device__
void mat_set(double *A, int n, double alpha=0, double beta=0)
{
    /*
    * Инициализация A значениями alpha на диагонали и
    * значениями beta вне диагонали
    */
    int i, j;
    for (i = 0; i < n; i++)
    {
        for (j = 0; j < n; j++)
        {
            if (i == j)
            {
                A[i * n + j] = alpha;
            }
            else
            {
                A[i * n + j] = beta;
            }
        }
    }
}

void mat_swap_rows(double *A, int n, int i1, int i2)
{
    int j;
    double tmp;
    for (j = 0; j < n; j++)
    {
        tmp = A[i1 * n + j];
        A[i1 * n + j] = A[i2 * n + j];
        A[i2 * n + j] = tmp;
    }
}

void mat_transpose(const double *A, double *A_t, int n)
{
    int i, j;
    for (i = 0; i < n; i++)
    {
        for (j = 0; j < n; j++)
        {
            A_t[i * n + j] = A[j * n + i];
        }
    }
}

bool LUP(double *A, int n, int *pi, double eps=1e-10)
{
    int i, j, k, k_;
    int tmp;

    for (i = 0; i < n; i++)
    {
        pi[i] = i;
    }

    for (k = 0; k < n; k++)
    {
        k_ = k;
        for (i = k + 1; i < n; i++)
        {
            if (fabs(A[i * n + k]) > fabs(A[k_ * n + k]))
            {
                k_ = i;
            }
        }

        if (fabs(A[k_ * n + k]) < eps)
        {
            return false;
        }

        if (k != k_)
        {
            tmp = pi[k];
            pi[k] = pi[k_];
            pi[k_] = tmp;
            mat_swap_rows(A, n, k, k_);
        }

        for (i = k + 1; i < n; i++)
        {
            A[i * n + k] /= A[k * n + k];
            for (j = k + 1; j < n; j++)
            {
                A[i * n + j] -= A[i * n + k] * A[k * n + j];
            }
        }
    }
    return true;
}

void LUP_solve(const double *LU, const int *pi, const double *b, int n, double *x, double *work)
{
    /*
    * work - вектор размерности n
    */
    int i, j;
    double sum;
    double *y = work;

    memset(y, 0, n * sizeof(double));

    for (i = 0; i < n; i++) // прямой ход
    {
        for (sum = 0, j = 0; j <= i - 1; sum += LU[i * n + j] * y[j], j++);
        y[i] = b[pi[i]] - sum;
    }

    for (i = n - 1; i >= 0; i--) // обратный ход
    {
        for (sum = 0, j = i + 1; j < n; sum += LU[i * n + j] * x[j], j++);
        x[i] = (y[i] - sum) / LU[i * n + i];
    }
}

void LUP_inv_mat(double *A, double *A_inv, int n, double *work, int *iwork, double eps=1e-10) // инвариантная матрица при помощи LUP разложения (украдено из Кормана)
{
    /*
    * work  - вектор длины n ^ 2 + 2 * n
    * iwork - вектор длины n
    */
    int i;
    double *X, *e, *space;
    int *pi;

    X = work;
    e = X + n * n;
    space = e + n;

    pi = iwork;

    memset(e, 0, n * sizeof(double));
    e[0] = 1;

    LUP(A, n, pi, eps);
    for (i = 0; i < n - 1; i++)
    {
        LUP_solve(A, pi, e, n, X + i * n, space);
        e[i] = 0;
        e[i + 1] = 1;
    }
    LUP_solve(A, pi, e, n, X + i * n, space);
    mat_transpose(X, A_inv, n);
}

__global__
void kernel(int nc, uchar4 *im, int w, int h)
{
    int i, j, jc, idx;
    int offset;
    double3 to_dot;
    double dist, dist_max;

    idx = blockDim.x * blockIdx.x + threadIdx.x; // абсолютный номер потока
    offset = blockDim.x * gridDim.x; // общее кол-во потоков

    for (i = idx; i < w * h; i += offset) // идем по всем пикселям
    {
        dist_max = -INFINITY;
        jc = 0;
        for (j = 0; j < nc; j++) // цикл по числу классов
        {

            to_dot = make_double3((double)im[i].x - AVG[j].x,
								(double)im[i].y - AVG[j].y,
								(double)im[i].z - AVG[j].z);
            dist = -dot(COV + j * 3 * 3, (double *)&to_dot, (double *)&to_dot, 3); // произведение из основной формулы
            if (dist > dist_max) // для каждого класса вычисляем dist и выбираем максимум
            {
                jc = j;
                dist_max = dist;
            }
        }
        im[i].w = jc; // записываем номер класса с max dist
    }
}

int main()
{
    int nc, np, i, j, x, y, w, h;
    uchar4 *im = NULL, *im_dev = NULL;
    std::vector<double3> v(0);
    double mat[3 * 3],
           work[3 * 3 + 2 * 3];
    int iwork[3];

    FILE *fp;
    char name_src_im[256], name_dst_im[256];

    dim3 blocks(256), threads(256);

    scanf("%s\n%s\n%d", name_src_im, name_dst_im, &nc);

    fp = fopen(name_src_im, "rb");
    if (fp == NULL)
    {
        fprintf(stderr, "Error: can't open %s\n", name_src_im);
        return 0;
    }
    fread(&w, sizeof(int), 1, fp);
    fread(&h, sizeof(int), 1, fp);

    im = (uchar4 *)malloc(w * h * sizeof(uchar4));
    if (im == NULL)
    {
        fprintf(stderr, "Error: not enough memory in CPU\n");
        goto FREE;
    }
    CSC(hipMalloc(&im_dev, w * h * sizeof(uchar4)));

    fread(im, sizeof(uchar4), w * h, fp); // считывание пикселей
    fclose(fp);

    for (j = 0; j < nc; j++) // цикл по числу классов
    {
        scanf("%d", &np); // количество пикселей в классе
        if(v.size() < np) v.resize(np);

        avg[j] = make_double3(0, 0, 0);
        for (i = 0; i < np; i++) // просчёт среднего по j-му классу
        {
            scanf("%d %d", &x, &y);
            v[i] = make_double3((double)im[y * w + x].x,
								(double)im[y * w + x].y,
								(double)im[y * w + x].z);

            avg[j].x += v[i].x;
            avg[j].y += v[i].y;
            avg[j].z += v[i].z;
        }
        avg[j].x /= np;
        avg[j].y /= np;
        avg[j].z /= np;

        if (np > 1)
        {
            mat_set(mat, 3, 0, 0);          // инициализация нулями
            for (i = 0; i < np; i++)        // просчёт ковариации по i-му классу
            {
                v[i].x = v[i].x - avg[j].x;
                v[i].y = v[i].y - avg[j].y;
                v[i].z = v[i].z - avg[j].z;
								// double3 приводим к вектору double
                mat_mul((double *)(v.data() + i), (double *)(v.data() + i), mat, 3, 1, 3, 1, 1); // умножение векторов как матриц 3х1 1х3 получаем 3х3
            }
            mat_mul_C(mat, mat, 1. / (np - 1), 3, 3); // делим матрицу на коэффициент
            LUP_inv_mat(mat, cov + j * 3 * 3, 3, work, iwork); // получаем матрицу ковариации для j-того класса (work, iwork -- место, чтобы внутри функции не делать malloc)
        }
        else
        {
            mat_set(cov + j * 3 * 3, 3, 1, 0); // единичная матрица
        }
    }

    /* Копирование в константную память */
    CSC(hipMemcpyToSymbol(HIP_SYMBOL(AVG), avg, nc * sizeof(double3), 0, hipMemcpyHostToDevice));
    CSC(hipMemcpyToSymbol(HIP_SYMBOL(COV), cov, nc * 3 * 3 * sizeof(double), 0, hipMemcpyHostToDevice));

    /* Копирование изображения */
    hipMemcpy(im_dev, im, w * h * sizeof(uchar4), hipMemcpyHostToDevice); // копирование изображения

    kernel<<<blocks, threads>>>(nc, im_dev, w, h);
    CSC(hipGetLastError());
    CSC(hipMemcpy(im, im_dev, w * h * sizeof(uchar4), hipMemcpyDeviceToHost));

    fp = fopen(name_dst_im, "wb");
    if (fp == NULL)
    {
        fprintf(stderr, "Error: can't open %s\n", name_dst_im);
        goto FREE;
    }
    fwrite(&w, sizeof(int), 1, fp);
    fwrite(&h, sizeof(int), 1, fp);
    fwrite(im, sizeof(uchar4), w * h, fp);
    fclose(fp);

FREE:
    free(im);
    hipFree(im_dev);
    return 0;
}

	.text
	.file	"lab3.hip"
	.globl	_Z3dotPKdS0_S0_i                # -- Begin function _Z3dotPKdS0_S0_i
	.type	_Z3dotPKdS0_S0_i,@function
_Z3dotPKdS0_S0_i:                       # @_Z3dotPKdS0_S0_i
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB0_1
# %bb.2:                                # %.preheader.lr.ph
	movl	%ecx, %eax
	xorpd	%xmm0, %xmm0
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
.LBB0_3:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_4 Depth 2
	movl	%r8d, %r10d
	leaq	(%rdi,%r10,8), %r10
	movsd	(%rsi,%r9,8), %xmm1             # xmm1 = mem[0],zero
	xorl	%r11d, %r11d
.LBB0_4:                                #   Parent Loop BB0_3 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%r10,%r11,8), %xmm2            # xmm2 = mem[0],zero
	mulsd	%xmm1, %xmm2
	mulsd	(%rdx,%r11,8), %xmm2
	addsd	%xmm2, %xmm0
	incq	%r11
	cmpq	%r11, %rax
	jne	.LBB0_4
# %bb.5:                                # %._crit_edge
                                        #   in Loop: Header=BB0_3 Depth=1
	incq	%r9
	addl	%ecx, %r8d
	cmpq	%rax, %r9
	jne	.LBB0_3
# %bb.6:                                # %._crit_edge22
	retq
.LBB0_1:
	xorps	%xmm0, %xmm0
	retq
.Lfunc_end0:
	.size	_Z3dotPKdS0_S0_i, .Lfunc_end0-_Z3dotPKdS0_S0_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z7mat_mulPKdS0_Pdiiidd         # -- Begin function _Z7mat_mulPKdS0_Pdiiidd
	.type	_Z7mat_mulPKdS0_Pdiiidd,@function
_Z7mat_mulPKdS0_Pdiiidd:                # @_Z7mat_mulPKdS0_Pdiiidd
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, -8(%rsp)                  # 8-byte Spill
	movq	%rdi, -16(%rsp)                 # 8-byte Spill
	testl	%ecx, %ecx
	jle	.LBB1_11
# %bb.1:                                # %.preheader35.lr.ph
	movslq	%r9d, %rax
	movl	%ecx, %ecx
	movl	%eax, %r10d
	movl	%r8d, %r11d
	movq	%rax, -24(%rsp)                 # 8-byte Spill
	leaq	(,%rax,8), %rbx
	xorl	%r14d, %r14d
	xorl	%r15d, %r15d
.LBB1_2:                                # %.preheader35
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_4 Depth 2
                                        #       Child Loop BB1_6 Depth 3
	testl	%r9d, %r9d
	jle	.LBB1_10
# %bb.3:                                # %.preheader.lr.ph
                                        #   in Loop: Header=BB1_2 Depth=1
	movl	%r14d, %eax
	movq	-16(%rsp), %rsi                 # 8-byte Reload
	leaq	(%rsi,%rax,8), %r12
	movq	%r15, %r13
	imulq	-24(%rsp), %r13                 # 8-byte Folded Reload
	movq	-8(%rsp), %rsi                  # 8-byte Reload
	xorl	%eax, %eax
.LBB1_4:                                # %.preheader
                                        #   Parent Loop BB1_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB1_6 Depth 3
	xorpd	%xmm2, %xmm2
	xorpd	%xmm3, %xmm3
	testl	%r8d, %r8d
	jle	.LBB1_7
# %bb.5:                                # %.lr.ph.preheader
                                        #   in Loop: Header=BB1_4 Depth=2
	movq	%rsi, %rbp
	xorl	%edi, %edi
.LBB1_6:                                # %.lr.ph
                                        #   Parent Loop BB1_2 Depth=1
                                        #     Parent Loop BB1_4 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movsd	(%r12,%rdi,8), %xmm4            # xmm4 = mem[0],zero
	mulsd	(%rbp), %xmm4
	addsd	%xmm4, %xmm3
	incq	%rdi
	addq	%rbx, %rbp
	cmpq	%rdi, %r11
	jne	.LBB1_6
.LBB1_7:                                # %._crit_edge
                                        #   in Loop: Header=BB1_4 Depth=2
	leaq	(%rax,%r13), %rdi
	ucomisd	%xmm2, %xmm1
	jne	.LBB1_8
	jnp	.LBB1_9
.LBB1_8:                                #   in Loop: Header=BB1_4 Depth=2
	movsd	(%rdx,%rdi,8), %xmm2            # xmm2 = mem[0],zero
	mulsd	%xmm1, %xmm2
.LBB1_9:                                # %._crit_edge._crit_edge
                                        #   in Loop: Header=BB1_4 Depth=2
	mulsd	%xmm0, %xmm3
	addsd	%xmm2, %xmm3
	movsd	%xmm3, (%rdx,%rdi,8)
	incq	%rax
	addq	$8, %rsi
	cmpq	%r10, %rax
	jne	.LBB1_4
.LBB1_10:                               # %._crit_edge39
                                        #   in Loop: Header=BB1_2 Depth=1
	incq	%r15
	addl	%r8d, %r14d
	cmpq	%rcx, %r15
	jne	.LBB1_2
.LBB1_11:                               # %._crit_edge41
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z7mat_mulPKdS0_Pdiiidd, .Lfunc_end1-_Z7mat_mulPKdS0_Pdiiidd
	.cfi_endproc
                                        # -- End function
	.globl	_Z9mat_mul_CPKdPddii            # -- Begin function _Z9mat_mul_CPKdPddii
	.type	_Z9mat_mul_CPKdPddii,@function
_Z9mat_mul_CPKdPddii:                   # @_Z9mat_mul_CPKdPddii
	.cfi_startproc
# %bb.0:
	testl	%edx, %edx
	jle	.LBB2_7
# %bb.1:                                # %.preheader.lr.ph
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movl	%edx, %eax
	movl	%ecx, %edx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
.LBB2_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_4 Depth 2
	testl	%ecx, %ecx
	jle	.LBB2_5
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB2_2 Depth=1
	movl	%r8d, %r11d
	leaq	(%rsi,%r11,8), %r10
	leaq	(%rdi,%r11,8), %r11
	xorl	%ebx, %ebx
.LBB2_4:                                #   Parent Loop BB2_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%r11,%rbx,8), %xmm1            # xmm1 = mem[0],zero
	mulsd	%xmm0, %xmm1
	movsd	%xmm1, (%r10,%rbx,8)
	incq	%rbx
	cmpq	%rbx, %rdx
	jne	.LBB2_4
.LBB2_5:                                # %._crit_edge
                                        #   in Loop: Header=BB2_2 Depth=1
	incq	%r9
	addl	%ecx, %r8d
	cmpq	%rax, %r9
	jne	.LBB2_2
# %bb.6:
	popq	%rbx
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
.LBB2_7:                                # %._crit_edge18
	retq
.Lfunc_end2:
	.size	_Z9mat_mul_CPKdPddii, .Lfunc_end2-_Z9mat_mul_CPKdPddii
	.cfi_endproc
                                        # -- End function
	.globl	_Z7mat_sumPKdS0_Pdiidd          # -- Begin function _Z7mat_sumPKdS0_Pdiidd
	.type	_Z7mat_sumPKdS0_Pdiidd,@function
_Z7mat_sumPKdS0_Pdiidd:                 # @_Z7mat_sumPKdS0_Pdiidd
	.cfi_startproc
# %bb.0:
	testl	%ecx, %ecx
	jle	.LBB3_7
# %bb.1:                                # %.preheader.lr.ph
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	%ecx, %eax
	movl	%r8d, %ecx
	xorl	%r9d, %r9d
	xorl	%r10d, %r10d
.LBB3_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_4 Depth 2
	testl	%r8d, %r8d
	jle	.LBB3_5
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB3_2 Depth=1
	movl	%r9d, %r14d
	leaq	(%rdx,%r14,8), %r11
	leaq	(%rsi,%r14,8), %rbx
	leaq	(%rdi,%r14,8), %r14
	xorl	%r15d, %r15d
.LBB3_4:                                #   Parent Loop BB3_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%r14,%r15,8), %xmm2            # xmm2 = mem[0],zero
	mulsd	%xmm0, %xmm2
	movsd	(%rbx,%r15,8), %xmm3            # xmm3 = mem[0],zero
	mulsd	%xmm1, %xmm3
	addsd	%xmm2, %xmm3
	movsd	%xmm3, (%r11,%r15,8)
	incq	%r15
	cmpq	%r15, %rcx
	jne	.LBB3_4
.LBB3_5:                                # %._crit_edge
                                        #   in Loop: Header=BB3_2 Depth=1
	incq	%r10
	addl	%r8d, %r9d
	cmpq	%rax, %r10
	jne	.LBB3_2
# %bb.6:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %r15
.LBB3_7:                                # %._crit_edge23
	retq
.Lfunc_end3:
	.size	_Z7mat_sumPKdS0_Pdiidd, .Lfunc_end3-_Z7mat_sumPKdS0_Pdiidd
	.cfi_endproc
                                        # -- End function
	.globl	_Z7mat_setPdidd                 # -- Begin function _Z7mat_setPdidd
	.type	_Z7mat_setPdidd,@function
_Z7mat_setPdidd:                        # @_Z7mat_setPdidd
	.cfi_startproc
# %bb.0:
	testl	%esi, %esi
	jle	.LBB4_8
# %bb.1:                                # %.preheader.lr.ph
	movl	%esi, %eax
	xorl	%ecx, %ecx
	xorl	%edx, %edx
.LBB4_2:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_3 Depth 2
	movl	%ecx, %r8d
	leaq	(%rdi,%r8,8), %r8
	movl	%edx, %r9d
	imull	%esi, %r9d
	leaq	(%rdi,%rdx,8), %r10
	xorl	%r11d, %r11d
.LBB4_3:                                #   Parent Loop BB4_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	cmpq	%r11, %rdx
	jne	.LBB4_5
# %bb.4:                                #   in Loop: Header=BB4_3 Depth=2
	movsd	%xmm0, (%r10,%r9,8)
	jmp	.LBB4_6
.LBB4_5:                                #   in Loop: Header=BB4_3 Depth=2
	movsd	%xmm1, (%r8,%r11,8)
.LBB4_6:                                #   in Loop: Header=BB4_3 Depth=2
	incq	%r11
	cmpq	%r11, %rax
	jne	.LBB4_3
# %bb.7:                                # %._crit_edge
                                        #   in Loop: Header=BB4_2 Depth=1
	incq	%rdx
	addl	%esi, %ecx
	cmpq	%rax, %rdx
	jne	.LBB4_2
.LBB4_8:                                # %._crit_edge21
	retq
.Lfunc_end4:
	.size	_Z7mat_setPdidd, .Lfunc_end4-_Z7mat_setPdidd
	.cfi_endproc
                                        # -- End function
	.globl	_Z13mat_swap_rowsPdiii          # -- Begin function _Z13mat_swap_rowsPdiii
	.type	_Z13mat_swap_rowsPdiii,@function
_Z13mat_swap_rowsPdiii:                 # @_Z13mat_swap_rowsPdiii
	.cfi_startproc
# %bb.0:
	testl	%esi, %esi
	jle	.LBB5_3
# %bb.1:                                # %.lr.ph
	imull	%esi, %edx
	imull	%esi, %ecx
	movslq	%edx, %rdx
	movslq	%ecx, %r8
	movl	%esi, %eax
	leaq	(%rdi,%rdx,8), %rcx
	leaq	(%rdi,%r8,8), %rdx
	xorl	%esi, %esi
.LBB5_2:                                # =>This Inner Loop Header: Depth=1
	movsd	(%rcx,%rsi,8), %xmm0            # xmm0 = mem[0],zero
	movsd	(%rdx,%rsi,8), %xmm1            # xmm1 = mem[0],zero
	movsd	%xmm1, (%rcx,%rsi,8)
	movsd	%xmm0, (%rdx,%rsi,8)
	incq	%rsi
	cmpq	%rsi, %rax
	jne	.LBB5_2
.LBB5_3:                                # %._crit_edge
	retq
.Lfunc_end5:
	.size	_Z13mat_swap_rowsPdiii, .Lfunc_end5-_Z13mat_swap_rowsPdiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z13mat_transposePKdPdi         # -- Begin function _Z13mat_transposePKdPdi
	.type	_Z13mat_transposePKdPdi,@function
_Z13mat_transposePKdPdi:                # @_Z13mat_transposePKdPdi
	.cfi_startproc
# %bb.0:
	testl	%edx, %edx
	jle	.LBB6_6
# %bb.1:                                # %.preheader.lr.ph
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movl	%edx, %eax
	leaq	(,%rax,8), %rcx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
.LBB6_2:                                # %.lr.ph
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_3 Depth 2
	movl	%r8d, %r10d
	leaq	(%rsi,%r10,8), %r10
	movq	%rdi, %r11
	xorl	%ebx, %ebx
.LBB6_3:                                #   Parent Loop BB6_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%r11), %xmm0                   # xmm0 = mem[0],zero
	movsd	%xmm0, (%r10,%rbx,8)
	incq	%rbx
	addq	%rcx, %r11
	cmpq	%rbx, %rax
	jne	.LBB6_3
# %bb.4:                                # %._crit_edge
                                        #   in Loop: Header=BB6_2 Depth=1
	incq	%r9
	addl	%edx, %r8d
	addq	$8, %rdi
	cmpq	%rax, %r9
	jne	.LBB6_2
# %bb.5:
	popq	%rbx
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
.LBB6_6:                                # %._crit_edge17
	retq
.Lfunc_end6:
	.size	_Z13mat_transposePKdPdi, .Lfunc_end6-_Z13mat_transposePKdPdi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function _Z3LUPPdiPid
.LCPI7_0:
	.quad	0x7fffffffffffffff              # double NaN
	.quad	0x7fffffffffffffff              # double NaN
	.text
	.globl	_Z3LUPPdiPid
	.type	_Z3LUPPdiPid,@function
_Z3LUPPdiPid:                           # @_Z3LUPPdiPid
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	testl	%esi, %esi
	jle	.LBB7_1
# %bb.3:                                # %.lr.ph.preheader
	movl	%esi, %ecx
	xorl	%eax, %eax
.LBB7_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	%eax, (%rdx,%rax,4)
	incq	%rax
	cmpq	%rax, %rcx
	jne	.LBB7_4
# %bb.5:                                # %.lr.ph88
	movq	%rdx, -24(%rsp)                 # 8-byte Spill
	leaq	(,%rcx,8), %r8
	leaq	(%rdi,%rcx,8), %r13
	leaq	8(,%rcx,8), %rax
	movq	%rax, -16(%rsp)                 # 8-byte Spill
	movl	$1, %eax
	movq	%rax, -48(%rsp)                 # 8-byte Spill
	xorl	%eax, %eax
	movq	%rax, -40(%rsp)                 # 8-byte Spill
	movapd	.LCPI7_0(%rip), %xmm1           # xmm1 = [NaN,NaN]
	movapd	.LCPI7_0(%rip), %xmm2           # xmm2 = [NaN,NaN]
	movq	%rcx, %r14
	movl	%esi, -60(%rsp)                 # 4-byte Spill
	movq	%rdi, %r15
	movq	%rdi, %r12
	xorl	%r10d, %r10d
	xorl	%eax, %eax
	movq	%rax, -56(%rsp)                 # 8-byte Spill
	movl	$0, -64(%rsp)                   # 4-byte Folded Spill
.LBB7_7:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB7_9 Depth 2
                                        #     Child Loop BB7_15 Depth 2
                                        #     Child Loop BB7_18 Depth 2
                                        #       Child Loop BB7_19 Depth 3
	leaq	1(%r10), %rdx
	leaq	(%rdi,%r10,8), %r9
	movl	%r10d, %eax
	movq	%rdx, -32(%rsp)                 # 8-byte Spill
	cmpq	%rcx, %rdx
	jae	.LBB7_12
# %bb.8:                                # %.lr.ph80.preheader
                                        #   in Loop: Header=BB7_7 Depth=1
	movq	%r13, %r11
	movq	-48(%rsp), %rbx                 # 8-byte Reload
	movl	%r10d, %ebp
.LBB7_9:                                # %.lr.ph80
                                        #   Parent Loop BB7_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%r11), %xmm3                   # xmm3 = mem[0],zero
	andpd	%xmm1, %xmm3
	movl	%ebp, %eax
	imull	%esi, %eax
	cltq
	movsd	(%r9,%rax,8), %xmm4             # xmm4 = mem[0],zero
	andpd	%xmm1, %xmm4
	ucomisd	%xmm4, %xmm3
	movl	%ebx, %eax
	ja	.LBB7_11
# %bb.10:                               # %.lr.ph80
                                        #   in Loop: Header=BB7_9 Depth=2
	movl	%ebp, %eax
.LBB7_11:                               # %.lr.ph80
                                        #   in Loop: Header=BB7_9 Depth=2
	incq	%rbx
	addq	%r8, %r11
	movl	%eax, %ebp
	cmpq	%rbx, %rcx
	jne	.LBB7_9
.LBB7_12:                               # %._crit_edge
                                        #   in Loop: Header=BB7_7 Depth=1
	movl	%eax, %edx
	imull	%esi, %edx
	movslq	%edx, %rdx
	movsd	(%r9,%rdx,8), %xmm3             # xmm3 = mem[0],zero
	andpd	%xmm2, %xmm3
	ucomisd	%xmm3, %xmm0
	movq	-32(%rsp), %rbp                 # 8-byte Reload
	ja	.LBB7_2
# %bb.13:                               #   in Loop: Header=BB7_7 Depth=1
	movq	%r13, -8(%rsp)                  # 8-byte Spill
	movl	%eax, %r11d
	cmpq	%r11, %r10
	je	.LBB7_16
# %bb.14:                               # %.lr.ph.i
                                        #   in Loop: Header=BB7_7 Depth=1
	movq	-24(%rsp), %r13                 # 8-byte Reload
	movl	(%r13,%r10,4), %r11d
	cltq
	movl	(%r13,%rax,4), %ebx
	movl	%ebx, (%r13,%r10,4)
	movl	%r11d, (%r13,%rax,4)
	leaq	(%rdi,%rdx,8), %rax
	xorl	%edx, %edx
.LBB7_15:                               #   Parent Loop BB7_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%r12,%rdx,8), %xmm3            # xmm3 = mem[0],zero
	movsd	(%rax,%rdx,8), %xmm4            # xmm4 = mem[0],zero
	movsd	%xmm4, (%r12,%rdx,8)
	movsd	%xmm3, (%rax,%rdx,8)
	incq	%rdx
	cmpq	%rdx, %rcx
	jne	.LBB7_15
.LBB7_16:                               # %_Z13mat_swap_rowsPdiii.exit
                                        #   in Loop: Header=BB7_7 Depth=1
	cmpq	%rcx, %rbp
	jae	.LBB7_6
# %bb.17:                               # %.lr.ph86
                                        #   in Loop: Header=BB7_7 Depth=1
	movl	-40(%rsp), %eax                 # 4-byte Reload
	leaq	(%r15,%rax,8), %rax
	movl	-64(%rsp), %edx                 # 4-byte Reload
	movl	%edx, %ebx
	imull	%esi, %ebx
	addl	%r10d, %ebx
	movl	-60(%rsp), %r10d                # 4-byte Reload
	movq	-48(%rsp), %rdx                 # 8-byte Reload
.LBB7_18:                               # %.lr.ph83
                                        #   Parent Loop BB7_7 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB7_19 Depth 3
	movl	%r10d, %r11d
	leaq	(%r15,%r11,8), %rbp
	movq	%rdx, %r11
	imulq	%rcx, %r11
	movsd	(%r9,%r11,8), %xmm3             # xmm3 = mem[0],zero
	divsd	(%rdi,%rbx,8), %xmm3
	movsd	%xmm3, (%r9,%r11,8)
	movl	$1, %r13d
.LBB7_19:                               #   Parent Loop BB7_7 Depth=1
                                        #     Parent Loop BB7_18 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movsd	(%r9,%r11,8), %xmm3             # xmm3 = mem[0],zero
	mulsd	(%rax,%r13,8), %xmm3
	movsd	(%rbp,%r13,8), %xmm4            # xmm4 = mem[0],zero
	subsd	%xmm3, %xmm4
	movsd	%xmm4, (%rbp,%r13,8)
	incq	%r13
	cmpq	%r13, %r14
	jne	.LBB7_19
# %bb.20:                               # %._crit_edge84
                                        #   in Loop: Header=BB7_18 Depth=2
	incq	%rdx
	addl	%esi, %r10d
	cmpq	%rcx, %rdx
	jne	.LBB7_18
.LBB7_6:                                # %.loopexit
                                        #   in Loop: Header=BB7_7 Depth=1
	incl	-64(%rsp)                       # 4-byte Folded Spill
	movq	-32(%rsp), %r10                 # 8-byte Reload
	cmpq	%rcx, %r10
	setae	%al
	movq	%rax, -56(%rsp)                 # 8-byte Spill
	incq	-48(%rsp)                       # 8-byte Folded Spill
	movq	-8(%rsp), %r13                  # 8-byte Reload
	addq	-16(%rsp), %r13                 # 8-byte Folded Reload
	addq	%r8, %r12
	addq	$8, %r15
	addl	%esi, -60(%rsp)                 # 4-byte Folded Spill
	movq	-40(%rsp), %rax                 # 8-byte Reload
	addl	%esi, %eax
	movq	%rax, -40(%rsp)                 # 8-byte Spill
	decq	%r14
	cmpq	%rcx, %r10
	jne	.LBB7_7
	jmp	.LBB7_2
.LBB7_1:
	movb	$1, %al
	movq	%rax, -56(%rsp)                 # 8-byte Spill
.LBB7_2:                                # %._crit_edge89
	movq	-56(%rsp), %rax                 # 8-byte Reload
	andb	$1, %al
                                        # kill: def $al killed $al killed $rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end7:
	.size	_Z3LUPPdiPid, .Lfunc_end7-_Z3LUPPdiPid
	.cfi_endproc
                                        # -- End function
	.globl	_Z9LUP_solvePKdPKiS0_iPdS3_     # -- Begin function _Z9LUP_solvePKdPKiS0_iPdS3_
	.type	_Z9LUP_solvePKdPKiS0_iPdS3_,@function
_Z9LUP_solvePKdPKiS0_iPdS3_:            # @_Z9LUP_solvePKdPKiS0_iPdS3_
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %rbx
	movq	%r8, 16(%rsp)                   # 8-byte Spill
	movl	%ecx, %ebp
	movq	%rdx, %r13
	movq	%rsi, %r14
	movq	%rdi, %r15
	movslq	%ecx, %rax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	leaq	(,%rax,8), %rdx
	movq	%rdx, (%rsp)                    # 8-byte Spill
	xorl	%r12d, %r12d
	movq	%r9, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	movq	(%rsp), %r11                    # 8-byte Reload
	movq	8(%rsp), %r10                   # 8-byte Reload
	movq	16(%rsp), %r9                   # 8-byte Reload
	testl	%r10d, %r10d
	jle	.LBB8_10
# %bb.1:                                # %.preheader52.preheader
	movl	%ebp, %eax
	xorl	%ecx, %ecx
.LBB8_2:                                # %.preheader52
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB8_12 Depth 2
	testq	%rcx, %rcx
	je	.LBB8_3
# %bb.11:                               # %.lr.ph
                                        #   in Loop: Header=BB8_2 Depth=1
	movl	%r12d, %edx
	leaq	(%r15,%rdx,8), %rdx
	xorpd	%xmm0, %xmm0
	xorl	%esi, %esi
.LBB8_12:                               #   Parent Loop BB8_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rdx,%rsi,8), %xmm1            # xmm1 = mem[0],zero
	mulsd	(%rbx,%rsi,8), %xmm1
	addsd	%xmm1, %xmm0
	incq	%rsi
	cmpq	%rsi, %rcx
	jne	.LBB8_12
	jmp	.LBB8_13
.LBB8_3:                                #   in Loop: Header=BB8_2 Depth=1
	xorpd	%xmm0, %xmm0
.LBB8_13:                               # %._crit_edge
                                        #   in Loop: Header=BB8_2 Depth=1
	movslq	(%r14,%rcx,4), %rdx
	movsd	(%r13,%rdx,8), %xmm1            # xmm1 = mem[0],zero
	subsd	%xmm0, %xmm1
	movsd	%xmm1, (%rbx,%rcx,8)
	incq	%rcx
	addl	%ebp, %r12d
	cmpq	%rax, %rcx
	jne	.LBB8_2
# %bb.4:                                # %.preheader51
	testl	%ebp, %ebp
	jle	.LBB8_10
# %bb.5:                                # %.preheader.lr.ph
	incl	%ebp
	movq	%r10, %rax
	imulq	%r10, %rax
	leaq	(%r15,%rax,8), %rax
	xorq	$-8, %r11
	leaq	(%r9,%r10,8), %rcx
	xorl	%edx, %edx
	movq	%r10, %rsi
.LBB8_6:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB8_8 Depth 2
	xorpd	%xmm0, %xmm0
	cmpq	%r10, %rsi
	jge	.LBB8_9
# %bb.7:                                # %.lr.ph59
                                        #   in Loop: Header=BB8_6 Depth=1
	xorl	%edi, %edi
.LBB8_8:                                #   Parent Loop BB8_6 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rax,%rdi,8), %xmm1            # xmm1 = mem[0],zero
	mulsd	(%rcx,%rdi,8), %xmm1
	addsd	%xmm1, %xmm0
	incq	%rdi
	cmpl	%edi, %edx
	jne	.LBB8_8
.LBB8_9:                                # %._crit_edge60
                                        #   in Loop: Header=BB8_6 Depth=1
	leaq	-1(%rsi), %rdi
	movsd	(%rbx,%rdi,8), %xmm1            # xmm1 = mem[0],zero
	subsd	%xmm0, %xmm1
	movl	%ebp, %r8d
	imull	%edi, %r8d
	movslq	%r8d, %r8
	divsd	(%r15,%r8,8), %xmm1
	movsd	%xmm1, (%r9,%rdi,8)
	addq	%r11, %rax
	addq	$-8, %rcx
	incq	%rdx
	cmpq	$1, %rsi
	movq	%rdi, %rsi
	jg	.LBB8_6
.LBB8_10:                               # %._crit_edge65
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end8:
	.size	_Z9LUP_solvePKdPKiS0_iPdS3_, .Lfunc_end8-_Z9LUP_solvePKdPKiS0_iPdS3_
	.cfi_endproc
                                        # -- End function
	.globl	_Z11LUP_inv_matPdS_iS_Pid       # -- Begin function _Z11LUP_inv_matPdS_iS_Pid
	.type	_Z11LUP_inv_matPdS_iS_Pid,@function
_Z11LUP_inv_matPdS_iS_Pid:              # @_Z11LUP_inv_matPdS_iS_Pid
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movsd	%xmm0, 8(%rsp)                  # 8-byte Spill
	movq	%r8, 16(%rsp)                   # 8-byte Spill
	movq	%rcx, %rbp
	movl	%edx, %r14d
	movq	%rsi, %r15
	movq	%rdi, 24(%rsp)                  # 8-byte Spill
	movl	%edx, %eax
	imull	%eax, %eax
	movq	%rax, 64(%rsp)                  # 8-byte Spill
	leaq	(%rcx,%rax,8), %r12
	movslq	%edx, %rbx
	leaq	(%r12,%rbx,8), %rax
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	leaq	(,%rbx,8), %rdx
	xorl	%r13d, %r13d
	movq	%r12, %rdi
	xorl	%esi, %esi
	movq	%rdx, 72(%rsp)                  # 8-byte Spill
	callq	memset@PLT
	movabsq	$4607182418800017408, %rax      # imm = 0x3FF0000000000000
	movq	%r12, 40(%rsp)                  # 8-byte Spill
	movq	%rax, (%r12)
	movq	24(%rsp), %r12                  # 8-byte Reload
	movq	%r12, %rdi
	movl	%ebx, %esi
	movq	16(%rsp), %rbx                  # 8-byte Reload
	movq	%rbx, %rdx
	movsd	8(%rsp), %xmm0                  # 8-byte Reload
                                        # xmm0 = mem[0],zero
	callq	_Z3LUPPdiPid
	cmpl	$2, %r14d
	jl	.LBB9_4
# %bb.1:                                # %.lr.ph.preheader
	movq	%r15, 56(%rsp)                  # 8-byte Spill
	leal	-1(%r14), %eax
	movl	%eax, 36(%rsp)                  # 4-byte Spill
	movl	%eax, %eax
	movq	%rax, 80(%rsp)                  # 8-byte Spill
	movq	64(%rsp), %rax                  # 8-byte Reload
	leaq	8(,%rax,8), %r15
	addq	%rbp, %r15
	movq	%r14, 8(%rsp)                   # 8-byte Spill
	xorl	%r12d, %r12d
	movq	%rbp, %rbx
	movq	40(%rsp), %r14                  # 8-byte Reload
	movq	72(%rsp), %r13                  # 8-byte Reload
.LBB9_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	24(%rsp), %rdi                  # 8-byte Reload
	movq	16(%rsp), %rsi                  # 8-byte Reload
	movq	%r14, %rdx
	movq	8(%rsp), %rcx                   # 8-byte Reload
                                        # kill: def $ecx killed $ecx killed $rcx
	movq	%rbx, %r8
	movq	48(%rsp), %r9                   # 8-byte Reload
	callq	_Z9LUP_solvePKdPKiS0_iPdS3_
	movq	$0, -8(%r15,%r12,8)
	movabsq	$4607182418800017408, %rax      # imm = 0x3FF0000000000000
	movq	%rax, (%r15,%r12,8)
	incq	%r12
	addq	%r13, %rbx
	cmpq	%r12, 80(%rsp)                  # 8-byte Folded Reload
	jne	.LBB9_2
# %bb.3:
	movq	56(%rsp), %r15                  # 8-byte Reload
	movq	16(%rsp), %rbx                  # 8-byte Reload
	movq	8(%rsp), %r14                   # 8-byte Reload
	movq	24(%rsp), %r12                  # 8-byte Reload
	movl	36(%rsp), %r13d                 # 4-byte Reload
.LBB9_4:                                # %._crit_edge
	imull	%r14d, %r13d
	movslq	%r13d, %rax
	leaq	(,%rax,8), %r8
	addq	%rbp, %r8
	movq	%r12, %rdi
	movq	%rbx, %rsi
	movq	40(%rsp), %rdx                  # 8-byte Reload
	movl	%r14d, %ecx
	movq	48(%rsp), %r9                   # 8-byte Reload
	callq	_Z9LUP_solvePKdPKiS0_iPdS3_
	testl	%r14d, %r14d
	jle	.LBB9_9
# %bb.5:                                # %.preheader.lr.ph.i
	movl	%r14d, %eax
	leaq	(,%rax,8), %rcx
	xorl	%edx, %edx
	xorl	%esi, %esi
.LBB9_6:                                # %.lr.ph.i
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB9_7 Depth 2
	movl	%edx, %edi
	leaq	(%r15,%rdi,8), %rdi
	movq	%rbp, %r8
	xorl	%r9d, %r9d
.LBB9_7:                                #   Parent Loop BB9_6 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%r8), %xmm0                    # xmm0 = mem[0],zero
	movsd	%xmm0, (%rdi,%r9)
	addq	$8, %r9
	addq	%rcx, %r8
	cmpq	%r9, %rcx
	jne	.LBB9_7
# %bb.8:                                # %._crit_edge.i
                                        #   in Loop: Header=BB9_6 Depth=1
	incq	%rsi
	addl	%r14d, %edx
	addq	$8, %rbp
	cmpq	%rax, %rsi
	jne	.LBB9_6
.LBB9_9:                                # %_Z13mat_transposePKdPdi.exit
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end9:
	.size	_Z11LUP_inv_matPdS_iS_Pid, .Lfunc_end9-_Z11LUP_inv_matPdS_iS_Pid
	.cfi_endproc
                                        # -- End function
	.globl	_Z21__device_stub__kerneliP15HIP_vector_typeIhLj4EEii # -- Begin function _Z21__device_stub__kerneliP15HIP_vector_typeIhLj4EEii
	.type	_Z21__device_stub__kerneliP15HIP_vector_typeIhLj4EEii,@function
_Z21__device_stub__kerneliP15HIP_vector_typeIhLj4EEii: # @_Z21__device_stub__kerneliP15HIP_vector_typeIhLj4EEii
	.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	20(%rsp), %rax
	movl	%edi, (%rax)
	leaq	40(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	16(%rsp), %rsi
	movl	%edx, (%rsi)
	leaq	12(%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	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z6kerneliP15HIP_vector_typeIhLj4EEii, %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_end10:
	.size	_Z21__device_stub__kerneliP15HIP_vector_typeIhLj4EEii, .Lfunc_end10-_Z21__device_stub__kerneliP15HIP_vector_typeIhLj4EEii
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI11_0:
	.quad	0x3ff0000000000000              # double 1
.LCPI11_1:
	.quad	0x3ddb7cdfd9d7bdbb              # double 1.0E-10
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0
.LCPI11_2:
	.zero	16
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
.Lfunc_begin0:
	.cfi_startproc
	.cfi_personality 3, __gxx_personality_v0
	.cfi_lsda 3, .Lexception0
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$808, %rsp                      # imm = 0x328
	.cfi_def_cfa_offset 864
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	xorl	%eax, %eax
	movq	%rax, 8(%rsp)
	movq	%rax, 32(%rsp)
	xorpd	%xmm0, %xmm0
	movapd	%xmm0, 16(%rsp)
.Ltmp0:
	leaq	432(%rsp), %rsi
	leaq	176(%rsp), %rdx
	movq	%rsp, %rcx
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	__isoc23_scanf
.Ltmp1:
# %bb.1:
	leaq	432(%rsp), %rdi
	movl	$.L.str.1, %esi
	callq	fopen
	testq	%rax, %rax
	je	.LBB11_2
# %bb.5:
	movq	%rax, %r14
	leaq	44(%rsp), %rbx
	movl	$4, %esi
	movl	$1, %edx
	movq	%rbx, %rdi
	movq	%rax, %rcx
	callq	fread
	leaq	40(%rsp), %r12
	movl	$4, %esi
	movl	$1, %edx
	movq	%r12, %rdi
	movq	%r14, %rcx
	callq	fread
	movslq	(%rbx), %r15
	movslq	(%r12), %r12
	movq	%r12, 64(%rsp)                  # 8-byte Spill
	imulq	%r15, %r12
	leaq	(,%r12,4), %r13
	movq	%r13, %rdi
	callq	malloc
	movq	%rax, %rbx
	testq	%rax, %rax
	je	.LBB11_6
# %bb.7:
.Ltmp2:
	leaq	8(%rsp), %rdi
	movq	%r13, 56(%rsp)                  # 8-byte Spill
	movq	%r13, %rsi
	callq	hipMalloc
.Ltmp3:
# %bb.8:                                # %_ZL9hipMallocI15HIP_vector_typeIhLj4EEE10hipError_tPPT_m.exit
	testl	%eax, %eax
	jne	.LBB11_9
# %bb.13:
	movl	$4, %esi
	movq	%rbx, %rdi
	movq	%r12, %rdx
	movq	%r14, %rcx
	callq	fread
	movq	%r14, %rdi
	callq	fclose
	movslq	(%rsp), %rax
	testq	%rax, %rax
	jle	.LBB11_49
# %bb.14:                               # %.lr.ph139.preheader
	movabsq	$4607182418800017408, %r13      # imm = 0x3FF0000000000000
	xorl	%r12d, %r12d
.LBB11_15:                              # %.lr.ph139
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB11_22 Depth 2
                                        #     Child Loop BB11_42 Depth 2
                                        #       Child Loop BB11_43 Depth 3
                                        #     Child Loop BB11_27 Depth 2
                                        #       Child Loop BB11_28 Depth 3
                                        #     Child Loop BB11_31 Depth 2
                                        #       Child Loop BB11_32 Depth 3
                                        #         Child Loop BB11_33 Depth 4
                                        #     Child Loop BB11_37 Depth 2
                                        #       Child Loop BB11_38 Depth 3
.Ltmp7:
	movl	$.L.str.6, %edi
	leaq	4(%rsp), %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
.Ltmp8:
# %bb.16:                               #   in Loop: Header=BB11_15 Depth=1
	movq	24(%rsp), %rcx
	subq	16(%rsp), %rcx
	sarq	$3, %rcx
	movabsq	$-6148914691236517205, %rax     # imm = 0xAAAAAAAAAAAAAAAB
	imulq	%rax, %rcx
	movl	4(%rsp), %eax
	movslq	%eax, %rsi
	cmpq	%rsi, %rcx
	jae	.LBB11_19
# %bb.17:                               #   in Loop: Header=BB11_15 Depth=1
.Ltmp9:
	leaq	16(%rsp), %rdi
	callq	_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE6resizeEm
.Ltmp10:
# %bb.18:                               # %._crit_edge152
                                        #   in Loop: Header=BB11_15 Depth=1
	movl	4(%rsp), %eax
.LBB11_19:                              #   in Loop: Header=BB11_15 Depth=1
	leaq	(,%r12,8), %rcx
	leaq	(%rcx,%rcx,2), %rbp
	addq	$avg, %rbp
	xorpd	%xmm0, %xmm0
	movupd	%xmm0, (%rbp)
	movq	$0, 16(%rbp)
	testl	%eax, %eax
	jle	.LBB11_20
# %bb.21:                               # %.lr.ph.preheader
                                        #   in Loop: Header=BB11_15 Depth=1
	movq	%r12, 72(%rsp)                  # 8-byte Spill
	movl	$16, %r12d
	xorl	%r14d, %r14d
.LBB11_22:                              # %.lr.ph
                                        #   Parent Loop BB11_15 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
.Ltmp12:
	movl	$.L.str.7, %edi
	leaq	52(%rsp), %rsi
	leaq	48(%rsp), %rdx
	xorl	%eax, %eax
	callq	__isoc23_scanf
.Ltmp13:
# %bb.23:                               #   in Loop: Header=BB11_22 Depth=2
	movslq	48(%rsp), %rax
	imulq	%r15, %rax
	movslq	52(%rsp), %rcx
	addq	%rax, %rcx
	movzbl	(%rbx,%rcx,4), %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	movzbl	1(%rbx,%rcx,4), %eax
	xorps	%xmm1, %xmm1
	cvtsi2sd	%eax, %xmm1
	movzbl	2(%rbx,%rcx,4), %eax
	xorps	%xmm2, %xmm2
	cvtsi2sd	%eax, %xmm2
	movq	16(%rsp), %rax
	movsd	%xmm0, -16(%rax,%r12)
	movsd	%xmm1, -8(%rax,%r12)
	movsd	%xmm2, (%rax,%r12)
	movq	16(%rsp), %rax
	movsd	-16(%rax,%r12), %xmm2           # xmm2 = mem[0],zero
	addsd	(%rbp), %xmm2
	movsd	%xmm2, (%rbp)
	movsd	-8(%rax,%r12), %xmm1            # xmm1 = mem[0],zero
	addsd	8(%rbp), %xmm1
	movsd	%xmm1, 8(%rbp)
	movsd	(%rax,%r12), %xmm0              # xmm0 = mem[0],zero
	addsd	16(%rbp), %xmm0
	movsd	%xmm0, 16(%rbp)
	incq	%r14
	movslq	4(%rsp), %rax
	addq	$24, %r12
	cmpq	%rax, %r14
	jl	.LBB11_22
# %bb.24:                               #   in Loop: Header=BB11_15 Depth=1
	movq	72(%rsp), %r12                  # 8-byte Reload
	jmp	.LBB11_25
.LBB11_20:                              #   in Loop: Header=BB11_15 Depth=1
	xorpd	%xmm0, %xmm0
	xorpd	%xmm1, %xmm1
	xorpd	%xmm2, %xmm2
.LBB11_25:                              # %._crit_edge
                                        #   in Loop: Header=BB11_15 Depth=1
	xorps	%xmm3, %xmm3
	cvtsi2sd	%eax, %xmm3
	divsd	%xmm3, %xmm2
	movsd	%xmm2, (%rbp)
	divsd	%xmm3, %xmm1
	movsd	%xmm1, 8(%rbp)
	divsd	%xmm3, %xmm0
	movsd	%xmm0, 16(%rbp)
	cmpl	$2, %eax
	jl	.LBB11_41
# %bb.26:                               # %.lr.ph.i.preheader
                                        #   in Loop: Header=BB11_15 Depth=1
	leaq	96(%rsp), %rcx
	xorl	%edx, %edx
.LBB11_27:                              # %.lr.ph.i
                                        #   Parent Loop BB11_15 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB11_28 Depth 3
	leaq	(%rsp,%rdx,8), %rsi
	addq	$96, %rsi
	leaq	(%rdx,%rdx,2), %rdi
	leaq	(%rsi,%rdi,8), %rsi
	movq	%rcx, %rdi
	xorl	%r8d, %r8d
.LBB11_28:                              #   Parent Loop BB11_15 Depth=1
                                        #     Parent Loop BB11_27 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	cmpq	%r8, %rdx
	movq	%rdi, %r9
	cmoveq	%rsi, %r9
	movq	$0, (%r9)
	incq	%r8
	addq	$8, %rdi
	cmpq	$3, %r8
	jne	.LBB11_28
# %bb.29:                               # %._crit_edge.i
                                        #   in Loop: Header=BB11_27 Depth=2
	incq	%rdx
	addq	$24, %rcx
	cmpq	$3, %rdx
	jne	.LBB11_27
# %bb.30:                               # %.lr.ph134.preheader
                                        #   in Loop: Header=BB11_15 Depth=1
	movl	%eax, %ecx
	movq	16(%rsp), %rdi
	xorl	%edx, %edx
	xorl	%esi, %esi
.LBB11_31:                              # %.lr.ph134
                                        #   Parent Loop BB11_15 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB11_32 Depth 3
                                        #         Child Loop BB11_33 Depth 4
	leaq	(,%rsi,8), %r8
	leaq	(%r8,%r8,2), %r8
	movsd	(%rdi,%r8), %xmm0               # xmm0 = mem[0],zero
	subsd	(%rbp), %xmm0
	movsd	%xmm0, (%rdi,%r8)
	movq	16(%rsp), %rdi
	movsd	8(%rdi,%r8), %xmm0              # xmm0 = mem[0],zero
	subsd	8(%rbp), %xmm0
	movsd	%xmm0, 8(%rdi,%r8)
	movq	16(%rsp), %rdi
	movsd	16(%rdi,%r8), %xmm0             # xmm0 = mem[0],zero
	subsd	16(%rbp), %xmm0
	movsd	%xmm0, 16(%rdi,%r8)
	movq	16(%rsp), %rdi
	addq	%rdi, %r8
	leaq	(%rdi,%rdx), %r9
	leaq	96(%rsp), %r10
	xorl	%r11d, %r11d
.LBB11_32:                              # %.preheader35.i
                                        #   Parent Loop BB11_15 Depth=1
                                        #     Parent Loop BB11_31 Depth=2
                                        # =>    This Loop Header: Depth=3
                                        #         Child Loop BB11_33 Depth 4
	movsd	(%r8,%r11,8), %xmm0             # xmm0 = mem[0],zero
	xorl	%r14d, %r14d
	xorpd	%xmm2, %xmm2
.LBB11_33:                              # %.preheader.i
                                        #   Parent Loop BB11_15 Depth=1
                                        #     Parent Loop BB11_31 Depth=2
                                        #       Parent Loop BB11_32 Depth=3
                                        # =>      This Inner Loop Header: Depth=4
	movsd	(%r9,%r14,8), %xmm1             # xmm1 = mem[0],zero
	mulsd	%xmm0, %xmm1
	addsd	%xmm2, %xmm1
	addsd	(%r10,%r14,8), %xmm1
	movsd	%xmm1, (%r10,%r14,8)
	incq	%r14
	cmpq	$3, %r14
	jne	.LBB11_33
# %bb.34:                               # %._crit_edge39.i
                                        #   in Loop: Header=BB11_32 Depth=3
	incq	%r11
	addq	$24, %r10
	cmpq	$3, %r11
	jne	.LBB11_32
# %bb.35:                               # %_Z7mat_mulPKdS0_Pdiiidd.exit
                                        #   in Loop: Header=BB11_31 Depth=2
	incq	%rsi
	addq	$24, %rdx
	cmpq	%rcx, %rsi
	jne	.LBB11_31
# %bb.36:                               # %_Z7mat_setPdidd.exit._crit_edge
                                        #   in Loop: Header=BB11_15 Depth=1
	decl	%eax
	xorps	%xmm1, %xmm1
	cvtsi2sd	%eax, %xmm1
	movsd	.LCPI11_0(%rip), %xmm0          # xmm0 = mem[0],zero
	divsd	%xmm1, %xmm0
	leaq	96(%rsp), %rax
	xorl	%ecx, %ecx
.LBB11_37:                              # %.preheader.i98
                                        #   Parent Loop BB11_15 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB11_38 Depth 3
	xorl	%edx, %edx
.LBB11_38:                              #   Parent Loop BB11_15 Depth=1
                                        #     Parent Loop BB11_37 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movsd	(%rax,%rdx,8), %xmm1            # xmm1 = mem[0],zero
	mulsd	%xmm0, %xmm1
	movsd	%xmm1, (%rax,%rdx,8)
	incq	%rdx
	cmpq	$3, %rdx
	jne	.LBB11_38
# %bb.39:                               # %._crit_edge.i103
                                        #   in Loop: Header=BB11_37 Depth=2
	incq	%rcx
	addq	$24, %rax
	cmpq	$3, %rcx
	jne	.LBB11_37
# %bb.40:                               # %_Z9mat_mul_CPKdPddii.exit
                                        #   in Loop: Header=BB11_15 Depth=1
	movq	%r12, %rax
	movabsq	$38654705664, %rcx              # imm = 0x900000000
	imulq	%rcx, %rax
	sarq	$29, %rax
	leaq	cov(%rax), %rsi
	leaq	96(%rsp), %rdi
	movl	$3, %edx
	leaq	688(%rsp), %rcx
	leaq	84(%rsp), %r8
	movsd	.LCPI11_1(%rip), %xmm0          # xmm0 = mem[0],zero
	callq	_Z11LUP_inv_matPdS_iS_Pid
	jmp	.LBB11_48
.LBB11_41:                              #   in Loop: Header=BB11_15 Depth=1
	movq	%r12, %rax
	movabsq	$38654705664, %rcx              # imm = 0x900000000
	imulq	%rcx, %rax
	sarq	$29, %rax
	leaq	cov(%rax), %rax
	movq	%rax, %rcx
	xorl	%edx, %edx
.LBB11_42:                              # %.lr.ph.i104
                                        #   Parent Loop BB11_15 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB11_43 Depth 3
	leaq	(%rdx,%rdx,2), %rsi
	leaq	(%rax,%rdx,8), %rdi
	xorl	%r8d, %r8d
.LBB11_43:                              #   Parent Loop BB11_15 Depth=1
                                        #     Parent Loop BB11_42 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	cmpq	%r8, %rdx
	jne	.LBB11_45
# %bb.44:                               #   in Loop: Header=BB11_43 Depth=3
	movq	%r13, (%rdi,%rsi,8)
	jmp	.LBB11_46
.LBB11_45:                              #   in Loop: Header=BB11_43 Depth=3
	movq	$0, (%rcx,%r8,8)
.LBB11_46:                              #   in Loop: Header=BB11_43 Depth=3
	incq	%r8
	cmpq	$3, %r8
	jne	.LBB11_43
# %bb.47:                               # %._crit_edge.i111
                                        #   in Loop: Header=BB11_42 Depth=2
	incq	%rdx
	addq	$24, %rcx
	cmpq	$3, %rdx
	jne	.LBB11_42
.LBB11_48:                              # %_Z7mat_setPdidd.exit114
                                        #   in Loop: Header=BB11_15 Depth=1
	incq	%r12
	movslq	(%rsp), %rax
	cmpq	%rax, %r12
	jl	.LBB11_15
.LBB11_49:                              # %._crit_edge140
	shlq	$3, %rax
	leaq	(%rax,%rax,2), %rdx
.Ltmp15:
	movl	$AVG, %edi
	movl	$avg, %esi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
.Ltmp16:
# %bb.50:                               # %_Z17hipMemcpyToSymbolIA32_15HIP_vector_typeIdLj3EEE10hipError_tRKT_PKvmm13hipMemcpyKind.exit
	testl	%eax, %eax
	jne	.LBB11_51
# %bb.54:
	movl	(%rsp), %eax
	leal	(%rax,%rax,8), %eax
	movslq	%eax, %rdx
	shlq	$3, %rdx
.Ltmp20:
	movl	$COV, %edi
	movl	$cov, %esi
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
.Ltmp21:
# %bb.55:                               # %_Z17hipMemcpyToSymbolIA288_dE10hipError_tRKT_PKvmm13hipMemcpyKind.exit
	testl	%eax, %eax
	jne	.LBB11_56
# %bb.59:
	movq	8(%rsp), %rdi
.Ltmp25:
	movq	%rbx, %rsi
	movq	56(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
.Ltmp26:
# %bb.60:
.Ltmp27:
	movabsq	$4294967552, %rdi               # imm = 0x100000100
	movl	$1, %esi
	movq	%rdi, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
.Ltmp28:
# %bb.61:
	testl	%eax, %eax
	jne	.LBB11_63
# %bb.62:
	movl	(%rsp), %edi
	movq	8(%rsp), %rsi
.Ltmp29:
	movl	%r15d, %edx
	movq	64(%rsp), %rcx                  # 8-byte Reload
                                        # kill: def $ecx killed $ecx killed $rcx
	callq	_Z21__device_stub__kerneliP15HIP_vector_typeIhLj4EEii
.Ltmp30:
.LBB11_63:
.Ltmp31:
	callq	hipGetLastError
.Ltmp32:
# %bb.64:
	testl	%eax, %eax
	movq	56(%rsp), %rdx                  # 8-byte Reload
	jne	.LBB11_65
# %bb.68:
	movq	8(%rsp), %rsi
.Ltmp36:
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
.Ltmp37:
# %bb.69:
	testl	%eax, %eax
	jne	.LBB11_70
# %bb.73:
	leaq	176(%rsp), %rdi
	movl	$.L.str.8, %esi
	callq	fopen
	testq	%rax, %rax
	je	.LBB11_74
# %bb.75:
	movq	%rax, %r14
	leaq	44(%rsp), %r15
	movl	$4, %esi
	movl	$1, %edx
	movq	%r15, %rdi
	movq	%rax, %rcx
	callq	fwrite
	leaq	40(%rsp), %r12
	movl	$4, %esi
	movl	$1, %edx
	movq	%r12, %rdi
	movq	%r14, %rcx
	callq	fwrite
	movslq	(%r15), %rax
	movslq	(%r12), %rdx
	imulq	%rax, %rdx
	movl	$4, %esi
	movq	%rbx, %rdi
	movq	%r14, %rcx
	callq	fwrite
	movq	%r14, %rdi
	callq	fclose
.LBB11_76:
	movq	%rbx, %rdi
	callq	free
	movq	8(%rsp), %rdi
.Ltmp41:
	callq	hipFree
.Ltmp42:
.LBB11_77:
	movq	16(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB11_79
# %bb.78:
	callq	_ZdlPv
.LBB11_79:                              # %_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EED2Ev.exit
	xorl	%eax, %eax
	addq	$808, %rsp                      # imm = 0x328
	.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
.LBB11_2:
	.cfi_def_cfa_offset 864
	movq	stderr(%rip), %rdi
	leaq	432(%rsp), %rdx
	movl	$.L.str.2, %esi
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB11_77
.LBB11_6:
	movq	stderr(%rip), %rcx
	movl	$.L.str.3, %edi
	movl	$32, %esi
	movl	$1, %edx
	callq	fwrite
	jmp	.LBB11_76
.LBB11_74:
	movq	stderr(%rip), %rdi
	leaq	176(%rsp), %rdx
	movl	$.L.str.2, %esi
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB11_76
.LBB11_9:
	movq	stderr(%rip), %rbx
.Ltmp4:
	movl	%eax, %edi
	callq	hipGetErrorString
.Ltmp5:
# %bb.10:
	movl	$.L.str.4, %esi
	movl	$.L.str.5, %edx
	movq	%rbx, %rdi
	movl	$317, %ecx                      # imm = 0x13D
	jmp	.LBB11_11
.LBB11_51:
	movq	stderr(%rip), %rbx
.Ltmp17:
	movl	%eax, %edi
	callq	hipGetErrorString
.Ltmp18:
# %bb.52:
	movl	$.L.str.4, %esi
	movl	$.L.str.5, %edx
	movq	%rbx, %rdi
	movl	$364, %ecx                      # imm = 0x16C
	jmp	.LBB11_11
.LBB11_56:
	movq	stderr(%rip), %rbx
.Ltmp22:
	movl	%eax, %edi
	callq	hipGetErrorString
.Ltmp23:
# %bb.57:
	movl	$.L.str.4, %esi
	movl	$.L.str.5, %edx
	movq	%rbx, %rdi
	movl	$365, %ecx                      # imm = 0x16D
	jmp	.LBB11_11
.LBB11_65:
	movq	stderr(%rip), %rbx
.Ltmp33:
	movl	%eax, %edi
	callq	hipGetErrorString
.Ltmp34:
# %bb.66:
	movl	$.L.str.4, %esi
	movl	$.L.str.5, %edx
	movq	%rbx, %rdi
	movl	$371, %ecx                      # imm = 0x173
	jmp	.LBB11_11
.LBB11_70:
	movq	stderr(%rip), %rbx
.Ltmp38:
	movl	%eax, %edi
	callq	hipGetErrorString
.Ltmp39:
# %bb.71:
	movl	$.L.str.4, %esi
	movl	$.L.str.5, %edx
	movq	%rbx, %rdi
	movl	$372, %ecx                      # imm = 0x174
.LBB11_11:
	movq	%rax, %r8
	xorl	%eax, %eax
	callq	fprintf
	xorl	%edi, %edi
	callq	exit
.LBB11_72:
.Ltmp40:
	jmp	.LBB11_81
.LBB11_67:
.Ltmp35:
	jmp	.LBB11_81
.LBB11_58:
.Ltmp24:
	jmp	.LBB11_81
.LBB11_53:
.Ltmp19:
	jmp	.LBB11_81
.LBB11_12:
.Ltmp6:
	jmp	.LBB11_81
.LBB11_4:                               # %.loopexit.split-lp.loopexit.split-lp
.Ltmp43:
	jmp	.LBB11_81
.LBB11_3:                               # %.loopexit.split-lp.loopexit
.Ltmp11:
	jmp	.LBB11_81
.LBB11_80:                              # %.loopexit
.Ltmp14:
.LBB11_81:                              # %.loopexit.split-lp
	movq	%rax, %rbx
	movq	16(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB11_83
# %bb.82:
	callq	_ZdlPv
.LBB11_83:                              # %_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EED2Ev.exit116
	movq	%rbx, %rdi
	callq	_Unwind_Resume@PLT
.Lfunc_end11:
	.size	main, .Lfunc_end11-main
	.cfi_endproc
	.section	.gcc_except_table,"a",@progbits
	.p2align	2, 0x0
GCC_except_table11:
.Lexception0:
	.byte	255                             # @LPStart Encoding = omit
	.byte	255                             # @TType Encoding = omit
	.byte	1                               # Call site Encoding = uleb128
	.uleb128 .Lcst_end0-.Lcst_begin0
.Lcst_begin0:
	.uleb128 .Ltmp0-.Lfunc_begin0           # >> Call Site 1 <<
	.uleb128 .Ltmp1-.Ltmp0                  #   Call between .Ltmp0 and .Ltmp1
	.uleb128 .Ltmp43-.Lfunc_begin0          #     jumps to .Ltmp43
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp2-.Lfunc_begin0           # >> Call Site 2 <<
	.uleb128 .Ltmp3-.Ltmp2                  #   Call between .Ltmp2 and .Ltmp3
	.uleb128 .Ltmp6-.Lfunc_begin0           #     jumps to .Ltmp6
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp7-.Lfunc_begin0           # >> Call Site 3 <<
	.uleb128 .Ltmp10-.Ltmp7                 #   Call between .Ltmp7 and .Ltmp10
	.uleb128 .Ltmp11-.Lfunc_begin0          #     jumps to .Ltmp11
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp12-.Lfunc_begin0          # >> Call Site 4 <<
	.uleb128 .Ltmp13-.Ltmp12                #   Call between .Ltmp12 and .Ltmp13
	.uleb128 .Ltmp14-.Lfunc_begin0          #     jumps to .Ltmp14
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp15-.Lfunc_begin0          # >> Call Site 5 <<
	.uleb128 .Ltmp16-.Ltmp15                #   Call between .Ltmp15 and .Ltmp16
	.uleb128 .Ltmp19-.Lfunc_begin0          #     jumps to .Ltmp19
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp20-.Lfunc_begin0          # >> Call Site 6 <<
	.uleb128 .Ltmp21-.Ltmp20                #   Call between .Ltmp20 and .Ltmp21
	.uleb128 .Ltmp24-.Lfunc_begin0          #     jumps to .Ltmp24
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp25-.Lfunc_begin0          # >> Call Site 7 <<
	.uleb128 .Ltmp30-.Ltmp25                #   Call between .Ltmp25 and .Ltmp30
	.uleb128 .Ltmp43-.Lfunc_begin0          #     jumps to .Ltmp43
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp31-.Lfunc_begin0          # >> Call Site 8 <<
	.uleb128 .Ltmp32-.Ltmp31                #   Call between .Ltmp31 and .Ltmp32
	.uleb128 .Ltmp35-.Lfunc_begin0          #     jumps to .Ltmp35
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp36-.Lfunc_begin0          # >> Call Site 9 <<
	.uleb128 .Ltmp37-.Ltmp36                #   Call between .Ltmp36 and .Ltmp37
	.uleb128 .Ltmp40-.Lfunc_begin0          #     jumps to .Ltmp40
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp41-.Lfunc_begin0          # >> Call Site 10 <<
	.uleb128 .Ltmp42-.Ltmp41                #   Call between .Ltmp41 and .Ltmp42
	.uleb128 .Ltmp43-.Lfunc_begin0          #     jumps to .Ltmp43
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp4-.Lfunc_begin0           # >> Call Site 11 <<
	.uleb128 .Ltmp5-.Ltmp4                  #   Call between .Ltmp4 and .Ltmp5
	.uleb128 .Ltmp6-.Lfunc_begin0           #     jumps to .Ltmp6
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp17-.Lfunc_begin0          # >> Call Site 12 <<
	.uleb128 .Ltmp18-.Ltmp17                #   Call between .Ltmp17 and .Ltmp18
	.uleb128 .Ltmp19-.Lfunc_begin0          #     jumps to .Ltmp19
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp22-.Lfunc_begin0          # >> Call Site 13 <<
	.uleb128 .Ltmp23-.Ltmp22                #   Call between .Ltmp22 and .Ltmp23
	.uleb128 .Ltmp24-.Lfunc_begin0          #     jumps to .Ltmp24
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp33-.Lfunc_begin0          # >> Call Site 14 <<
	.uleb128 .Ltmp34-.Ltmp33                #   Call between .Ltmp33 and .Ltmp34
	.uleb128 .Ltmp35-.Lfunc_begin0          #     jumps to .Ltmp35
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp38-.Lfunc_begin0          # >> Call Site 15 <<
	.uleb128 .Ltmp39-.Ltmp38                #   Call between .Ltmp38 and .Ltmp39
	.uleb128 .Ltmp40-.Lfunc_begin0          #     jumps to .Ltmp40
	.byte	0                               #   On action: cleanup
	.uleb128 .Ltmp39-.Lfunc_begin0          # >> Call Site 16 <<
	.uleb128 .Lfunc_end11-.Ltmp39           #   Call between .Ltmp39 and .Lfunc_end11
	.byte	0                               #     has no landing pad
	.byte	0                               #   On action: cleanup
.Lcst_end0:
	.p2align	2, 0x0
                                        # -- End function
	.section	.text._ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE6resizeEm,"axG",@progbits,_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE6resizeEm,comdat
	.weak	_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE6resizeEm # -- Begin function _ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE6resizeEm
	.p2align	1, 0x90
	.type	_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE6resizeEm,@function
_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE6resizeEm: # @_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE6resizeEm
	.cfi_startproc
# %bb.0:
	movq	(%rdi), %rdx
	movq	8(%rdi), %rcx
	movq	%rcx, %rax
	subq	%rdx, %rax
	sarq	$3, %rax
	movabsq	$-6148914691236517205, %r8      # imm = 0xAAAAAAAAAAAAAAAB
	imulq	%rax, %r8
	movq	%rsi, %rax
	subq	%r8, %rax
	jbe	.LBB12_1
# %bb.5:
	movq	%rax, %rsi
	jmp	_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE17_M_default_appendEm # TAILCALL
.LBB12_1:
	jae	.LBB12_4
# %bb.2:
	leaq	(%rsi,%rsi,2), %rax
	leaq	(%rdx,%rax,8), %rax
	cmpq	%rax, %rcx
	je	.LBB12_4
# %bb.3:
	movq	%rax, 8(%rdi)
.LBB12_4:                               # %_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE15_M_erase_at_endEPS1_.exit
	retq
.Lfunc_end12:
	.size	_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE6resizeEm, .Lfunc_end12-_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE6resizeEm
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt15__new_allocatorI15HIP_vector_typeIdLj3EEE8allocateEmPKv,"axG",@progbits,_ZNSt15__new_allocatorI15HIP_vector_typeIdLj3EEE8allocateEmPKv,comdat
	.weak	_ZNSt15__new_allocatorI15HIP_vector_typeIdLj3EEE8allocateEmPKv # -- Begin function _ZNSt15__new_allocatorI15HIP_vector_typeIdLj3EEE8allocateEmPKv
	.p2align	1, 0x90
	.type	_ZNSt15__new_allocatorI15HIP_vector_typeIdLj3EEE8allocateEmPKv,@function
_ZNSt15__new_allocatorI15HIP_vector_typeIdLj3EEE8allocateEmPKv: # @_ZNSt15__new_allocatorI15HIP_vector_typeIdLj3EEE8allocateEmPKv
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movabsq	$384307168202282326, %rax       # imm = 0x555555555555556
	cmpq	%rax, %rsi
	jae	.LBB13_1
# %bb.3:
	shlq	$3, %rsi
	leaq	(%rsi,%rsi,2), %rdi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	_Znwm                           # TAILCALL
.LBB13_1:
	.cfi_def_cfa_offset 16
	movabsq	$768614336404564651, %rax       # imm = 0xAAAAAAAAAAAAAAB
	cmpq	%rax, %rsi
	jb	.LBB13_2
# %bb.4:
	callq	_ZSt28__throw_bad_array_new_lengthv
.LBB13_2:
	callq	_ZSt17__throw_bad_allocv
.Lfunc_end13:
	.size	_ZNSt15__new_allocatorI15HIP_vector_typeIdLj3EEE8allocateEmPKv, .Lfunc_end13-_ZNSt15__new_allocatorI15HIP_vector_typeIdLj3EEE8allocateEmPKv
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE17_M_default_appendEm,"axG",@progbits,_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE17_M_default_appendEm,comdat
	.weak	_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE17_M_default_appendEm # -- Begin function _ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE17_M_default_appendEm
	.p2align	1, 0x90
	.type	_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE17_M_default_appendEm,@function
_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE17_M_default_appendEm: # @_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE17_M_default_appendEm
	.cfi_startproc
# %bb.0:
	testq	%rsi, %rsi
	je	.LBB14_13
# %bb.1:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %r14
	movq	%rdi, %rbx
	movq	8(%rdi), %r15
	movabsq	$-6148914691236517205, %rax     # imm = 0xAAAAAAAAAAAAAAAB
	movq	16(%rdi), %rcx
	subq	%r15, %rcx
	sarq	$3, %rcx
	imulq	%rax, %rcx
	cmpq	%rsi, %rcx
	jae	.LBB14_2
# %bb.3:
	movq	(%rbx), %r12
	movq	%r15, %r13
	subq	%r12, %r13
	sarq	$3, %r13
	imulq	%rax, %r13
	movl	$.L.str.10, %edx
	movq	%rbx, %rdi
	movq	%r14, %rsi
	callq	_ZNKSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE12_M_check_lenEmPKc
	testq	%rax, %rax
	movq	%rax, (%rsp)                    # 8-byte Spill
	je	.LBB14_4
# %bb.5:
	movq	%rbx, %rdi
	movq	%rax, %rsi
	xorl	%edx, %edx
	callq	_ZNSt15__new_allocatorI15HIP_vector_typeIdLj3EEE8allocateEmPKv
	movq	%rax, %rbp
	jmp	.LBB14_6
.LBB14_2:                               # %_ZSt27__uninitialized_default_n_aIP15HIP_vector_typeIdLj3EEmS1_ET_S3_T0_RSaIT1_E.exit
	shlq	$3, %r14
	leaq	(%r14,%r14,2), %r14
	movq	%r15, %rdi
	xorl	%esi, %esi
	movq	%r14, %rdx
	callq	memset@PLT
	addq	%r14, %r15
	movq	%r15, 8(%rbx)
	jmp	.LBB14_12
.LBB14_4:
	xorl	%ebp, %ebp
.LBB14_6:
	leaq	(,%r13,2), %rax
	addq	%r13, %rax
	leaq	(%rbp,%rax,8), %r13
	leaq	(,%r14,8), %rax
	leaq	(%rax,%rax,2), %rdx
	movq	%r13, %rdi
	xorl	%esi, %esi
	callq	memset@PLT
	cmpq	%r15, %r12
	je	.LBB14_9
# %bb.7:                                # %.lr.ph.i.i.i.preheader
	movq	%rbp, %rax
	movq	%r12, %rcx
.LBB14_8:                               # %.lr.ph.i.i.i
                                        # =>This Inner Loop Header: Depth=1
	movq	16(%rcx), %rdx
	movq	%rdx, 16(%rax)
	movups	(%rcx), %xmm0
	movups	%xmm0, (%rax)
	addq	$24, %rcx
	addq	$24, %rax
	cmpq	%r15, %rcx
	jne	.LBB14_8
.LBB14_9:                               # %_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE11_S_relocateEPS1_S4_S4_RS2_.exit
	testq	%r12, %r12
	je	.LBB14_11
# %bb.10:
	movq	%r12, %rdi
	callq	_ZdlPv
.LBB14_11:                              # %_ZNSt12_Vector_baseI15HIP_vector_typeIdLj3EESaIS1_EE13_M_deallocateEPS1_m.exit37
	movq	%rbp, (%rbx)
	leaq	(%r14,%r14,2), %rax
	leaq	(,%rax,8), %rax
	addq	%r13, %rax
	movq	%rax, 8(%rbx)
	movq	(%rsp), %rax                    # 8-byte Reload
	leaq	(%rax,%rax,2), %rax
	leaq	(,%rax,8), %rax
	addq	%rbp, %rax
	movq	%rax, 16(%rbx)
.LBB14_12:
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r12
	.cfi_restore %r13
	.cfi_restore %r14
	.cfi_restore %r15
	.cfi_restore %rbp
.LBB14_13:
	retq
.Lfunc_end14:
	.size	_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE17_M_default_appendEm, .Lfunc_end14-_ZNSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE17_M_default_appendEm
	.cfi_endproc
                                        # -- End function
	.section	.text._ZNKSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE12_M_check_lenEmPKc,"axG",@progbits,_ZNKSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE12_M_check_lenEmPKc,comdat
	.weak	_ZNKSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE12_M_check_lenEmPKc # -- Begin function _ZNKSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE12_M_check_lenEmPKc
	.p2align	1, 0x90
	.type	_ZNKSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE12_M_check_lenEmPKc,@function
_ZNKSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE12_M_check_lenEmPKc: # @_ZNKSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE12_M_check_lenEmPKc
	.cfi_startproc
# %bb.0:
	movabsq	$384307168202282325, %rcx       # imm = 0x555555555555555
	movq	8(%rdi), %rax
	subq	(%rdi), %rax
	sarq	$3, %rax
	movabsq	$-6148914691236517205, %rdi     # imm = 0xAAAAAAAAAAAAAAAB
	imulq	%rax, %rdi
	movq	%rcx, %rax
	subq	%rdi, %rax
	cmpq	%rsi, %rax
	jb	.LBB15_2
# %bb.1:
	cmpq	%rsi, %rdi
	cmovaq	%rdi, %rsi
	leaq	(%rsi,%rdi), %rax
	cmpq	%rcx, %rax
	cmovaeq	%rcx, %rax
	addq	%rdi, %rsi
	cmovbq	%rcx, %rax
	retq
.LBB15_2:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	%rdx, %rdi
	callq	_ZSt20__throw_length_errorPKc
.Lfunc_end15:
	.size	_ZNKSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE12_M_check_lenEmPKc, .Lfunc_end15-_ZNKSt6vectorI15HIP_vector_typeIdLj3EESaIS1_EE12_M_check_lenEmPKc
	.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
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB16_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB16_2:
	subq	$32, %rsp
	.cfi_adjust_cfa_offset 32
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z6kerneliP15HIP_vector_typeIhLj4EEii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	movl	$AVG, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movl	$768, %r9d                      # imm = 0x300
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$1
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$COV, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movl	$2304, %r9d                     # imm = 0x900
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$1
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$__hip_module_dtor, %edi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end16:
	.size	__hip_module_ctor, .Lfunc_end16-__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	.LBB17_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
.LBB17_2:
	retq
.Lfunc_end17:
	.size	__hip_module_dtor, .Lfunc_end17-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	avg,@object                     # @avg
	.bss
	.globl	avg
	.p2align	4, 0x0
avg:
	.zero	768
	.size	avg, 768

	.type	cov,@object                     # @cov
	.globl	cov
	.p2align	4, 0x0
cov:
	.zero	2304
	.size	cov, 2304

	.type	AVG,@object                     # @AVG
	.local	AVG
	.comm	AVG,768,16
	.type	COV,@object                     # @COV
	.local	COV
	.comm	COV,2304,16
	.type	_Z6kerneliP15HIP_vector_typeIhLj4EEii,@object # @_Z6kerneliP15HIP_vector_typeIhLj4EEii
	.section	.rodata,"a",@progbits
	.globl	_Z6kerneliP15HIP_vector_typeIhLj4EEii
	.p2align	3, 0x0
_Z6kerneliP15HIP_vector_typeIhLj4EEii:
	.quad	_Z21__device_stub__kerneliP15HIP_vector_typeIhLj4EEii
	.size	_Z6kerneliP15HIP_vector_typeIhLj4EEii, 8

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

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

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Error: can't open %s\n"
	.size	.L.str.2, 22

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Error: not enough memory in CPU\n"
	.size	.L.str.3, 33

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"ERROR in %s:%d. Message: %s\n"
	.size	.L.str.4, 29

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/iDneprov/university/master/PGP-POD/3/lab3.hip"
	.size	.L.str.5, 108

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"%d"
	.size	.L.str.6, 3

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"%d %d"
	.size	.L.str.7, 6

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"wb"
	.size	.L.str.8, 3

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"vector::_M_default_append"
	.size	.L.str.10, 26

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z6kerneliP15HIP_vector_typeIhLj4EEii"
	.size	.L__unnamed_1, 38

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"AVG"
	.size	.L__unnamed_2, 4

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"COV"
	.size	.L__unnamed_3, 4

	.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 _Z21__device_stub__kerneliP15HIP_vector_typeIhLj4EEii
	.addrsig_sym __gxx_personality_v0
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Unwind_Resume
	.addrsig_sym avg
	.addrsig_sym cov
	.addrsig_sym AVG
	.addrsig_sym COV
	.addrsig_sym _Z6kerneliP15HIP_vector_typeIhLj4EEii
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z6kerneliP15HIP_vector_typeIhLj4EEii ; -- Begin function _Z6kerneliP15HIP_vector_typeIhLj4EEii
	.globl	_Z6kerneliP15HIP_vector_typeIhLj4EEii
	.p2align	8
	.type	_Z6kerneliP15HIP_vector_typeIhLj4EEii,@function
_Z6kerneliP15HIP_vector_typeIhLj4EEii:  ; @_Z6kerneliP15HIP_vector_typeIhLj4EEii
; %bb.0:
	s_clause 0x1
	s_load_b32 s8, s[2:3], 0x24
	s_load_b64 s[4:5], s[2:3], 0x10
	v_and_b32_e32 v3, 0x3ff, v0
	s_add_u32 s6, s2, 24
	s_addc_u32 s7, s3, 0
	s_waitcnt lgkmcnt(0)
	s_and_b32 s11, s8, 0xffff
	s_mul_i32 s8, s5, s4
	v_mad_u64_u32 v[1:2], null, s15, s11, v[3:4]
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_i32_e64 s8, v1
	s_cbranch_execz .LBB0_11
; %bb.1:                                ; %.preheader.lr.ph
	s_load_b64 s[0:1], s[0:1], 0x4
	s_clause 0x1
	s_load_b64 s[4:5], s[2:3], 0x8
	s_load_b32 s9, s[2:3], 0x0
	v_bfe_u32 v2, v0, 10, 10
	s_load_b32 s2, s[6:7], 0x0
	v_bfe_u32 v0, v0, 20, 10
	s_mov_b32 s12, 0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s0, s0, 16
	v_mul_u32_u24_e32 v2, s1, v2
	s_mul_i32 s0, s0, s1
	s_add_u32 s6, s4, 3
	v_mul_lo_u32 v3, s0, v3
	s_addc_u32 s7, s5, 0
	s_cmp_gt_i32 s9, 0
	s_mul_i32 s11, s2, s11
	s_cselect_b32 s10, -1, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v0, v3, v2, v0
	v_mul_lo_u32 v0, v0, 24
.LBB0_2:                                ; %.preheader
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
                                        ;       Child Loop BB0_5 Depth 3
                                        ;         Child Loop BB0_6 Depth 4
	v_ashrrev_i32_e32 v2, 31, v1
	s_and_not1_b32 vcc_lo, exec_lo, s10
	s_delay_alu instid0(VALU_DEP_1)
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	s_cbranch_vccnz .LBB0_9
; %bb.3:                                ; %.lr.ph
                                        ;   in Loop: Header=BB0_2 Depth=1
	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
	s_mov_b32 s0, 0
	s_mov_b32 s13, 0
	v_mov_b32_e32 v10, 0
	global_load_b32 v4, v[4:5], off
	v_dual_mov_b32 v11, 0xfff00000 :: v_dual_mov_b32 v16, 0
	s_waitcnt vmcnt(0)
	v_and_b32_e32 v5, 0xff, v4
	v_bfe_u32 v6, v4, 8, 8
	v_bfe_u32 v8, v4, 16, 8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_cvt_f64_u32_e32 v[4:5], v5
	v_cvt_f64_u32_e32 v[6:7], v6
	s_delay_alu instid0(VALU_DEP_3)
	v_cvt_f64_u32_e32 v[8:9], v8
.LBB0_4:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_5 Depth 3
                                        ;         Child Loop BB0_6 Depth 4
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, COV@rel32@lo+4
	s_addc_u32 s3, s3, COV@rel32@hi+12
	s_ashr_i32 s1, s0, 31
	s_mul_hi_u32 s16, s13, 24
	s_lshl_b64 s[14:15], s[0:1], 3
	v_mov_b32_e32 v12, 0
	s_add_u32 s1, s14, s2
	s_addc_u32 s14, s15, s3
	s_mul_i32 s15, s13, 24
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, AVG@rel32@lo+4
	s_addc_u32 s3, s3, AVG@rel32@hi+12
	s_add_u32 s2, s15, s2
	s_addc_u32 s3, s16, s3
	s_clause 0x1
	s_load_b128 s[16:19], s[2:3], 0x0
	s_load_b64 s[2:3], s[2:3], 0x10
	v_mov_b32_e32 v13, 0
	s_mov_b32 s15, 0
	s_waitcnt lgkmcnt(0)
	v_add_f64 v[14:15], v[4:5], -s[16:17]
	v_add_f64 v[17:18], v[6:7], -s[18:19]
	v_add_f64 v[19:20], v[8:9], -s[2:3]
	ds_store_2addr_b64 v0, v[14:15], v[17:18] offset1:1
	ds_store_b64 v0, v[19:20] offset:16
.LBB0_5:                                ; %.preheader.i
                                        ;   Parent Loop BB0_2 Depth=1
                                        ;     Parent Loop BB0_4 Depth=2
                                        ; =>    This Loop Header: Depth=3
                                        ;         Child Loop BB0_6 Depth 4
	v_lshl_add_u32 v14, s15, 3, v0
	v_mov_b32_e32 v17, v0
	s_mov_b64 s[2:3], 0
	ds_load_b64 v[14:15], v14
.LBB0_6:                                ;   Parent Loop BB0_2 Depth=1
                                        ;     Parent Loop BB0_4 Depth=2
                                        ;       Parent Loop BB0_5 Depth=3
                                        ; =>      This Inner Loop Header: Depth=4
	s_add_u32 s16, s1, s2
	s_addc_u32 s17, s14, s3
	ds_load_b64 v[20:21], v17
	s_load_b64 s[16:17], s[16:17], 0x0
	v_add_nc_u32_e32 v17, 8, v17
	s_add_u32 s2, s2, 8
	s_addc_u32 s3, s3, 0
	s_cmp_lg_u32 s2, 24
	s_waitcnt lgkmcnt(0)
	v_mul_f64 v[18:19], v[14:15], s[16:17]
	s_delay_alu instid0(VALU_DEP_1)
	v_fma_f64 v[12:13], v[18:19], v[20:21], v[12:13]
	s_cbranch_scc1 .LBB0_6
; %bb.7:                                ;   in Loop: Header=BB0_5 Depth=3
	s_add_i32 s15, s15, 1
	s_add_u32 s1, s1, 24
	s_addc_u32 s14, s14, 0
	s_cmp_lg_u32 s15, 3
	s_cbranch_scc1 .LBB0_5
; %bb.8:                                ; %_Z3dotPKdS0_S0_i.exit
                                        ;   in Loop: Header=BB0_4 Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cmp_lt_f64_e64 vcc_lo, v[10:11], -v[12:13]
	v_xor_b32_e32 v14, 0x80000000, v13
	s_add_i32 s0, s0, 9
	v_dual_cndmask_b32 v11, v11, v14 :: v_dual_cndmask_b32 v10, v10, v12
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_cndmask_b32_e64 v16, v16, s13, vcc_lo
	s_add_i32 s13, s13, 1
	s_cmp_lg_u32 s13, s9
	s_cbranch_scc1 .LBB0_4
	s_branch .LBB0_10
.LBB0_9:                                ;   in Loop: Header=BB0_2 Depth=1
	v_mov_b32_e32 v16, 0
.LBB0_10:                               ; %._crit_edge
                                        ;   in Loop: Header=BB0_2 Depth=1
	v_add_nc_u32_e32 v1, s11, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_co_u32 v2, s0, s6, v2
	v_add_co_ci_u32_e64 v3, s0, s7, v3, s0
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cmp_le_i32_e32 vcc_lo, s8, v1
	global_store_b8 v[2:3], v16, off
	s_or_b32 s12, vcc_lo, s12
	s_and_not1_b32 exec_lo, exec_lo, s12
	s_cbranch_execnz .LBB0_2
.LBB0_11:                               ; %._crit_edge35
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z6kerneliP15HIP_vector_typeIhLj4EEii
		.amdhsa_group_segment_fixed_size 24576
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 1
		.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 2
		.amdhsa_next_free_vgpr 22
		.amdhsa_next_free_sgpr 20
		.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	_Z6kerneliP15HIP_vector_typeIhLj4EEii, .Lfunc_end0-_Z6kerneliP15HIP_vector_typeIhLj4EEii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 708
; NumSgprs: 22
; NumVgprs: 22
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 24576 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 2
; NumSGPRsForWavesPerEU: 22
; NumVGPRsForWavesPerEU: 22
; 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: 2
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	AVG                     ; @AVG
	.type	AVG,@object
	.section	.bss,"aw",@nobits
	.globl	AVG
	.p2align	4, 0x0
AVG:
	.zero	768
	.size	AVG, 768

	.protected	COV                     ; @COV
	.type	COV,@object
	.globl	COV
	.p2align	4, 0x0
COV:
	.zero	2304
	.size	COV, 2304

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.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 AVG
	.addrsig_sym COV
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         20
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 24576
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z6kerneliP15HIP_vector_typeIhLj4EEii
    .private_segment_fixed_size: 0
    .sgpr_count:     22
    .sgpr_spill_count: 0
    .symbol:         _Z6kerneliP15HIP_vector_typeIhLj4EEii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     22
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
