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

#define MAX_DIM 20
#define SCALING_FACTOR 256
#define NUM_THREADS 1024
#define MOD_BASE 10007

int NUM_DIM;
int * MAT_DIM;
int ** MAT_LIST;

int find_max(int * arr, int num_elem)
{
    int max = 0;
    for (int i = 0; i < num_elem; i++) {
        if (arr[i] > max) {
            max = arr[i];
        }
    }
    return max;
}

int * def_mat_dim(int k)
{
    int * dim = (int *) malloc(k * sizeof(int));
    int i;
    srand(time(NULL));

    for (i = 0; i < k; i++)
    {
        dim[i] = (rand() % MAX_DIM) + 1;
        //printf("%d\n", dim[i]);
    }
    return dim;
}

int * equipartition(int k, int blocks) 
{
    float div = float(k) / float(blocks);
    int div_int = int(div);
    float rem = div - float(div_int);

    int * partition = (int *) malloc((blocks + 1) * sizeof(int));
    srand(time(NULL));

    partition[0] = 0;
    partition[blocks] = k-1;

    int cur_index = 0;
    float round_factor = 0.0;
    for (int i = 1; i < blocks; i++) {
        cur_index += div;
        round_factor = float(rand()) / float(RAND_MAX);
        if (round_factor < rem) {
            cur_index += 1;
        }
        partition[i] = cur_index;
    }
    return partition;
}

int * creat_mat(int dimX, int dimY)
{
    int x;
    int * mat = (int *) malloc(dimX * dimY * sizeof(int));

    srand(time(NULL));

    for (x = 0; x < dimX * dimY; x++) {
        mat[x] = rand() % MOD_BASE;
        //mat[x] = (rand() % MAX_DIM) * SCALING_FACTOR;
        //printf("%d ", mat[x]);
    }
    return mat;
}

void if_mats_equal(int * A,  int * B, int rows, int cols)
{
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            if (A[i * cols + j] != B[i * cols + j]) { 
                printf("Matrices are not equal\n"); 
                //printf("%d %d\n", i, j);
                return;
            }
        }
    }
    printf("Matrices are equal\n");
}

void cpu_mat_mul(int* A, int* B, int* C, int ARows, int ACols, int BRows, int BCols)
{
    int sum = 0;
    for (int i = 0; i < ARows; i++) {
        for (int j = 0; j < BCols; j++) {
            for (int k = 0; k < ACols; k++) {
                sum += (A[i * ACols + k] * B[k * BCols + j]) % MOD_BASE;
                //sum += A[i * ACols + k] * B[k * BCols + j];
            }
            //printf("%d %d\n", i, j);
            C[i * BCols + j] = sum % MOD_BASE;
            //C[i * BCols + j] = sum;
            sum = 0;
        }
    }
}

void print_mat(int * mat, int dimX, int dimY)
{
    for (int i = 0; i < dimX; i++) {
        for (int j = 0; j < dimY; j++) {
            printf("%d  ", mat[i * dimY + j]);
        }
        printf("\n");
    }
}

int * cpu_multi_mat_mult(int num_dim, int * dim_list, int ** mat_list, int start = 0) {

    int max_dim = find_max(dim_list, num_dim);
    int * output_mat1 = (int *) calloc(max_dim * max_dim, sizeof(int));
    int * output_mat2 = (int *) calloc(max_dim * max_dim, sizeof(int));

    cpu_mat_mul(mat_list[start], mat_list[start + 1], output_mat1, dim_list[start], dim_list[start + 1], dim_list[start + 1], dim_list[start + 2]);
    int num_rows = dim_list[start];
    int num_cols = dim_list[start + 2];

    //print_mat(output_mat1, num_rows, num_cols);
    int num_mult;
    for (num_mult = 1; num_mult < num_dim - 2; num_mult++) {
        //printf("multiplied %d matrices\n", num_mult + 1);
        if (num_mult % 2 == 1) {
            cpu_mat_mul(output_mat1, mat_list[start + num_mult + 1], output_mat2, num_rows, num_cols, dim_list[start + num_mult + 1] , dim_list[start + num_mult + 2]);
        }
        else {
            cpu_mat_mul(output_mat2, mat_list[start + num_mult + 1], output_mat1, num_rows, num_cols, dim_list[start + num_mult + 1] , dim_list[start + num_mult + 2]);
        }
        num_cols = dim_list[start + num_mult + 2];
    }

    int * output_mat = (int *) malloc(dim_list[0] * dim_list[num_dim - 1] * sizeof(int));
    //printf("%d %d\n", num_rows, num_cols);
    if (num_mult % 2 == 1) {
        memcpy(output_mat, output_mat1, dim_list[0] * dim_list[num_dim - 1] * sizeof(int));
    }
    else {
        memcpy(output_mat, output_mat2, dim_list[0] * dim_list[num_dim - 1] * sizeof(int));
    }
    free(output_mat1);
    free(output_mat2);
    return output_mat;
}

__device__
void matmult(int* A, int* B, int* C, int ARows, int ACols, int BRows, int BCols)
{
    int num_elem_output = ARows * BCols;
    int C_elem_row = 0;
    int C_elem_col = 0;
    int sum = 0;

    for (int n = threadIdx.x; n < num_elem_output; n+=NUM_THREADS) {
        C_elem_col = n % BCols;
        C_elem_row = (n + (BCols - C_elem_col)) / BCols - 1;
        
        for (int i = 0; i < ACols; i++) {
            sum += (A[C_elem_row * ACols + i] * B[i * BCols + C_elem_col]) % MOD_BASE;
            //sum += A[C_elem_row * ACols + i] * B[i * BCols + C_elem_col];
        }
        C[C_elem_row * BCols + C_elem_col] = sum % MOD_BASE;
        //C[C_elem_row * BCols + C_elem_col] = sum;
        sum = 0;
        __syncthreads();
    }
    __syncthreads();
}

__global__
void gpu_seq_multi_matmult(int num_dim, int * dim_list, int ** mat_list, int * output_mat1, int * output_mat2)
{

    matmult(mat_list[0], mat_list[1], output_mat1, dim_list[0], dim_list[1], dim_list[1], dim_list[2]);
    __syncthreads();

    int num_mult;
    int num_rows = dim_list[0];
    int num_cols = dim_list[2];
    for (num_mult = 1; num_mult < num_dim - 2; num_mult++) {
        if (num_mult % 2 == 1) {
            matmult(output_mat1, mat_list[num_mult + 1], output_mat2, num_rows, num_cols, dim_list[num_mult + 1], dim_list[num_mult + 2]);
        } else {
            matmult(output_mat2, mat_list[num_mult + 1], output_mat1, num_rows, num_cols, dim_list[num_mult + 1], dim_list[num_mult + 2]);
        }
        num_cols = dim_list[num_mult + 2];
        __syncthreads();
    }
}

__global__
void gpu_par_multi_matmult(int start_dim_idx, int end_dim_idx, int * dim_list, int ** mat_list, int * output_mat1, int * output_mat2)
{

    matmult(mat_list[start_dim_idx], mat_list[start_dim_idx + 1], output_mat1, dim_list[start_dim_idx], dim_list[start_dim_idx + 1], dim_list[start_dim_idx + 1], dim_list[start_dim_idx + 2]);
    __syncthreads();

    //int num_mult;
    int num_rows = dim_list[start_dim_idx];
    int num_cols = dim_list[start_dim_idx + 2];
//    for (num_mult = start_dim_idx + 1; count < end_dim_idx - start_dim_idx - 2; num_mult++) {
    for (int count = 1; count < (end_dim_idx - start_dim_idx - 2); count++) {
        if (count % 2 == 1) {
            matmult(output_mat1, mat_list[start_dim_idx + count + 1], output_mat2, num_rows, num_cols, dim_list[start_dim_idx + count + 1], dim_list[start_dim_idx + count + 2]);
        } else {
            matmult(output_mat2, mat_list[start_dim_idx + count + 1], output_mat1, num_rows, num_cols, dim_list[start_dim_idx + count + 1], dim_list[start_dim_idx + count + 2]);
            //matmult(output_mat2, mat_list[num_mult + 1], output_mat1, num_rows, num_cols, dim_list[num_mult + 1], dim_list[num_mult + 2]);
        }
        num_cols = dim_list[start_dim_idx + count + 2];
        __syncthreads();
    }
}

int * cpu_algo(int num_dim, int * mat_dim, int ** mat_list)
{

    int * cpu_mat = cpu_multi_mat_mult(num_dim, mat_dim, mat_list);
    //printf("%d %d\n", mat_dim[0], mat_dim[num_dim-1]);
    printf("row: %d  col: %d\n", mat_dim[0], mat_dim[num_dim-1]);
    printf("printing cpu sequential result\n");
    print_mat(cpu_mat, mat_dim[0], mat_dim[num_dim-1]);
    printf("\n");
    return cpu_mat;

}

int * gpu_one_block_algo(int num_dim, int * mat_dim, int ** mat_list)
{

    int num_mat = num_dim - 1;
    int max_dim = find_max(mat_dim, num_dim);

    printf("Copying matrix dimensions to device\n");

    int * d_mat_dim;
    cudaMalloc((void **)&d_mat_dim, num_dim * sizeof(int));
    cudaMemcpy(d_mat_dim, mat_dim, num_dim * sizeof(int), cudaMemcpyHostToDevice);

    printf("Allocating space to store output matrix\n");
    int * out_mat = (int *) calloc(max_dim * max_dim, sizeof(int));
    int * d_out_mat1, * d_out_mat2;
    cudaMalloc((void **) &d_out_mat1, max_dim * max_dim * sizeof(int));
    cudaMalloc((void **) &d_out_mat2, max_dim * max_dim * sizeof(int));
    cudaMemcpy(d_out_mat1, out_mat, max_dim * max_dim * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_out_mat2, out_mat, max_dim * max_dim * sizeof(int), cudaMemcpyHostToDevice);

    printf("Allocating space for each matrix, and storing pointer address of matrices on the host\n");
    int ** int_mat_list = (int **) malloc(num_mat * sizeof(int *));
    for (int k = 0; k < num_mat; k++) {
        cudaMalloc((void **)&int_mat_list[k], mat_dim[k] * mat_dim[k+1] * sizeof(int));
        cudaMemcpy(int_mat_list[k], mat_list[k], mat_dim[k] * mat_dim[k+1] * sizeof(int), cudaMemcpyHostToDevice);
    }

    printf("Copying pointer addresses of matrices from host to device\n");
    int ** d_mat_list;
    cudaMalloc(&d_mat_list, num_mat * sizeof(int *));
    cudaMemcpy(d_mat_list, int_mat_list, num_mat * sizeof(int *), cudaMemcpyHostToDevice);



    gpu_seq_multi_matmult<<<1, NUM_THREADS>>>(num_dim, d_mat_dim, d_mat_list, d_out_mat1, d_out_mat2);
    cudaDeviceSynchronize();


    if (num_dim % 2 == 1) {
        cudaMemcpy(out_mat, d_out_mat1, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int), cudaMemcpyDeviceToHost);
    } else {
        cudaMemcpy(out_mat, d_out_mat2, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int), cudaMemcpyDeviceToHost);
    }

    printf("row: %d  col: %d\n", mat_dim[0], mat_dim[num_dim-1]);
    printf("printing gpu (one thread block) result\n");
    print_mat(out_mat, mat_dim[0], mat_dim[num_dim-1]);
    printf("\n");

    int * output_mat = (int *) malloc(mat_dim[0] * mat_dim[num_dim-1] * sizeof(int)); 
    memcpy(output_mat, out_mat, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int));

    free(out_mat);
    cudaFree(d_mat_dim);
    cudaFree(int_mat_list);
    cudaFree(d_mat_list);
    cudaFree(d_out_mat1);
    cudaFree(d_out_mat2);

    for (int k = 0; k < num_mat; k++) {
        cudaFree(int_mat_list[k]);
    }
    free(int_mat_list);

    return output_mat;
}

int * gpu_multi_block_algo(int num_dim, int * mat_dim, int ** mat_list)
{

    int num_mat = num_dim - 1;
    int nblocks = int(double(num_mat) / double(sqrt(num_mat)));
    int max_dim = find_max(mat_dim, num_dim);
    printf("using %d blocks for %d matrices\n", nblocks, num_mat);

    printf("Copying matrix dimensions to device\n");

    int * d_mat_dim;
    cudaMalloc((void **)&d_mat_dim, num_dim * sizeof(int));
    cudaMemcpy(d_mat_dim, mat_dim, num_dim * sizeof(int), cudaMemcpyHostToDevice);

    printf("Allocating space to store output matrix\n");
    int * out_mat = (int *) calloc(max_dim * max_dim, sizeof(int));
    int * d_out_mat1, * d_out_mat2;
    cudaMalloc((void **) &d_out_mat1, max_dim * max_dim * sizeof(int));
    cudaMalloc((void **) &d_out_mat2, max_dim * max_dim * sizeof(int));
    cudaMemcpy(d_out_mat1, out_mat, max_dim * max_dim * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_out_mat2, out_mat, max_dim * max_dim * sizeof(int), cudaMemcpyHostToDevice);

    printf("Allocating space for each matrix, and storing pointer address of matrices on the host\n");
    int ** int_mat_list = (int **) malloc(num_mat * sizeof(int *));
    for (int k = 0; k < num_mat; k++) {
        cudaMalloc((void **)&int_mat_list[k], mat_dim[k] * mat_dim[k+1] * sizeof(int));
        cudaMemcpy(int_mat_list[k], mat_list[k], mat_dim[k] * mat_dim[k+1] * sizeof(int), cudaMemcpyHostToDevice);
    }

    printf("Copying pointer addresses of matrices from host to device\n");
    int ** d_mat_list;
    cudaMalloc(&d_mat_list, num_mat * sizeof(int *));
    cudaMemcpy(d_mat_list, int_mat_list, num_mat * sizeof(int *), cudaMemcpyHostToDevice);



    printf("Allocating a set of intermediate arrays to store partial results\n");
    int ** int_mat1, ** int_mat2;
    int_mat1 = (int **) malloc(nblocks * sizeof(int *));
    int_mat2 = (int **) malloc(nblocks * sizeof(int *));
    for (int k = 0; k < nblocks; k++) {
        cudaMalloc((void **)&int_mat1[k], max_dim * max_dim * sizeof(int));
        cudaMalloc((void **)&int_mat2[k], max_dim * max_dim * sizeof(int));
        cudaMemcpy(int_mat1[k], out_mat, max_dim * max_dim * sizeof(int), cudaMemcpyHostToDevice);
        cudaMemcpy(int_mat2[k], out_mat, max_dim * max_dim * sizeof(int), cudaMemcpyHostToDevice);
    }

    printf("Copying pointer addresses of intermediate products from host to device\n");
    int ** int_mat_final = (int **) malloc(nblocks * sizeof(int *));
    int ** d_int_mat_final;
    cudaMalloc((void **)&d_int_mat_final, nblocks * sizeof(int *));


    printf("Defining a partition to split work evenly\n");
    int * mat_list_partition = equipartition(num_dim, nblocks);
    //print_mat(mat_list_partition, 1, nblocks + 1);

    printf("Getting the dimensions of the partial products matrices\n");
    //print_mat(mat_dim, 1, num_dim);
    int * int_mat_dim = (int *) malloc((nblocks + 1) * sizeof(int));
    for (int i = 0; i < nblocks + 1; i++) {
        int_mat_dim[i] = mat_dim[mat_list_partition[i]];
    }
    //printf("Moving the dimensions of partial products from host to device\n");
    //print_mat(int_mat_dim, 1, nblocks + 1);
    //printf("\n");

    printf("Moving the dimensions of partial products from host to device\n");
    int * d_int_mat_dim;
    cudaMalloc((void **)&d_int_mat_dim, (nblocks + 1) * sizeof(int));
    cudaMemcpy(d_int_mat_dim, int_mat_dim, (nblocks + 1) * sizeof(int), cudaMemcpyHostToDevice);

    printf("Defining two output arrays which store so that one array storing partial results can be used as input of another output, which is store in the new output array\n");
    int * d_int_output_mat1, * d_int_output_mat2;
    cudaMalloc((void **)&d_int_output_mat1, max_dim * max_dim * sizeof(int));
    cudaMalloc((void **)&d_int_output_mat2, max_dim * max_dim * sizeof(int));
    cudaMemcpy(d_int_output_mat1, out_mat, max_dim * max_dim * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_int_output_mat2, out_mat, max_dim * max_dim * sizeof(int), cudaMemcpyHostToDevice);


    printf("Calling kernels to calculate a subset of input matrices asynchronously\n");
    for (int i = 0; i < nblocks; i++) {

        //printf("nblocks %d, i %d\n", nblocks, i);
        //printf("mat_list_partition %d %d\n", mat_list_partition[i], mat_list_partition[i+1] - 1);
        //printf("end mat_dim %d %d\n", mat_dim[mat_list_partition[i]], mat_dim[mat_list_partition[i+1]]);
        gpu_par_multi_matmult<<<1, NUM_THREADS>>>(mat_list_partition[i], mat_list_partition[i+1] + 1, d_mat_dim, d_mat_list, int_mat1[i], int_mat2[i]);
    }
    cudaDeviceSynchronize();

    printf("output arrays accounted for, will be used to aggregate results\n");
    for (int i = 0; i < nblocks; i++) {
        cudaMemcpy(out_mat, int_mat1[i], mat_dim[mat_list_partition[i]] * mat_dim[mat_list_partition[i+1]] * sizeof(int), cudaMemcpyDeviceToHost);
        if ((mat_list_partition[i+1] - mat_list_partition[i]) % 2 == 1) {
            int_mat_final[i] = int_mat1[i];
        } else {
            int_mat_final[i] = int_mat2[i];
        }
    }

    cudaMemcpy(d_int_mat_final, int_mat_final, nblocks * sizeof(int *), cudaMemcpyHostToDevice);
    
    printf("Calling last kernel to aggregate partial results\n");
    gpu_par_multi_matmult<<<1, NUM_THREADS>>>(0, nblocks + 1, d_int_mat_dim, d_int_mat_final, d_int_output_mat1, d_int_output_mat2); //, d_iii);
    cudaDeviceSynchronize();

    if (nblocks % 2 == 1) {
        cudaMemcpy(out_mat, d_int_output_mat1, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int), cudaMemcpyDeviceToHost);   
    } else {
        cudaMemcpy(out_mat, d_int_output_mat2, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int), cudaMemcpyDeviceToHost);   
    }

    printf("rows:%d  cols: %d\n", mat_dim[mat_list_partition[0]], mat_dim[mat_list_partition[nblocks]]);
    printf("printing gpu (using %d blocks) results\n", nblocks);
    print_mat(out_mat, mat_dim[0], mat_dim[num_dim-1]);
    printf("\n");


    int * output_mat = (int *) malloc(mat_dim[0] * mat_dim[num_dim-1] * sizeof(int)); 
    memcpy(output_mat, out_mat, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int));
    

    free(out_mat);
    cudaFree(d_mat_dim);
    cudaFree(int_mat_list);
    cudaFree(d_mat_list);
    cudaFree(d_out_mat1);
    cudaFree(d_out_mat2);

    for (int k = 0; k < num_mat; k++) {
        cudaFree(int_mat_list[k]);
    }
    free(int_mat_list);

    for (int k = 0; k < nblocks; k++) {
        cudaFree(int_mat1[k]);
        cudaFree(int_mat2[k]);
    }
    free(int_mat1);
    free(int_mat2);

    // STILL MORE THINGS TO FREE


    return output_mat;

}

void read_file(char * filename)//, int * num_dim, int * mat_dim, int ** mat_list)
{

    FILE * fp;
    fp = fopen(filename, "r");

    fscanf(fp, "%d", &NUM_DIM);

    MAT_DIM = (int *) malloc(NUM_DIM * sizeof(int));
    for (int i = 0; i < NUM_DIM; i++) {
        fscanf(fp, "%d", &MAT_DIM[i]);
    }

    MAT_LIST = (int **) malloc((NUM_DIM - 1) * sizeof(int *));

    for (int arr_i = 0; arr_i < (NUM_DIM - 1); arr_i++) {
        MAT_LIST[arr_i] = (int *) malloc(MAT_DIM[arr_i] * MAT_DIM[arr_i + 1] * sizeof(int));
        for (int i = 0; i < (MAT_DIM[arr_i] * MAT_DIM[arr_i + 1]); i++) {
            fscanf(fp, "%d", &MAT_LIST[arr_i][i]);
        }
    }
}


void usage() {

    printf("usage\n ./multi_matmult <filename> <algorithm>\n");
    printf("algorithm:\n\t 0 - cpu (sequential), 1 - gpu (parallel, one thread block), 2 - (parallel, multiple thread blocks)\n");
    //printf("num_blocks : specify only when algorithm == 2. num_blocks should be no more than 1/10th the number of matrices you use\n");

}

int main(int argc, char ** argv)
{
    
    int algorithm = -1;
    int * product_mat;
    char filename[256];

    if (argc == 3) {
        
        strcpy(filename, argv[1]);

        if ((atoi(argv[2]) >= 0) && (atoi(argv[2]) < 3)) {
            algorithm = atoi(argv[2]);
        } else { 
            usage(); 
            return 0;
            }
    } else { 
        usage(); 
        return 0;
        }

    read_file(filename);
/*
    printf("%d\n", NUM_DIM);

    for (int i = 0; i < NUM_DIM; i++) {
        printf("%d ", MAT_DIM[i]);
    }
    printf("\n");

    for (int arr_i = 0; arr_i < (NUM_DIM - 1); arr_i++) {
        for (int i = 0; i < MAT_DIM[arr_i]; i++) {
            for (int j = 0; j < MAT_DIM[arr_i + 1]; j++) {
                printf("%d ", MAT_LIST[arr_i][(i * MAT_DIM[arr_i + 1]) + j]);
            }
            printf("\n");
        }
    }
*/

/*
    for (int i = 0; i < num_dim; i++) {
        printf("%d ", mat_dim[i]);
    }
    printf("\n");
*/

/*
    int num_dim = 45;
    int num_mat = num_dim - 1;
    int * mat_dim = def_mat_dim(num_dim);
    int ** mat_list = (int **) malloc((num_mat) * sizeof(int *));
    int max_dim = find_max(mat_dim, num_dim);

    for (int k = 0; k < num_mat; k++) {
        //printf("================= MATRIX %d ====================\n", k);
        //printf("%d %d\n", mat_dim[k], mat_dim[k+1]);
        mat_list[k] = creat_mat(mat_dim[k], mat_dim[k+1]);
        //printf("%d %d\n", mat_dim[k], mat_dim[k+1]);
        //print_mat(mat_list[k], mat_dim[k], mat_dim[k+1]);
    }
*/    


    if (algorithm == 0) {
        product_mat = cpu_algo(NUM_DIM, MAT_DIM, MAT_LIST);
    } 
    else if (algorithm == 1) {
        product_mat = gpu_one_block_algo(NUM_DIM, MAT_DIM, MAT_LIST);
    } else {
        product_mat = gpu_multi_block_algo(NUM_DIM, MAT_DIM, MAT_LIST);
    }

    return 0;

}

	.file	"tmpxft_0038f332_00000000-6_multi_matmult_int.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2044:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2044:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z8find_maxPii
	.type	_Z8find_maxPii, @function
_Z8find_maxPii:
.LFB2027:
	.cfi_startproc
	endbr64
	xorl	%edx, %edx
	xorl	%eax, %eax
.L3:
	cmpl	%edx, %esi
	jle	.L6
	movl	(%rdi,%rdx,4), %ecx
	cmpl	%ecx, %eax
	cmovl	%ecx, %eax
	incq	%rdx
	jmp	.L3
.L6:
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z8find_maxPii, .-_Z8find_maxPii
	.globl	_Z11def_mat_dimi
	.type	_Z11def_mat_dimi, @function
_Z11def_mat_dimi:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movslq	%edi, %rdi
	movl	$20, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rdi, %r12
	salq	$2, %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	pushq	%rcx
	.cfi_def_cfa_offset 48
	call	malloc@PLT
	xorl	%edi, %edi
	movq	%rax, %rbx
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
.L8:
	cmpl	%ebp, %r12d
	jle	.L11
	call	rand@PLT
	cltd
	idivl	%r13d
	incl	%edx
	movl	%edx, (%rbx,%rbp,4)
	incq	%rbp
	jmp	.L8
.L11:
	popq	%rdx
	.cfi_def_cfa_offset 40
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z11def_mat_dimi, .-_Z11def_mat_dimi
	.globl	_Z13equipartitionii
	.type	_Z13equipartitionii, @function
_Z13equipartitionii:
.LFB2029:
	.cfi_startproc
	endbr64
	cvtsi2ssl	%esi, %xmm0
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	cvtsi2ssl	%edi, %xmm1
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movl	%esi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movl	%edi, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	leal	1(%rsi), %ebp
	decl	%r12d
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movslq	%ebp, %rbp
	divss	%xmm0, %xmm1
	salq	$2, %rbp
	movq	%rbp, %rdi
	subq	$16, %rsp
	.cfi_def_cfa_offset 64
	cvttss2sil	%xmm1, %eax
	movaps	%xmm1, %xmm2
	movss	%xmm1, 12(%rsp)
	cvtsi2ssl	%eax, %xmm0
	subss	%xmm0, %xmm2
	movd	%xmm2, %r14d
	call	malloc@PLT
	xorl	%edi, %edi
	movq	%rax, %rbx
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	xorl	%r9d, %r9d
	movl	%r9d, (%rbx)
	movl	%r12d, -4(%rbx,%rbp)
	movl	$1, %r12d
	xorl	%ebp, %ebp
.L13:
	cmpl	%r12d, %r13d
	jle	.L19
	cvtsi2ssl	%ebp, %xmm0
	addss	12(%rsp), %xmm0
	cvttss2sil	%xmm0, %ebp
	call	rand@PLT
	movd	%r14d, %xmm4
	cvtsi2ssl	%eax, %xmm0
	mulss	.LC0(%rip), %xmm0
	comiss	%xmm0, %xmm4
	jbe	.L14
	incl	%ebp
.L14:
	movl	%ebp, (%rbx,%r12,4)
	incq	%r12
	jmp	.L13
.L19:
	addq	$16, %rsp
	.cfi_def_cfa_offset 48
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z13equipartitionii, .-_Z13equipartitionii
	.globl	_Z9creat_matii
	.type	_Z9creat_matii, @function
_Z9creat_matii:
.LFB2030:
	.cfi_startproc
	endbr64
	imull	%esi, %edi
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	movl	$10007, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movl	%edi, %r12d
	movslq	%edi, %rdi
	pushq	%rcx
	.cfi_def_cfa_offset 48
	salq	$2, %rdi
	call	malloc@PLT
	xorl	%edi, %edi
	movq	%rax, %rbx
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
.L21:
	cmpl	%ebp, %r12d
	jle	.L24
	call	rand@PLT
	cltd
	idivl	%r13d
	movl	%edx, (%rbx,%rbp,4)
	incq	%rbp
	jmp	.L21
.L24:
	popq	%rdx
	.cfi_def_cfa_offset 40
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2030:
	.size	_Z9creat_matii, .-_Z9creat_matii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Matrices are not equal\n"
.LC2:
	.string	"Matrices are equal\n"
	.text
	.globl	_Z13if_mats_equalPiS_ii
	.type	_Z13if_mats_equalPiS_ii, @function
_Z13if_mats_equalPiS_ii:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	%rdi, %r10
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
.L26:
	cmpl	%edx, %r9d
	jge	.L29
	movslq	%r8d, %rax
	xorl	%edi, %edi
	salq	$2, %rax
.L30:
	cmpl	%ecx, %edi
	jge	.L32
	movl	(%r10,%rax), %r11d
	addq	$4, %rax
	cmpl	-4(%rsi,%rax), %r11d
	je	.L27
	leaq	.LC1(%rip), %rsi
	jmp	.L31
.L27:
	incl	%edi
	jmp	.L30
.L32:
	incl	%r9d
	addl	%ecx, %r8d
	jmp	.L26
.L29:
	leaq	.LC2(%rip), %rsi
.L31:
	movl	$2, %edi
	xorl	%eax, %eax
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2031:
	.size	_Z13if_mats_equalPiS_ii, .-_Z13if_mats_equalPiS_ii
	.globl	_Z11cpu_mat_mulPiS_S_iiii
	.type	_Z11cpu_mat_mulPiS_S_iiii, @function
_Z11cpu_mat_mulPiS_S_iiii:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%r9d, %r9d
	xorl	%r11d, %r11d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%ecx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movslq	56(%rsp), %rax
	xorl	%ebx, %ebx
	salq	$2, %rax
	movq	%rdi, -16(%rsp)
	movq	%rdx, -8(%rsp)
	movq	%rax, -24(%rsp)
.L34:
	cmpl	%r12d, %ebx
	jge	.L33
	movq	-8(%rsp), %rdi
	movslq	%r11d, %rax
	movq	-16(%rsp), %rdx
	xorl	%ecx, %ecx
	leaq	(%rdi,%rax,4), %r15
	movslq	%r9d, %rax
	leaq	(%rdx,%rax,4), %r14
.L39:
	cmpl	%ecx, 56(%rsp)
	jle	.L36
	leaq	0(%rbp,%rcx,4), %r13
	xorl	%edi, %edi
	xorl	%esi, %esi
.L37:
	cmpl	%edi, %r8d
	jle	.L41
	movl	(%r14,%rdi,4), %eax
	imull	0(%r13), %eax
	movl	$10007, %r10d
	incq	%rdi
	cltd
	idivl	%r10d
	movq	-24(%rsp), %rax
	addq	%rax, %r13
	addl	%edx, %esi
	jmp	.L37
.L41:
	movl	%esi, %eax
	movl	$10007, %esi
	cltd
	idivl	%esi
	movl	%edx, (%r15,%rcx,4)
	incq	%rcx
	jmp	.L39
.L36:
	incl	%ebx
	addl	56(%rsp), %r11d
	addl	%r8d, %r9d
	jmp	.L34
.L33:
	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
.LFE2032:
	.size	_Z11cpu_mat_mulPiS_S_iiii, .-_Z11cpu_mat_mulPiS_S_iiii
	.section	.rodata.str1.1
.LC3:
	.string	"%d  "
.LC4:
	.string	"\n"
	.text
	.globl	_Z9print_matPiii
	.type	_Z9print_matPiii, @function
_Z9print_matPiii:
.LFB2033:
	.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
	movl	%esi, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	xorl	%r13d, %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
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%edx, %ebx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
.L43:
	cmpl	%r14d, %r13d
	jge	.L42
	movslq	%r12d, %rax
	xorl	%r15d, %r15d
	leaq	0(%rbp,%rax,4), %r8
.L46:
	cmpl	%r15d, %ebx
	jle	.L48
	movl	(%r8,%r15,4), %edx
	leaq	.LC3(%rip), %rsi
	xorl	%eax, %eax
	incq	%r15
	movl	$2, %edi
	movq	%r8, 8(%rsp)
	call	__printf_chk@PLT
	movq	8(%rsp), %r8
	jmp	.L46
.L48:
	leaq	.LC4(%rip), %rsi
	xorl	%eax, %eax
	incl	%r13d
	addl	%ebx, %r12d
	movl	$2, %edi
	call	__printf_chk@PLT
	jmp	.L43
.L42:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2033:
	.size	_Z9print_matPiii, .-_Z9print_matPiii
	.globl	_Z18cpu_multi_mat_multiPiPS_i
	.type	_Z18cpu_multi_mat_multiPiPS_i, @function
_Z18cpu_multi_mat_multiPiPS_i:
.LFB2034:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$1, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movslq	%edi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	movl	%ebp, %esi
	movq	%rbx, %rdi
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movl	%ecx, 36(%rsp)
	call	_Z8find_maxPii
	movl	$4, %esi
	imull	%eax, %eax
	movslq	%eax, %r14
	movq	%r14, %rdi
	call	calloc@PLT
	movl	$4, %esi
	movq	%r14, %rdi
	movq	%rax, 8(%rsp)
	call	calloc@PLT
	movslq	36(%rsp), %rdx
	movq	%rax, 16(%rsp)
	leaq	2(%rdx), %rax
	incq	%rdx
	leaq	0(,%rax,4), %rcx
	movl	(%rbx,%rax,4), %r11d
	movq	(%r12,%rdx,8), %rsi
	movl	-8(%rbx,%rcx), %eax
	movl	-4(%rbx,%rcx), %r8d
	leaq	-8(%rcx), %r13
	movq	-8(%r12,%rdx,8), %rdi
	movl	%eax, 24(%rsp)
	pushq	%r9
	.cfi_def_cfa_offset 120
	movl	%r8d, %r9d
	pushq	%r11
	.cfi_def_cfa_offset 128
	movl	40(%rsp), %ecx
	movq	24(%rsp), %rdx
	movl	%r11d, 60(%rsp)
	call	_Z11cpu_mat_mulPiS_S_iiii
	popq	%r10
	.cfi_def_cfa_offset 120
	leal	-2(%rbp), %eax
	popq	%r11
	.cfi_def_cfa_offset 112
	movl	%eax, 40(%rsp)
	movl	44(%rsp), %r11d
.L50:
	movl	%r15d, %r14d
	andl	$1, %r14d
	cmpl	%r15d, 40(%rsp)
	jle	.L59
	testl	%r14d, %r14d
	je	.L51
	movl	36(%rsp), %eax
	movl	%r11d, %r8d
	leal	(%rax,%r15), %edx
	movslq	%edx, %rdx
	leaq	2(%rdx), %rcx
	movq	8(%r12,%rdx,8), %rsi
	movl	-4(%rbx,%rcx,4), %r9d
	pushq	%rdi
	.cfi_remember_state
	.cfi_def_cfa_offset 120
	movl	(%rbx,%rcx,4), %edx
	pushq	%rdx
	.cfi_def_cfa_offset 128
	movl	40(%rsp), %ecx
	movq	32(%rsp), %rdx
	movq	24(%rsp), %rdi
	jmp	.L57
.L51:
	.cfi_restore_state
	movl	8(%rbx,%r13), %r9d
	movq	16(%r12,%r13,2), %rsi
	pushq	%rcx
	.cfi_def_cfa_offset 120
	movl	%r11d, %r8d
	movl	12(%rbx,%r13), %edx
	pushq	%rdx
	.cfi_def_cfa_offset 128
	movl	40(%rsp), %ecx
	movq	24(%rsp), %rdx
	movq	32(%rsp), %rdi
.L57:
	call	_Z11cpu_mat_mulPiS_S_iiii
	incl	%r15d
	movl	12(%rbx,%r13), %r11d
	popq	%rax
	.cfi_def_cfa_offset 120
	addq	$4, %r13
	popq	%rdx
	.cfi_def_cfa_offset 112
	jmp	.L50
.L59:
	leaq	-4(%rbx,%rbp,4), %r12
	movl	(%rbx), %ecx
	imull	(%r12), %ecx
	movslq	%ecx, %rcx
	salq	$2, %rcx
	movq	%rcx, %rdi
	movq	%rcx, 24(%rsp)
	call	malloc@PLT
	testl	%r14d, %r14d
	movl	(%rbx), %edx
	movq	24(%rsp), %rcx
	movq	%rax, %rbp
	je	.L54
	imull	(%r12), %edx
	movq	8(%rsp), %rsi
	movslq	%edx, %rdx
	salq	$2, %rdx
	jmp	.L58
.L54:
	imull	(%r12), %edx
	movq	16(%rsp), %rsi
	movslq	%edx, %rdx
	salq	$2, %rdx
.L58:
	movq	%rbp, %rdi
	call	__memcpy_chk@PLT
	movq	8(%rsp), %rdi
	call	free@PLT
	movq	16(%rsp), %rdi
	call	free@PLT
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbp, %rax
	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
.LFE2034:
	.size	_Z18cpu_multi_mat_multiPiPS_i, .-_Z18cpu_multi_mat_multiPiPS_i
	.globl	_Z7matmultPiS_S_iiii
	.type	_Z7matmultPiS_S_iiii, @function
_Z7matmultPiS_S_iiii:
.LFB2035:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2035:
	.size	_Z7matmultPiS_S_iiii, .-_Z7matmultPiS_S_iiii
	.section	.rodata.str1.1
.LC5:
	.string	"row: %d  col: %d\n"
.LC6:
	.string	"printing cpu sequential result\n"
	.text
	.globl	_Z8cpu_algoiPiPS_
	.type	_Z8cpu_algoiPiPS_, @function
_Z8cpu_algoiPiPS_:
.LFB2036:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	xorl	%ecx, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movslq	%edi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rsi, %rbx
	movl	%ebp, %edi
	leaq	-4(%rbx,%rbp,4), %rbp
	call	_Z18cpu_multi_mat_multiPiPS_i
	movl	0(%rbp), %ecx
	movl	(%rbx), %edx
	movl	$2, %edi
	movq	%rax, %r12
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	0(%rbp), %edx
	movl	(%rbx), %esi
	movq	%r12, %rdi
	call	_Z9print_matPiii
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2036:
	.size	_Z8cpu_algoiPiPS_, .-_Z8cpu_algoiPiPS_
	.section	.rodata.str1.1
.LC7:
	.string	"r"
.LC8:
	.string	"%d"
	.text
	.globl	_Z9read_filePc
	.type	_Z9read_filePc, @function
_Z9read_filePc:
.LFB2039:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC7(%rip), %rsi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	leaq	.LC8(%rip), %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	call	fopen@PLT
	leaq	NUM_DIM(%rip), %rdx
	movq	%rbp, %rsi
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	call	__isoc23_fscanf@PLT
	movslq	NUM_DIM(%rip), %rdi
	salq	$2, %rdi
	call	malloc@PLT
	movq	%rax, MAT_DIM(%rip)
.L65:
	movl	NUM_DIM(%rip), %edi
	cmpl	%r12d, %edi
	jle	.L72
	movq	MAT_DIM(%rip), %rax
	movq	8(%rsp), %rdi
	movq	%rbp, %rsi
	leaq	(%rax,%r12,4), %rdx
	xorl	%eax, %eax
	incq	%r12
	call	__isoc23_fscanf@PLT
	jmp	.L65
.L72:
	decl	%edi
	xorl	%r15d, %r15d
	movslq	%edi, %rdi
	salq	$3, %rdi
	call	malloc@PLT
	movq	%rax, MAT_LIST(%rip)
.L67:
	movl	NUM_DIM(%rip), %eax
	decl	%eax
	cmpl	%r15d, %eax
	jle	.L73
	movq	MAT_DIM(%rip), %rax
	leaq	0(,%r15,4), %rbx
	movq	MAT_LIST(%rip), %r14
	leaq	0(,%r15,8), %r12
	leaq	4(%rbx), %r13
	movl	(%rax,%r15,4), %edi
	imull	4(%rax,%rbx), %edi
	addq	%r12, %r14
	movslq	%edi, %rdi
	salq	$2, %rdi
	call	malloc@PLT
	movq	%rax, (%r14)
	xorl	%r14d, %r14d
.L68:
	movq	MAT_DIM(%rip), %rdx
	movl	(%rdx,%rbx), %eax
	imull	(%rdx,%r13), %eax
	cmpl	%r14d, %eax
	jle	.L74
	movq	MAT_LIST(%rip), %rax
	movq	8(%rsp), %rdi
	movq	%rbp, %rsi
	movq	(%rax,%r12), %rax
	leaq	(%rax,%r14,4), %rdx
	xorl	%eax, %eax
	incq	%r14
	call	__isoc23_fscanf@PLT
	jmp	.L68
.L74:
	incq	%r15
	jmp	.L67
.L73:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2039:
	.size	_Z9read_filePc, .-_Z9read_filePc
	.section	.rodata.str1.1
.LC9:
	.string	"usage\n ./multi_matmult <filename> <algorithm>\n"
.LC10:
	.string	"algorithm:\n\t 0 - cpu (sequential), 1 - gpu (parallel, one thread block), 2 - (parallel, multiple thread blocks)\n"
	.text
	.globl	_Z5usagev
	.type	_Z5usagev, @function
_Z5usagev:
.LFB2040:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC10(%rip), %rsi
	xorl	%eax, %eax
	popq	%rdx
	.cfi_def_cfa_offset 8
	movl	$2, %edi
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2040:
	.size	_Z5usagev, .-_Z5usagev
	.globl	_Z49__device_stub__Z21gpu_seq_multi_matmultiPiPS_S_S_iPiPS_S_S_
	.type	_Z49__device_stub__Z21gpu_seq_multi_matmultiPiPS_S_S_iPiPS_S_S_, @function
_Z49__device_stub__Z21gpu_seq_multi_matmultiPiPS_S_S_iPiPS_S_S_:
.LFB2066:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rcx, 16(%rsp)
	leaq	56(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L77
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z21gpu_seq_multi_matmultiPiPS_S_S_(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 192
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	128(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L77:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L79
	call	__stack_chk_fail@PLT
.L79:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2066:
	.size	_Z49__device_stub__Z21gpu_seq_multi_matmultiPiPS_S_S_iPiPS_S_S_, .-_Z49__device_stub__Z21gpu_seq_multi_matmultiPiPS_S_S_iPiPS_S_S_
	.globl	_Z21gpu_seq_multi_matmultiPiPS_S_S_
	.type	_Z21gpu_seq_multi_matmultiPiPS_S_S_, @function
_Z21gpu_seq_multi_matmultiPiPS_S_S_:
.LFB2067:
	.cfi_startproc
	endbr64
	jmp	_Z49__device_stub__Z21gpu_seq_multi_matmultiPiPS_S_S_iPiPS_S_S_
	.cfi_endproc
.LFE2067:
	.size	_Z21gpu_seq_multi_matmultiPiPS_S_S_, .-_Z21gpu_seq_multi_matmultiPiPS_S_S_
	.section	.rodata.str1.1
.LC11:
	.string	"Copying matrix dimensions to device\n"
.LC12:
	.string	"Allocating space to store output matrix\n"
.LC13:
	.string	"Allocating space for each matrix, and storing pointer address of matrices on the host\n"
.LC14:
	.string	"Copying pointer addresses of matrices from host to device\n"
.LC15:
	.string	"printing gpu (one thread block) result\n"
	.text
	.globl	_Z18gpu_one_block_algoiPiPS_
	.type	_Z18gpu_one_block_algoiPiPS_, @function
_Z18gpu_one_block_algoiPiPS_:
.LFB2037:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdx, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	leal	-1(%rdi), %r14d
	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
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	movl	%edi, %esi
	subq	$104, %rsp
	.cfi_def_cfa_offset 160
	movl	%edi, 12(%rsp)
	movq	%rbx, %rdi
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	call	_Z8find_maxPii
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	movl	%eax, %r12d
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movslq	12(%rsp), %rax
	leaq	32(%rsp), %rdi
	imull	%r12d, %r12d
	salq	$2, %rax
	movq	%rax, %rsi
	movq	%rax, 16(%rsp)
	call	cudaMalloc@PLT
	movq	16(%rsp), %rdx
	movq	%rbx, %rsi
	movslq	%r12d, %r12
	movq	32(%rsp), %rdi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%r12, %rdi
	movl	$4, %esi
	salq	$2, %r12
	call	calloc@PLT
	movq	%r12, %rsi
	leaq	40(%rsp), %rdi
	movq	%rax, (%rsp)
	call	cudaMalloc@PLT
	movq	%r12, %rsi
	leaq	48(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	movq	40(%rsp), %rdi
	movq	%r12, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	(%rsp), %rsi
	movq	48(%rsp), %rdi
	movq	%r12, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movslq	%r14d, %rax
	salq	$3, %rax
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
	call	malloc@PLT
	movq	%rax, %r12
	movq	%rax, %r13
.L83:
	cmpl	%ebp, %r14d
	jle	.L93
	movl	(%rbx,%rbp,4), %esi
	imull	4(%rbx,%rbp,4), %esi
	movq	%r13, %rdi
	addq	$8, %r13
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	(%rbx,%rbp,4), %edx
	imull	4(%rbx,%rbp,4), %edx
	movl	$1, %ecx
	movq	(%r15,%rbp,8), %rsi
	movq	-8(%r13), %rdi
	incq	%rbp
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	jmp	.L83
.L93:
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	24(%rsp), %rsi
	leaq	56(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	24(%rsp), %rdx
	movq	56(%rsp), %rdi
	movq	%r12, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$10, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L85
	movq	48(%rsp), %r8
	movq	40(%rsp), %rcx
	movq	56(%rsp), %rdx
	movq	32(%rsp), %rsi
	movl	12(%rsp), %edi
	call	_Z49__device_stub__Z21gpu_seq_multi_matmultiPiPS_S_S_iPiPS_S_S_
.L85:
	call	cudaDeviceSynchronize@PLT
	movq	16(%rsp), %rax
	movl	$2, %ecx
	leaq	-4(%rbx,%rax), %r13
	movl	12(%rsp), %eax
	cltd
	idivl	%ecx
	decl	%edx
	movl	(%rbx), %edx
	jne	.L86
	imull	0(%r13), %edx
	movq	40(%rsp), %rsi
	movslq	%edx, %rdx
	salq	$2, %rdx
	jmp	.L92
.L86:
	imull	0(%r13), %edx
	movq	48(%rsp), %rsi
	movl	$2, %ecx
	movslq	%edx, %rdx
	salq	$2, %rdx
.L92:
	movq	(%rsp), %rdi
	call	cudaMemcpy@PLT
	movl	0(%r13), %ecx
	movl	(%rbx), %edx
	movl	$2, %edi
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	0(%r13), %edx
	movl	(%rbx), %esi
	movq	(%rsp), %rdi
	call	_Z9print_matPiii
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	(%rbx), %r15d
	imull	0(%r13), %r15d
	movslq	%r15d, %r15
	salq	$2, %r15
	movq	%r15, %rdi
	call	malloc@PLT
	movl	(%rbx), %edx
	imull	0(%r13), %edx
	movq	%r15, %rcx
	movq	(%rsp), %rsi
	movq	%rax, %rdi
	movq	%rax, %rbp
	xorl	%ebx, %ebx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	__memcpy_chk@PLT
	movq	(%rsp), %rdi
	call	free@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	%r12, %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
.L88:
	cmpl	%ebx, %r14d
	jle	.L94
	movq	(%r12,%rbx,8), %rdi
	incq	%rbx
	call	cudaFree@PLT
	jmp	.L88
.L94:
	movq	%r12, %rdi
	call	free@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L90
	call	__stack_chk_fail@PLT
.L90:
	addq	$104, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbp, %rax
	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
.LFE2037:
	.size	_Z18gpu_one_block_algoiPiPS_, .-_Z18gpu_one_block_algoiPiPS_
	.globl	_Z50__device_stub__Z21gpu_par_multi_matmultiiPiPS_S_S_iiPiPS_S_S_
	.type	_Z50__device_stub__Z21gpu_par_multi_matmultiiPiPS_S_S_iiPiPS_S_S_, @function
_Z50__device_stub__Z21gpu_par_multi_matmultiiPiPS_S_S_iiPiPS_S_S_:
.LFB2068:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movl	%edi, 44(%rsp)
	leaq	72(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 32(%rsp)
	leaq	56(%rsp), %rdx
	movq	%rcx, 24(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 16(%rsp)
	movq	%r9, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L95
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z21gpu_par_multi_matmultiiPiPS_S_S_(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L95:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L97
	call	__stack_chk_fail@PLT
.L97:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2068:
	.size	_Z50__device_stub__Z21gpu_par_multi_matmultiiPiPS_S_S_iiPiPS_S_S_, .-_Z50__device_stub__Z21gpu_par_multi_matmultiiPiPS_S_S_iiPiPS_S_S_
	.globl	_Z21gpu_par_multi_matmultiiPiPS_S_S_
	.type	_Z21gpu_par_multi_matmultiiPiPS_S_S_, @function
_Z21gpu_par_multi_matmultiiPiPS_S_S_:
.LFB2069:
	.cfi_startproc
	endbr64
	jmp	_Z50__device_stub__Z21gpu_par_multi_matmultiiPiPS_S_S_iiPiPS_S_S_
	.cfi_endproc
.LFE2069:
	.size	_Z21gpu_par_multi_matmultiiPiPS_S_S_, .-_Z21gpu_par_multi_matmultiiPiPS_S_S_
	.section	.rodata.str1.1
.LC16:
	.string	"using %d blocks for %d matrices\n"
.LC17:
	.string	"Allocating a set of intermediate arrays to store partial results\n"
.LC18:
	.string	"Copying pointer addresses of intermediate products from host to device\n"
.LC19:
	.string	"Defining a partition to split work evenly\n"
.LC20:
	.string	"Getting the dimensions of the partial products matrices\n"
.LC21:
	.string	"Moving the dimensions of partial products from host to device\n"
.LC22:
	.string	"Defining two output arrays which store so that one array storing partial results can be used as input of another output, which is store in the new output array\n"
.LC23:
	.string	"Calling kernels to calculate a subset of input matrices asynchronously\n"
.LC24:
	.string	"output arrays accounted for, will be used to aggregate results\n"
.LC25:
	.string	"Calling last kernel to aggregate partial results\n"
.LC26:
	.string	"rows:%d  cols: %d\n"
.LC27:
	.string	"printing gpu (using %d blocks) results\n"
	.text
	.globl	_Z20gpu_multi_block_algoiPiPS_
	.type	_Z20gpu_multi_block_algoiPiPS_, @function
_Z20gpu_multi_block_algoiPiPS_:
.LFB2038:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%r15d, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	%edi, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	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	$168, %rsp
	.cfi_def_cfa_offset 224
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leal	-1(%rdi), %eax
	cvtsi2sdl	%eax, %xmm2
	movl	%eax, 56(%rsp)
	movaps	%xmm2, %xmm0
	movq	%xmm2, %rbp
	call	sqrt@PLT
	movl	%r13d, %esi
	movq	%rbx, %rdi
	movaps	%xmm0, %xmm1
	movq	%rbp, %xmm0
	divsd	%xmm1, %xmm0
	cvttsd2sil	%xmm0, %eax
	movl	%eax, 20(%rsp)
	call	_Z8find_maxPii
	movl	20(%rsp), %edx
	leal	-1(%r13), %ecx
	leaq	.LC16(%rip), %rsi
	movl	%eax, %ebp
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$2, %edi
	imull	%ebp, %ebp
	xorl	%eax, %eax
	leaq	.LC11(%rip), %rsi
	call	__printf_chk@PLT
	movslq	%r13d, %rax
	leaq	64(%rsp), %rdi
	salq	$2, %rax
	movslq	%ebp, %rbp
	movq	%rax, %rsi
	movq	%rax, 48(%rsp)
	call	cudaMalloc@PLT
	movq	48(%rsp), %rdx
	movq	64(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbp, %rdi
	movl	$4, %esi
	call	calloc@PLT
	leaq	72(%rsp), %rdi
	movq	%rax, 8(%rsp)
	leaq	0(,%rbp,4), %rax
	movq	%rax, %rsi
	movq	%rax, (%rsp)
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	leaq	80(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	$1, %ecx
	movq	72(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	$1, %ecx
	movq	80(%rsp), %rdi
	call	cudaMemcpy@PLT
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movslq	56(%rsp), %rbp
	salq	$3, %rbp
	movq	%rbp, %rdi
	call	malloc@PLT
	movq	%rax, 32(%rsp)
	movq	%rax, %r12
.L101:
	cmpl	%r15d, 56(%rsp)
	jle	.L124
	movl	(%rbx,%r15,4), %esi
	imull	4(%rbx,%r15,4), %esi
	movq	%r12, %rdi
	addq	$8, %r12
	movslq	%esi, %rsi
	salq	$2, %rsi
	call	cudaMalloc@PLT
	movl	(%rbx,%r15,4), %edx
	imull	4(%rbx,%r15,4), %edx
	movl	$1, %ecx
	movq	(%r14,%r15,8), %rsi
	movq	-8(%r12), %rdi
	incq	%r15
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	jmp	.L101
.L124:
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	leaq	88(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	32(%rsp), %rsi
	movq	88(%rsp), %rdi
	movq	%rbp, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	leaq	.LC17(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movslq	20(%rsp), %rax
	salq	$3, %rax
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
	call	malloc@PLT
	movq	24(%rsp), %rdi
	movq	%rax, %rbp
	call	malloc@PLT
	movq	%rbp, %r14
	xorl	%r8d, %r8d
	movq	%rax, %r12
	movq	%rax, %r15
.L103:
	movl	20(%rsp), %eax
	cmpl	%eax, %r8d
	jge	.L125
	movq	(%rsp), %rsi
	movq	%r14, %rdi
	movl	%r8d, 40(%rsp)
	addq	$8, %r14
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	movq	%r15, %rdi
	addq	$8, %r15
	call	cudaMalloc@PLT
	movq	-8(%r14), %rdi
	movq	(%rsp), %rdx
	movl	$1, %ecx
	movq	8(%rsp), %rsi
	call	cudaMemcpy@PLT
	movq	-8(%r15), %rdi
	movq	(%rsp), %rdx
	movl	$1, %ecx
	movq	8(%rsp), %rsi
	call	cudaMemcpy@PLT
	movl	40(%rsp), %r8d
	incl	%r8d
	jmp	.L103
.L125:
	leaq	.LC18(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	24(%rsp), %rdi
	call	malloc@PLT
	movq	24(%rsp), %rsi
	leaq	96(%rsp), %rdi
	movq	%rax, %r15
	call	cudaMalloc@PLT
	leaq	.LC19(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	20(%rsp), %esi
	movl	%r13d, %edi
	call	_Z13equipartitionii
	leaq	.LC20(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %r13
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	20(%rsp), %eax
	incl	%eax
	movl	%eax, 60(%rsp)
	cltq
	salq	$2, %rax
	movq	%rax, %rdi
	movq	%rax, 40(%rsp)
	call	malloc@PLT
	movq	%rax, %r14
	xorl	%eax, %eax
.L105:
	cmpl	%eax, 20(%rsp)
	jl	.L126
	movslq	0(%r13,%rax,4), %rdx
	movl	(%rbx,%rdx,4), %edx
	movl	%edx, (%r14,%rax,4)
	incq	%rax
	jmp	.L105
.L126:
	leaq	.LC21(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	40(%rsp), %rsi
	leaq	104(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	40(%rsp), %rdx
	movq	104(%rsp), %rdi
	movq	%r14, %rsi
	movl	$1, %ecx
	xorl	%r14d, %r14d
	call	cudaMemcpy@PLT
	leaq	.LC22(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	(%rsp), %rsi
	leaq	112(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	leaq	120(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	$1, %ecx
	movq	112(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	8(%rsp), %rsi
	movl	$1, %ecx
	movq	120(%rsp), %rdi
	call	cudaMemcpy@PLT
	leaq	.LC23(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L107:
	cmpl	%r14d, 20(%rsp)
	jle	.L127
	movl	$4194305, %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	salq	$10, %rax
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movl	$1, 148(%rsp)
	movq	%rax, %rdx
	movq	%rax, 140(%rsp)
	movq	%rdi, 128(%rsp)
	movl	$1, 136(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L108
	movl	4(%r13,%r14,4), %eax
	movl	0(%r13,%r14,4), %edi
	movq	(%r12,%r14,8), %r9
	movq	0(%rbp,%r14,8), %r8
	movq	88(%rsp), %rcx
	movq	64(%rsp), %rdx
	leal	1(%rax), %esi
	call	_Z50__device_stub__Z21gpu_par_multi_matmultiiPiPS_S_S_iiPiPS_S_S_
.L108:
	incq	%r14
	jmp	.L107
.L127:
	call	cudaDeviceSynchronize@PLT
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%r14d, %r14d
	leaq	.LC24(%rip), %rsi
	call	__printf_chk@PLT
.L110:
	cmpl	%r14d, 20(%rsp)
	jle	.L128
	movslq	0(%r13,%r14,4), %rdx
	movslq	4(%r13,%r14,4), %rax
	movl	$2, %ecx
	movq	0(%rbp,%r14,8), %rsi
	movq	8(%rsp), %rdi
	movl	(%rbx,%rdx,4), %edx
	imull	(%rbx,%rax,4), %edx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	cudaMemcpy@PLT
	movl	4(%r13,%r14,4), %eax
	subl	0(%r13,%r14,4), %eax
	movl	$2, %ecx
	cltd
	idivl	%ecx
	decl	%edx
	jne	.L111
	movq	0(%rbp,%r14,8), %rax
	jmp	.L112
.L111:
	movq	(%r12,%r14,8), %rax
.L112:
	movq	%rax, (%r15,%r14,8)
	incq	%r14
	jmp	.L110
.L128:
	movq	24(%rsp), %rdx
	movq	96(%rsp), %rdi
	movl	$1, %ecx
	movq	%r15, %rsi
	call	cudaMemcpy@PLT
	leaq	.LC25(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$10, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	movq	%rdx, 140(%rsp)
	movl	$1, 148(%rsp)
	movq	%rdi, 128(%rsp)
	movl	$1, 136(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L114
	movq	120(%rsp), %r9
	movq	112(%rsp), %r8
	xorl	%edi, %edi
	movq	96(%rsp), %rcx
	movq	104(%rsp), %rdx
	movl	60(%rsp), %esi
	call	_Z50__device_stub__Z21gpu_par_multi_matmultiiPiPS_S_S_iiPiPS_S_S_
.L114:
	call	cudaDeviceSynchronize@PLT
	movq	48(%rsp), %rax
	movl	$2, %ecx
	leaq	-4(%rbx,%rax), %r14
	movl	20(%rsp), %eax
	cltd
	idivl	%ecx
	decl	%edx
	movl	(%rbx), %edx
	jne	.L115
	imull	(%r14), %edx
	movq	112(%rsp), %rsi
	movslq	%edx, %rdx
	salq	$2, %rdx
	jmp	.L123
.L115:
	imull	(%r14), %edx
	movq	120(%rsp), %rsi
	movl	$2, %ecx
	movslq	%edx, %rdx
	salq	$2, %rdx
.L123:
	movq	8(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	40(%rsp), %rax
	movl	$2, %edi
	leaq	.LC26(%rip), %rsi
	movslq	-4(%r13,%rax), %rax
	movl	(%rbx,%rax,4), %ecx
	movslq	0(%r13), %rax
	movl	(%rbx,%rax,4), %edx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	20(%rsp), %edx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC27(%rip), %rsi
	call	__printf_chk@PLT
	movl	(%r14), %edx
	movl	(%rbx), %esi
	movq	8(%rsp), %rdi
	call	_Z9print_matPiii
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	(%rbx), %r15d
	imull	(%r14), %r15d
	movslq	%r15d, %r15
	salq	$2, %r15
	movq	%r15, %rdi
	call	malloc@PLT
	movl	(%rbx), %edx
	imull	(%r14), %edx
	movq	%r15, %rcx
	movq	8(%rsp), %rsi
	movq	%rax, %rdi
	movq	%rax, %r13
	xorl	%ebx, %ebx
	movslq	%edx, %rdx
	salq	$2, %rdx
	call	__memcpy_chk@PLT
	movq	8(%rsp), %rdi
	call	free@PLT
	movq	64(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rdi
	call	cudaFree@PLT
	movq	80(%rsp), %rdi
	call	cudaFree@PLT
.L117:
	cmpl	%ebx, 56(%rsp)
	jle	.L129
	movq	32(%rsp), %rax
	movq	(%rax,%rbx,8), %rdi
	incq	%rbx
	call	cudaFree@PLT
	jmp	.L117
.L129:
	movq	32(%rsp), %rdi
	xorl	%ebx, %ebx
	call	free@PLT
.L119:
	cmpl	%ebx, 20(%rsp)
	jle	.L130
	movq	0(%rbp,%rbx,8), %rdi
	call	cudaFree@PLT
	movq	(%r12,%rbx,8), %rdi
	incq	%rbx
	call	cudaFree@PLT
	jmp	.L119
.L130:
	movq	%rbp, %rdi
	call	free@PLT
	movq	%r12, %rdi
	call	free@PLT
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L121
	call	__stack_chk_fail@PLT
.L121:
	addq	$168, %rsp
	.cfi_def_cfa_offset 56
	movq	%r13, %rax
	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
.LFE2038:
	.size	_Z20gpu_multi_block_algoiPiPS_, .-_Z20gpu_multi_block_algoiPiPS_
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2041:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	subq	$280, %rsp
	.cfi_def_cfa_offset 304
	movq	%fs:40, %rax
	movq	%rax, 264(%rsp)
	xorl	%eax, %eax
	cmpl	$3, %edi
	jne	.L132
	movq	%rsi, %rbx
	leaq	8(%rsp), %rbp
	movq	8(%rsi), %rsi
	movl	$256, %edx
	movq	%rbp, %rdi
	call	__strcpy_chk@PLT
	movq	16(%rbx), %rdi
	call	atoi@PLT
	movl	%eax, %ebx
	cmpl	$2, %eax
	ja	.L132
	movq	%rbp, %rdi
	call	_Z9read_filePc
	testl	%ebx, %ebx
	movq	MAT_LIST(%rip), %rdx
	movq	MAT_DIM(%rip), %rsi
	movl	NUM_DIM(%rip), %edi
	jne	.L138
	call	_Z8cpu_algoiPiPS_
	jmp	.L135
.L132:
	call	_Z5usagev
	jmp	.L135
.L138:
	decl	%ebx
	jne	.L136
	call	_Z18gpu_one_block_algoiPiPS_
	jmp	.L135
.L136:
	call	_Z20gpu_multi_block_algoiPiPS_
.L135:
	movq	264(%rsp), %rax
	subq	%fs:40, %rax
	je	.L137
	call	__stack_chk_fail@PLT
.L137:
	addq	$280, %rsp
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2041:
	.size	main, .-main
	.section	.rodata.str1.1
.LC28:
	.string	"_Z21gpu_par_multi_matmultiiPiPS_S_S_"
.LC29:
	.string	"_Z21gpu_seq_multi_matmultiPiPS_S_S_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2071:
	.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	.LC28(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z21gpu_par_multi_matmultiiPiPS_S_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC29(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z21gpu_seq_multi_matmultiPiPS_S_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2071:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	MAT_LIST
	.bss
	.align 8
	.type	MAT_LIST, @object
	.size	MAT_LIST, 8
MAT_LIST:
	.zero	8
	.globl	MAT_DIM
	.align 8
	.type	MAT_DIM, @object
	.size	MAT_DIM, 8
MAT_DIM:
	.zero	8
	.globl	NUM_DIM
	.align 4
	.type	NUM_DIM, @object
	.size	NUM_DIM, 4
NUM_DIM:
	.zero	4
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC0:
	.long	805306368
	.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 <string.h>
#include <time.h>
#include <math.h>

#define MAX_DIM 20
#define SCALING_FACTOR 256
#define NUM_THREADS 1024
#define MOD_BASE 10007

int NUM_DIM;
int * MAT_DIM;
int ** MAT_LIST;

int find_max(int * arr, int num_elem)
{
    int max = 0;
    for (int i = 0; i < num_elem; i++) {
        if (arr[i] > max) {
            max = arr[i];
        }
    }
    return max;
}

int * def_mat_dim(int k)
{
    int * dim = (int *) malloc(k * sizeof(int));
    int i;
    srand(time(NULL));

    for (i = 0; i < k; i++)
    {
        dim[i] = (rand() % MAX_DIM) + 1;
        //printf("%d\n", dim[i]);
    }
    return dim;
}

int * equipartition(int k, int blocks) 
{
    float div = float(k) / float(blocks);
    int div_int = int(div);
    float rem = div - float(div_int);

    int * partition = (int *) malloc((blocks + 1) * sizeof(int));
    srand(time(NULL));

    partition[0] = 0;
    partition[blocks] = k-1;

    int cur_index = 0;
    float round_factor = 0.0;
    for (int i = 1; i < blocks; i++) {
        cur_index += div;
        round_factor = float(rand()) / float(RAND_MAX);
        if (round_factor < rem) {
            cur_index += 1;
        }
        partition[i] = cur_index;
    }
    return partition;
}

int * creat_mat(int dimX, int dimY)
{
    int x;
    int * mat = (int *) malloc(dimX * dimY * sizeof(int));

    srand(time(NULL));

    for (x = 0; x < dimX * dimY; x++) {
        mat[x] = rand() % MOD_BASE;
        //mat[x] = (rand() % MAX_DIM) * SCALING_FACTOR;
        //printf("%d ", mat[x]);
    }
    return mat;
}

void if_mats_equal(int * A,  int * B, int rows, int cols)
{
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            if (A[i * cols + j] != B[i * cols + j]) { 
                printf("Matrices are not equal\n"); 
                //printf("%d %d\n", i, j);
                return;
            }
        }
    }
    printf("Matrices are equal\n");
}

void cpu_mat_mul(int* A, int* B, int* C, int ARows, int ACols, int BRows, int BCols)
{
    int sum = 0;
    for (int i = 0; i < ARows; i++) {
        for (int j = 0; j < BCols; j++) {
            for (int k = 0; k < ACols; k++) {
                sum += (A[i * ACols + k] * B[k * BCols + j]) % MOD_BASE;
                //sum += A[i * ACols + k] * B[k * BCols + j];
            }
            //printf("%d %d\n", i, j);
            C[i * BCols + j] = sum % MOD_BASE;
            //C[i * BCols + j] = sum;
            sum = 0;
        }
    }
}

void print_mat(int * mat, int dimX, int dimY)
{
    for (int i = 0; i < dimX; i++) {
        for (int j = 0; j < dimY; j++) {
            printf("%d  ", mat[i * dimY + j]);
        }
        printf("\n");
    }
}

int * cpu_multi_mat_mult(int num_dim, int * dim_list, int ** mat_list, int start = 0) {

    int max_dim = find_max(dim_list, num_dim);
    int * output_mat1 = (int *) calloc(max_dim * max_dim, sizeof(int));
    int * output_mat2 = (int *) calloc(max_dim * max_dim, sizeof(int));

    cpu_mat_mul(mat_list[start], mat_list[start + 1], output_mat1, dim_list[start], dim_list[start + 1], dim_list[start + 1], dim_list[start + 2]);
    int num_rows = dim_list[start];
    int num_cols = dim_list[start + 2];

    //print_mat(output_mat1, num_rows, num_cols);
    int num_mult;
    for (num_mult = 1; num_mult < num_dim - 2; num_mult++) {
        //printf("multiplied %d matrices\n", num_mult + 1);
        if (num_mult % 2 == 1) {
            cpu_mat_mul(output_mat1, mat_list[start + num_mult + 1], output_mat2, num_rows, num_cols, dim_list[start + num_mult + 1] , dim_list[start + num_mult + 2]);
        }
        else {
            cpu_mat_mul(output_mat2, mat_list[start + num_mult + 1], output_mat1, num_rows, num_cols, dim_list[start + num_mult + 1] , dim_list[start + num_mult + 2]);
        }
        num_cols = dim_list[start + num_mult + 2];
    }

    int * output_mat = (int *) malloc(dim_list[0] * dim_list[num_dim - 1] * sizeof(int));
    //printf("%d %d\n", num_rows, num_cols);
    if (num_mult % 2 == 1) {
        memcpy(output_mat, output_mat1, dim_list[0] * dim_list[num_dim - 1] * sizeof(int));
    }
    else {
        memcpy(output_mat, output_mat2, dim_list[0] * dim_list[num_dim - 1] * sizeof(int));
    }
    free(output_mat1);
    free(output_mat2);
    return output_mat;
}

__device__
void matmult(int* A, int* B, int* C, int ARows, int ACols, int BRows, int BCols)
{
    int num_elem_output = ARows * BCols;
    int C_elem_row = 0;
    int C_elem_col = 0;
    int sum = 0;

    for (int n = threadIdx.x; n < num_elem_output; n+=NUM_THREADS) {
        C_elem_col = n % BCols;
        C_elem_row = (n + (BCols - C_elem_col)) / BCols - 1;
        
        for (int i = 0; i < ACols; i++) {
            sum += (A[C_elem_row * ACols + i] * B[i * BCols + C_elem_col]) % MOD_BASE;
            //sum += A[C_elem_row * ACols + i] * B[i * BCols + C_elem_col];
        }
        C[C_elem_row * BCols + C_elem_col] = sum % MOD_BASE;
        //C[C_elem_row * BCols + C_elem_col] = sum;
        sum = 0;
        __syncthreads();
    }
    __syncthreads();
}

__global__
void gpu_seq_multi_matmult(int num_dim, int * dim_list, int ** mat_list, int * output_mat1, int * output_mat2)
{

    matmult(mat_list[0], mat_list[1], output_mat1, dim_list[0], dim_list[1], dim_list[1], dim_list[2]);
    __syncthreads();

    int num_mult;
    int num_rows = dim_list[0];
    int num_cols = dim_list[2];
    for (num_mult = 1; num_mult < num_dim - 2; num_mult++) {
        if (num_mult % 2 == 1) {
            matmult(output_mat1, mat_list[num_mult + 1], output_mat2, num_rows, num_cols, dim_list[num_mult + 1], dim_list[num_mult + 2]);
        } else {
            matmult(output_mat2, mat_list[num_mult + 1], output_mat1, num_rows, num_cols, dim_list[num_mult + 1], dim_list[num_mult + 2]);
        }
        num_cols = dim_list[num_mult + 2];
        __syncthreads();
    }
}

__global__
void gpu_par_multi_matmult(int start_dim_idx, int end_dim_idx, int * dim_list, int ** mat_list, int * output_mat1, int * output_mat2)
{

    matmult(mat_list[start_dim_idx], mat_list[start_dim_idx + 1], output_mat1, dim_list[start_dim_idx], dim_list[start_dim_idx + 1], dim_list[start_dim_idx + 1], dim_list[start_dim_idx + 2]);
    __syncthreads();

    //int num_mult;
    int num_rows = dim_list[start_dim_idx];
    int num_cols = dim_list[start_dim_idx + 2];
//    for (num_mult = start_dim_idx + 1; count < end_dim_idx - start_dim_idx - 2; num_mult++) {
    for (int count = 1; count < (end_dim_idx - start_dim_idx - 2); count++) {
        if (count % 2 == 1) {
            matmult(output_mat1, mat_list[start_dim_idx + count + 1], output_mat2, num_rows, num_cols, dim_list[start_dim_idx + count + 1], dim_list[start_dim_idx + count + 2]);
        } else {
            matmult(output_mat2, mat_list[start_dim_idx + count + 1], output_mat1, num_rows, num_cols, dim_list[start_dim_idx + count + 1], dim_list[start_dim_idx + count + 2]);
            //matmult(output_mat2, mat_list[num_mult + 1], output_mat1, num_rows, num_cols, dim_list[num_mult + 1], dim_list[num_mult + 2]);
        }
        num_cols = dim_list[start_dim_idx + count + 2];
        __syncthreads();
    }
}

int * cpu_algo(int num_dim, int * mat_dim, int ** mat_list)
{

    int * cpu_mat = cpu_multi_mat_mult(num_dim, mat_dim, mat_list);
    //printf("%d %d\n", mat_dim[0], mat_dim[num_dim-1]);
    printf("row: %d  col: %d\n", mat_dim[0], mat_dim[num_dim-1]);
    printf("printing cpu sequential result\n");
    print_mat(cpu_mat, mat_dim[0], mat_dim[num_dim-1]);
    printf("\n");
    return cpu_mat;

}

int * gpu_one_block_algo(int num_dim, int * mat_dim, int ** mat_list)
{

    int num_mat = num_dim - 1;
    int max_dim = find_max(mat_dim, num_dim);

    printf("Copying matrix dimensions to device\n");

    int * d_mat_dim;
    hipMalloc((void **)&d_mat_dim, num_dim * sizeof(int));
    hipMemcpy(d_mat_dim, mat_dim, num_dim * sizeof(int), hipMemcpyHostToDevice);

    printf("Allocating space to store output matrix\n");
    int * out_mat = (int *) calloc(max_dim * max_dim, sizeof(int));
    int * d_out_mat1, * d_out_mat2;
    hipMalloc((void **) &d_out_mat1, max_dim * max_dim * sizeof(int));
    hipMalloc((void **) &d_out_mat2, max_dim * max_dim * sizeof(int));
    hipMemcpy(d_out_mat1, out_mat, max_dim * max_dim * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_out_mat2, out_mat, max_dim * max_dim * sizeof(int), hipMemcpyHostToDevice);

    printf("Allocating space for each matrix, and storing pointer address of matrices on the host\n");
    int ** int_mat_list = (int **) malloc(num_mat * sizeof(int *));
    for (int k = 0; k < num_mat; k++) {
        hipMalloc((void **)&int_mat_list[k], mat_dim[k] * mat_dim[k+1] * sizeof(int));
        hipMemcpy(int_mat_list[k], mat_list[k], mat_dim[k] * mat_dim[k+1] * sizeof(int), hipMemcpyHostToDevice);
    }

    printf("Copying pointer addresses of matrices from host to device\n");
    int ** d_mat_list;
    hipMalloc(&d_mat_list, num_mat * sizeof(int *));
    hipMemcpy(d_mat_list, int_mat_list, num_mat * sizeof(int *), hipMemcpyHostToDevice);



    gpu_seq_multi_matmult<<<1, NUM_THREADS>>>(num_dim, d_mat_dim, d_mat_list, d_out_mat1, d_out_mat2);
    hipDeviceSynchronize();


    if (num_dim % 2 == 1) {
        hipMemcpy(out_mat, d_out_mat1, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int), hipMemcpyDeviceToHost);
    } else {
        hipMemcpy(out_mat, d_out_mat2, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int), hipMemcpyDeviceToHost);
    }

    printf("row: %d  col: %d\n", mat_dim[0], mat_dim[num_dim-1]);
    printf("printing gpu (one thread block) result\n");
    print_mat(out_mat, mat_dim[0], mat_dim[num_dim-1]);
    printf("\n");

    int * output_mat = (int *) malloc(mat_dim[0] * mat_dim[num_dim-1] * sizeof(int)); 
    memcpy(output_mat, out_mat, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int));

    free(out_mat);
    hipFree(d_mat_dim);
    hipFree(int_mat_list);
    hipFree(d_mat_list);
    hipFree(d_out_mat1);
    hipFree(d_out_mat2);

    for (int k = 0; k < num_mat; k++) {
        hipFree(int_mat_list[k]);
    }
    free(int_mat_list);

    return output_mat;
}

int * gpu_multi_block_algo(int num_dim, int * mat_dim, int ** mat_list)
{

    int num_mat = num_dim - 1;
    int nblocks = int(double(num_mat) / double(sqrt(num_mat)));
    int max_dim = find_max(mat_dim, num_dim);
    printf("using %d blocks for %d matrices\n", nblocks, num_mat);

    printf("Copying matrix dimensions to device\n");

    int * d_mat_dim;
    hipMalloc((void **)&d_mat_dim, num_dim * sizeof(int));
    hipMemcpy(d_mat_dim, mat_dim, num_dim * sizeof(int), hipMemcpyHostToDevice);

    printf("Allocating space to store output matrix\n");
    int * out_mat = (int *) calloc(max_dim * max_dim, sizeof(int));
    int * d_out_mat1, * d_out_mat2;
    hipMalloc((void **) &d_out_mat1, max_dim * max_dim * sizeof(int));
    hipMalloc((void **) &d_out_mat2, max_dim * max_dim * sizeof(int));
    hipMemcpy(d_out_mat1, out_mat, max_dim * max_dim * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_out_mat2, out_mat, max_dim * max_dim * sizeof(int), hipMemcpyHostToDevice);

    printf("Allocating space for each matrix, and storing pointer address of matrices on the host\n");
    int ** int_mat_list = (int **) malloc(num_mat * sizeof(int *));
    for (int k = 0; k < num_mat; k++) {
        hipMalloc((void **)&int_mat_list[k], mat_dim[k] * mat_dim[k+1] * sizeof(int));
        hipMemcpy(int_mat_list[k], mat_list[k], mat_dim[k] * mat_dim[k+1] * sizeof(int), hipMemcpyHostToDevice);
    }

    printf("Copying pointer addresses of matrices from host to device\n");
    int ** d_mat_list;
    hipMalloc(&d_mat_list, num_mat * sizeof(int *));
    hipMemcpy(d_mat_list, int_mat_list, num_mat * sizeof(int *), hipMemcpyHostToDevice);



    printf("Allocating a set of intermediate arrays to store partial results\n");
    int ** int_mat1, ** int_mat2;
    int_mat1 = (int **) malloc(nblocks * sizeof(int *));
    int_mat2 = (int **) malloc(nblocks * sizeof(int *));
    for (int k = 0; k < nblocks; k++) {
        hipMalloc((void **)&int_mat1[k], max_dim * max_dim * sizeof(int));
        hipMalloc((void **)&int_mat2[k], max_dim * max_dim * sizeof(int));
        hipMemcpy(int_mat1[k], out_mat, max_dim * max_dim * sizeof(int), hipMemcpyHostToDevice);
        hipMemcpy(int_mat2[k], out_mat, max_dim * max_dim * sizeof(int), hipMemcpyHostToDevice);
    }

    printf("Copying pointer addresses of intermediate products from host to device\n");
    int ** int_mat_final = (int **) malloc(nblocks * sizeof(int *));
    int ** d_int_mat_final;
    hipMalloc((void **)&d_int_mat_final, nblocks * sizeof(int *));


    printf("Defining a partition to split work evenly\n");
    int * mat_list_partition = equipartition(num_dim, nblocks);
    //print_mat(mat_list_partition, 1, nblocks + 1);

    printf("Getting the dimensions of the partial products matrices\n");
    //print_mat(mat_dim, 1, num_dim);
    int * int_mat_dim = (int *) malloc((nblocks + 1) * sizeof(int));
    for (int i = 0; i < nblocks + 1; i++) {
        int_mat_dim[i] = mat_dim[mat_list_partition[i]];
    }
    //printf("Moving the dimensions of partial products from host to device\n");
    //print_mat(int_mat_dim, 1, nblocks + 1);
    //printf("\n");

    printf("Moving the dimensions of partial products from host to device\n");
    int * d_int_mat_dim;
    hipMalloc((void **)&d_int_mat_dim, (nblocks + 1) * sizeof(int));
    hipMemcpy(d_int_mat_dim, int_mat_dim, (nblocks + 1) * sizeof(int), hipMemcpyHostToDevice);

    printf("Defining two output arrays which store so that one array storing partial results can be used as input of another output, which is store in the new output array\n");
    int * d_int_output_mat1, * d_int_output_mat2;
    hipMalloc((void **)&d_int_output_mat1, max_dim * max_dim * sizeof(int));
    hipMalloc((void **)&d_int_output_mat2, max_dim * max_dim * sizeof(int));
    hipMemcpy(d_int_output_mat1, out_mat, max_dim * max_dim * sizeof(int), hipMemcpyHostToDevice);
    hipMemcpy(d_int_output_mat2, out_mat, max_dim * max_dim * sizeof(int), hipMemcpyHostToDevice);


    printf("Calling kernels to calculate a subset of input matrices asynchronously\n");
    for (int i = 0; i < nblocks; i++) {

        //printf("nblocks %d, i %d\n", nblocks, i);
        //printf("mat_list_partition %d %d\n", mat_list_partition[i], mat_list_partition[i+1] - 1);
        //printf("end mat_dim %d %d\n", mat_dim[mat_list_partition[i]], mat_dim[mat_list_partition[i+1]]);
        gpu_par_multi_matmult<<<1, NUM_THREADS>>>(mat_list_partition[i], mat_list_partition[i+1] + 1, d_mat_dim, d_mat_list, int_mat1[i], int_mat2[i]);
    }
    hipDeviceSynchronize();

    printf("output arrays accounted for, will be used to aggregate results\n");
    for (int i = 0; i < nblocks; i++) {
        hipMemcpy(out_mat, int_mat1[i], mat_dim[mat_list_partition[i]] * mat_dim[mat_list_partition[i+1]] * sizeof(int), hipMemcpyDeviceToHost);
        if ((mat_list_partition[i+1] - mat_list_partition[i]) % 2 == 1) {
            int_mat_final[i] = int_mat1[i];
        } else {
            int_mat_final[i] = int_mat2[i];
        }
    }

    hipMemcpy(d_int_mat_final, int_mat_final, nblocks * sizeof(int *), hipMemcpyHostToDevice);
    
    printf("Calling last kernel to aggregate partial results\n");
    gpu_par_multi_matmult<<<1, NUM_THREADS>>>(0, nblocks + 1, d_int_mat_dim, d_int_mat_final, d_int_output_mat1, d_int_output_mat2); //, d_iii);
    hipDeviceSynchronize();

    if (nblocks % 2 == 1) {
        hipMemcpy(out_mat, d_int_output_mat1, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int), hipMemcpyDeviceToHost);   
    } else {
        hipMemcpy(out_mat, d_int_output_mat2, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int), hipMemcpyDeviceToHost);   
    }

    printf("rows:%d  cols: %d\n", mat_dim[mat_list_partition[0]], mat_dim[mat_list_partition[nblocks]]);
    printf("printing gpu (using %d blocks) results\n", nblocks);
    print_mat(out_mat, mat_dim[0], mat_dim[num_dim-1]);
    printf("\n");


    int * output_mat = (int *) malloc(mat_dim[0] * mat_dim[num_dim-1] * sizeof(int)); 
    memcpy(output_mat, out_mat, mat_dim[0] * mat_dim[num_dim-1] * sizeof(int));
    

    free(out_mat);
    hipFree(d_mat_dim);
    hipFree(int_mat_list);
    hipFree(d_mat_list);
    hipFree(d_out_mat1);
    hipFree(d_out_mat2);

    for (int k = 0; k < num_mat; k++) {
        hipFree(int_mat_list[k]);
    }
    free(int_mat_list);

    for (int k = 0; k < nblocks; k++) {
        hipFree(int_mat1[k]);
        hipFree(int_mat2[k]);
    }
    free(int_mat1);
    free(int_mat2);

    // STILL MORE THINGS TO FREE


    return output_mat;

}

void read_file(char * filename)//, int * num_dim, int * mat_dim, int ** mat_list)
{

    FILE * fp;
    fp = fopen(filename, "r");

    fscanf(fp, "%d", &NUM_DIM);

    MAT_DIM = (int *) malloc(NUM_DIM * sizeof(int));
    for (int i = 0; i < NUM_DIM; i++) {
        fscanf(fp, "%d", &MAT_DIM[i]);
    }

    MAT_LIST = (int **) malloc((NUM_DIM - 1) * sizeof(int *));

    for (int arr_i = 0; arr_i < (NUM_DIM - 1); arr_i++) {
        MAT_LIST[arr_i] = (int *) malloc(MAT_DIM[arr_i] * MAT_DIM[arr_i + 1] * sizeof(int));
        for (int i = 0; i < (MAT_DIM[arr_i] * MAT_DIM[arr_i + 1]); i++) {
            fscanf(fp, "%d", &MAT_LIST[arr_i][i]);
        }
    }
}


void usage() {

    printf("usage\n ./multi_matmult <filename> <algorithm>\n");
    printf("algorithm:\n\t 0 - cpu (sequential), 1 - gpu (parallel, one thread block), 2 - (parallel, multiple thread blocks)\n");
    //printf("num_blocks : specify only when algorithm == 2. num_blocks should be no more than 1/10th the number of matrices you use\n");

}

int main(int argc, char ** argv)
{
    
    int algorithm = -1;
    int * product_mat;
    char filename[256];

    if (argc == 3) {
        
        strcpy(filename, argv[1]);

        if ((atoi(argv[2]) >= 0) && (atoi(argv[2]) < 3)) {
            algorithm = atoi(argv[2]);
        } else { 
            usage(); 
            return 0;
            }
    } else { 
        usage(); 
        return 0;
        }

    read_file(filename);
/*
    printf("%d\n", NUM_DIM);

    for (int i = 0; i < NUM_DIM; i++) {
        printf("%d ", MAT_DIM[i]);
    }
    printf("\n");

    for (int arr_i = 0; arr_i < (NUM_DIM - 1); arr_i++) {
        for (int i = 0; i < MAT_DIM[arr_i]; i++) {
            for (int j = 0; j < MAT_DIM[arr_i + 1]; j++) {
                printf("%d ", MAT_LIST[arr_i][(i * MAT_DIM[arr_i + 1]) + j]);
            }
            printf("\n");
        }
    }
*/

/*
    for (int i = 0; i < num_dim; i++) {
        printf("%d ", mat_dim[i]);
    }
    printf("\n");
*/

/*
    int num_dim = 45;
    int num_mat = num_dim - 1;
    int * mat_dim = def_mat_dim(num_dim);
    int ** mat_list = (int **) malloc((num_mat) * sizeof(int *));
    int max_dim = find_max(mat_dim, num_dim);

    for (int k = 0; k < num_mat; k++) {
        //printf("================= MATRIX %d ====================\n", k);
        //printf("%d %d\n", mat_dim[k], mat_dim[k+1]);
        mat_list[k] = creat_mat(mat_dim[k], mat_dim[k+1]);
        //printf("%d %d\n", mat_dim[k], mat_dim[k+1]);
        //print_mat(mat_list[k], mat_dim[k], mat_dim[k+1]);
    }
*/    


    if (algorithm == 0) {
        product_mat = cpu_algo(NUM_DIM, MAT_DIM, MAT_LIST);
    } 
    else if (algorithm == 1) {
        product_mat = gpu_one_block_algo(NUM_DIM, MAT_DIM, MAT_LIST);
    } else {
        product_mat = gpu_multi_block_algo(NUM_DIM, MAT_DIM, MAT_LIST);
    }

    return 0;

}

	.text
	.file	"multi_matmult_int.hip"
	.globl	_Z8find_maxPii                  # -- Begin function _Z8find_maxPii
	.type	_Z8find_maxPii,@function
_Z8find_maxPii:                         # @_Z8find_maxPii
	.cfi_startproc
# %bb.0:
	testl	%esi, %esi
	jle	.LBB0_1
# %bb.3:                                # %.lr.ph.preheader
	movl	%esi, %ecx
	xorl	%edx, %edx
	xorl	%eax, %eax
.LBB0_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rdi,%rdx,4), %esi
	cmpl	%eax, %esi
	cmovgl	%esi, %eax
	incq	%rdx
	cmpq	%rdx, %rcx
	jne	.LBB0_4
# %bb.2:                                # %._crit_edge
	retq
.LBB0_1:
	xorl	%eax, %eax
	retq
.Lfunc_end0:
	.size	_Z8find_maxPii, .Lfunc_end0-_Z8find_maxPii
	.cfi_endproc
                                        # -- End function
	.globl	_Z11def_mat_dimi                # -- Begin function _Z11def_mat_dimi
	.type	_Z11def_mat_dimi,@function
_Z11def_mat_dimi:                       # @_Z11def_mat_dimi
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edi, %ebp
	movslq	%edi, %r14
	leaq	(,%r14,4), %rdi
	callq	malloc
	movq	%rax, %rbx
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	testl	%r14d, %r14d
	jle	.LBB1_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ebp, %r14d
	xorl	%r15d, %r15d
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	cltq
	imulq	$1717986919, %rax, %rcx         # imm = 0x66666667
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$35, %rcx
	addl	%edx, %ecx
	shll	$2, %ecx
	leal	(%rcx,%rcx,4), %ecx
	negl	%ecx
	addl	%ecx, %eax
	incl	%eax
	movl	%eax, (%rbx,%r15,4)
	incq	%r15
	cmpq	%r15, %r14
	jne	.LBB1_2
.LBB1_3:                                # %._crit_edge
	movq	%rbx, %rax
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z11def_mat_dimi, .Lfunc_end1-_Z11def_mat_dimi
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z13equipartitionii
.LCPI2_0:
	.long	0x30000000                      # float 4.65661287E-10
	.text
	.globl	_Z13equipartitionii
	.type	_Z13equipartitionii,@function
_Z13equipartitionii:                    # @_Z13equipartitionii
	.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	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$32, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%esi, %ebp
	movl	%edi, %r14d
	movslq	%esi, %r12
	leaq	4(,%r12,4), %rdi
	callq	malloc
	movq	%rax, %rbx
	xorl	%r15d, %r15d
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	movl	$0, (%rbx)
	leal	-1(%r14), %eax
	movl	%eax, (%rbx,%r12,4)
	cmpl	$2, %ebp
	jl	.LBB2_3
# %bb.1:                                # %.lr.ph.preheader
	cvtsi2ss	%r14d, %xmm1
	cvtsi2ss	%ebp, %xmm0
	divss	%xmm0, %xmm1
	cvttps2dq	%xmm1, %xmm0
	cvtdq2ps	%xmm0, %xmm0
	movaps	%xmm1, 16(%rsp)                 # 16-byte Spill
	subss	%xmm0, %xmm1
	movaps	%xmm1, (%rsp)                   # 16-byte Spill
	movl	%ebp, %r14d
	movl	$1, %r12d
.LBB2_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2ss	%r15d, %xmm0
	addss	16(%rsp), %xmm0                 # 16-byte Folded Reload
	cvttss2si	%xmm0, %ebp
	callq	rand
	xorps	%xmm0, %xmm0
	cvtsi2ss	%eax, %xmm0
	mulss	.LCPI2_0(%rip), %xmm0
	xorl	%r15d, %r15d
	movaps	(%rsp), %xmm1                   # 16-byte Reload
	ucomiss	%xmm0, %xmm1
	seta	%r15b
	addl	%ebp, %r15d
	movl	%r15d, (%rbx,%r12,4)
	incq	%r12
	cmpq	%r12, %r14
	jne	.LBB2_2
.LBB2_3:                                # %._crit_edge
	movq	%rbx, %rax
	addq	$32, %rsp
	.cfi_def_cfa_offset 48
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z13equipartitionii, .Lfunc_end2-_Z13equipartitionii
	.cfi_endproc
                                        # -- End function
	.globl	_Z9creat_matii                  # -- Begin function _Z9creat_matii
	.type	_Z9creat_matii,@function
_Z9creat_matii:                         # @_Z9creat_matii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edi, %ebp
	imull	%esi, %ebp
	movslq	%ebp, %r14
	leaq	(,%r14,4), %rdi
	callq	malloc
	movq	%rax, %rbx
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	testl	%r14d, %r14d
	jle	.LBB3_3
# %bb.1:                                # %.lr.ph.preheader
	movl	%ebp, %r14d
	xorl	%r15d, %r15d
.LBB3_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	callq	rand
	cltq
	imulq	$1757988013, %rax, %rcx         # imm = 0x68C8C4AD
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$44, %rcx
	addl	%edx, %ecx
	imull	$10007, %ecx, %ecx              # imm = 0x2717
	subl	%ecx, %eax
	movl	%eax, (%rbx,%r15,4)
	incq	%r15
	cmpq	%r15, %r14
	jne	.LBB3_2
.LBB3_3:                                # %._crit_edge
	movq	%rbx, %rax
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z9creat_matii, .Lfunc_end3-_Z9creat_matii
	.cfi_endproc
                                        # -- End function
	.globl	_Z13if_mats_equalPiS_ii         # -- Begin function _Z13if_mats_equalPiS_ii
	.type	_Z13if_mats_equalPiS_ii,@function
_Z13if_mats_equalPiS_ii:                # @_Z13if_mats_equalPiS_ii
	.cfi_startproc
# %bb.0:
	movq	%rdi, %rax
	movl	$.Lstr.1, %edi
	testl	%edx, %edx
	jle	puts@PLT                        # TAILCALL
# %bb.1:                                # %.preheader.lr.ph
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %rbp, -16
	movl	%edx, %edx
	movl	%ecx, %r8d
	xorl	%r9d, %r9d
	xorl	%r10d, %r10d
.LBB4_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_5 Depth 2
	testl	%ecx, %ecx
	jle	.LBB4_8
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB4_2 Depth=1
	movl	%r9d, %ebx
	leaq	(%rsi,%rbx,4), %r11
	leaq	(%rax,%rbx,4), %rbx
	xorl	%r14d, %r14d
.LBB4_5:                                #   Parent Loop BB4_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%rbx,%r14,4), %ebp
	cmpl	(%r11,%r14,4), %ebp
	jne	.LBB4_6
# %bb.4:                                #   in Loop: Header=BB4_5 Depth=2
	incq	%r14
	cmpq	%r14, %r8
	jne	.LBB4_5
.LBB4_8:                                # %._crit_edge
                                        #   in Loop: Header=BB4_2 Depth=1
	incq	%r10
	addl	%ecx, %r9d
	cmpq	%rdx, %r10
	jne	.LBB4_2
	jmp	.LBB4_7
.LBB4_6:
	movl	$.Lstr, %edi
.LBB4_7:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %rbp
	jmp	puts@PLT                        # TAILCALL
.Lfunc_end4:
	.size	_Z13if_mats_equalPiS_ii, .Lfunc_end4-_Z13if_mats_equalPiS_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z11cpu_mat_mulPiS_S_iiii       # -- Begin function _Z11cpu_mat_mulPiS_S_iiii
	.type	_Z11cpu_mat_mulPiS_S_iiii,@function
_Z11cpu_mat_mulPiS_S_iiii:              # @_Z11cpu_mat_mulPiS_S_iiii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdx, -16(%rsp)                 # 8-byte Spill
	movq	%rsi, -24(%rsp)                 # 8-byte Spill
	movq	%rdi, -32(%rsp)                 # 8-byte Spill
	testl	%ecx, %ecx
	jle	.LBB5_7
# %bb.1:                                # %.preheader27.lr.ph
	movslq	56(%rsp), %rax
	movl	%ecx, %ecx
	movq	%rcx, -8(%rsp)                  # 8-byte Spill
	movl	%eax, %r10d
	movl	%r8d, %r11d
	movq	%rax, -40(%rsp)                 # 8-byte Spill
	leaq	(,%rax,4), %rbx
	xorl	%eax, %eax
	movq	%rax, -48(%rsp)                 # 8-byte Spill
	xorl	%r15d, %r15d
.LBB5_2:                                # %.preheader27
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_4 Depth 2
                                        #       Child Loop BB5_9 Depth 3
	cmpl	$0, 56(%rsp)
	jle	.LBB5_6
# %bb.3:                                # %.preheader.lr.ph
                                        #   in Loop: Header=BB5_2 Depth=1
	movl	-48(%rsp), %eax                 # 4-byte Reload
	movq	-32(%rsp), %rcx                 # 8-byte Reload
	leaq	(%rcx,%rax,4), %r12
	movq	%r15, %rax
	imulq	-40(%rsp), %rax                 # 8-byte Folded Reload
	movq	-16(%rsp), %rcx                 # 8-byte Reload
	leaq	(%rcx,%rax,4), %r13
	movq	-24(%rsp), %rsi                 # 8-byte Reload
	xorl	%eax, %eax
.LBB5_4:                                # %.preheader
                                        #   Parent Loop BB5_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB5_9 Depth 3
	testl	%r8d, %r8d
	jle	.LBB5_5
# %bb.8:                                # %.lr.ph.preheader
                                        #   in Loop: Header=BB5_4 Depth=2
	xorl	%edx, %edx
	movq	%rsi, %rbp
	xorl	%r9d, %r9d
.LBB5_9:                                # %.lr.ph
                                        #   Parent Loop BB5_2 Depth=1
                                        #     Parent Loop BB5_4 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movl	(%rbp), %edi
	imull	(%r12,%rdx,4), %edi
	movslq	%edi, %rdi
	imulq	$1757988013, %rdi, %rcx         # imm = 0x68C8C4AD
	movq	%rcx, %r14
	shrq	$63, %r14
	sarq	$44, %rcx
	addl	%r14d, %ecx
	imull	$10007, %ecx, %ecx              # imm = 0x2717
	subl	%ecx, %edi
	addl	%edi, %r9d
	incq	%rdx
	addq	%rbx, %rbp
	cmpq	%rdx, %r11
	jne	.LBB5_9
	jmp	.LBB5_10
.LBB5_5:                                #   in Loop: Header=BB5_4 Depth=2
	xorl	%r9d, %r9d
.LBB5_10:                               # %._crit_edge
                                        #   in Loop: Header=BB5_4 Depth=2
	movslq	%r9d, %rdx
	imulq	$1757988013, %rdx, %r9          # imm = 0x68C8C4AD
	movq	%r9, %rbp
	shrq	$63, %rbp
	sarq	$44, %r9
	addl	%ebp, %r9d
	imull	$10007, %r9d, %r9d              # imm = 0x2717
	subl	%r9d, %edx
	movl	%edx, (%r13,%rax,4)
	incq	%rax
	addq	$4, %rsi
	cmpq	%r10, %rax
	jne	.LBB5_4
.LBB5_6:                                # %._crit_edge31
                                        #   in Loop: Header=BB5_2 Depth=1
	incq	%r15
	movq	-48(%rsp), %rax                 # 8-byte Reload
	addl	%r8d, %eax
	movq	%rax, -48(%rsp)                 # 8-byte Spill
	cmpq	-8(%rsp), %r15                  # 8-byte Folded Reload
	jne	.LBB5_2
.LBB5_7:                                # %._crit_edge33
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	_Z11cpu_mat_mulPiS_S_iiii, .Lfunc_end5-_Z11cpu_mat_mulPiS_S_iiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z9print_matPiii                # -- Begin function _Z9print_matPiii
	.type	_Z9print_matPiii,@function
_Z9print_matPiii:                       # @_Z9print_matPiii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, 8(%rsp)                   # 8-byte Spill
	testl	%esi, %esi
	jle	.LBB6_6
# %bb.1:                                # %.preheader.lr.ph
	movl	%edx, %ebx
	movl	%esi, %eax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	movl	%edx, %r12d
	xorl	%ebp, %ebp
	xorl	%r13d, %r13d
.LBB6_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB6_4 Depth 2
	testl	%ebx, %ebx
	jle	.LBB6_5
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB6_2 Depth=1
	movl	%ebp, %eax
	movq	8(%rsp), %rcx                   # 8-byte Reload
	leaq	(%rcx,%rax,4), %r14
	xorl	%r15d, %r15d
.LBB6_4:                                #   Parent Loop BB6_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%r14,%r15,4), %esi
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
	incq	%r15
	cmpq	%r15, %r12
	jne	.LBB6_4
.LBB6_5:                                # %._crit_edge
                                        #   in Loop: Header=BB6_2 Depth=1
	movl	$10, %edi
	callq	putchar@PLT
	incq	%r13
	addl	%ebx, %ebp
	cmpq	16(%rsp), %r13                  # 8-byte Folded Reload
	jne	.LBB6_2
.LBB6_6:                                # %._crit_edge13
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	_Z9print_matPiii, .Lfunc_end6-_Z9print_matPiii
	.cfi_endproc
                                        # -- End function
	.globl	_Z18cpu_multi_mat_multiPiPS_i   # -- Begin function _Z18cpu_multi_mat_multiPiPS_i
	.type	_Z18cpu_multi_mat_multiPiPS_i,@function
_Z18cpu_multi_mat_multiPiPS_i:          # @_Z18cpu_multi_mat_multiPiPS_i
	.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	$56, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%ecx, %ebp
	movq	%rdx, %r14
	movq	%rsi, %r15
	movl	%edi, %r13d
	testl	%edi, %edi
	jle	.LBB7_1
# %bb.2:                                # %.lr.ph.preheader.i
	movl	%r13d, %eax
	xorl	%ecx, %ecx
	xorl	%r12d, %r12d
.LBB7_3:                                # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	movl	(%r15,%rcx,4), %edx
	cmpl	%r12d, %edx
	cmovgl	%edx, %r12d
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB7_3
	jmp	.LBB7_4
.LBB7_1:
	xorl	%r12d, %r12d
.LBB7_4:                                # %_Z8find_maxPii.exit
	imull	%r12d, %r12d
	movl	$4, %esi
	movq	%r12, %rdi
	callq	calloc
	movl	$4, %esi
	movq	%r12, %rdi
	movq	%rax, %r12
	callq	calloc
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movslq	%ebp, %rbx
	movq	(%r14,%rbx,8), %rdi
	movq	8(%r14,%rbx,8), %rsi
	movl	(%r15,%rbx,4), %ecx
	movl	4(%r15,%rbx,4), %r8d
	movl	8(%r15,%rbx,4), %ebp
	movl	%ebp, (%rsp)
	movq	%r12, %rdx
	movl	%ecx, 20(%rsp)                  # 4-byte Spill
	callq	_Z11cpu_mat_mulPiS_S_iiii
	movq	%r14, %rax
	movq	%r12, %r14
	cmpl	$4, %r13d
	jl	.LBB7_11
# %bb.5:                                # %.lr.ph.preheader
	movq	%r12, 24(%rsp)                  # 8-byte Spill
	movq	%r13, 40(%rsp)                  # 8-byte Spill
	leal	-2(%r13), %r14d
	leaq	(%rax,%rbx,8), %r12
	addq	$16, %r12
	movq	%r15, 48(%rsp)                  # 8-byte Spill
	leaq	(%r15,%rbx,4), %r15
	addq	$12, %r15
	movl	$1, %ebx
.LBB7_6:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	-8(%r12,%rbx,8), %rsi
	movl	-4(%r15,%rbx,4), %r13d
	testb	$1, %bl
	movq	32(%rsp), %rax                  # 8-byte Reload
	jne	.LBB7_7
# %bb.8:                                #   in Loop: Header=BB7_6 Depth=1
	movl	%r13d, (%rsp)
	movq	%rax, %rdi
	movq	24(%rsp), %rdx                  # 8-byte Reload
	jmp	.LBB7_9
.LBB7_7:                                #   in Loop: Header=BB7_6 Depth=1
	movl	%r13d, (%rsp)
	movq	24(%rsp), %rdi                  # 8-byte Reload
	movq	%rax, %rdx
.LBB7_9:                                #   in Loop: Header=BB7_6 Depth=1
	movl	20(%rsp), %ecx                  # 4-byte Reload
	movl	%ebp, %r8d
	callq	_Z11cpu_mat_mulPiS_S_iiii
	incq	%rbx
	movl	%r13d, %ebp
	cmpq	%rbx, %r14
	jne	.LBB7_6
# %bb.10:                               # %._crit_edge.loopexit
	movq	40(%rsp), %r13                  # 8-byte Reload
	testb	$1, %r13b
	movq	24(%rsp), %r12                  # 8-byte Reload
	movq	%r12, %r14
	cmoveq	32(%rsp), %r14                  # 8-byte Folded Reload
	movq	48(%rsp), %r15                  # 8-byte Reload
.LBB7_11:                               # %._crit_edge
	movslq	(%r15), %rax
	movslq	%r13d, %rcx
	movslq	-4(%r15,%rcx,4), %rbx
	imulq	%rax, %rbx
	shlq	$2, %rbx
	movq	%rbx, %rdi
	callq	malloc
	movq	%rax, %r13
	movq	%rax, %rdi
	movq	%r14, %rsi
	movq	%rbx, %rdx
	callq	memcpy@PLT
	movq	%r12, %rdi
	callq	free
	movq	32(%rsp), %rdi                  # 8-byte Reload
	callq	free
	movq	%r13, %rax
	addq	$56, %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_end7:
	.size	_Z18cpu_multi_mat_multiPiPS_i, .Lfunc_end7-_Z18cpu_multi_mat_multiPiPS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z36__device_stub__gpu_seq_multi_matmultiPiPS_S_S_ # -- Begin function _Z36__device_stub__gpu_seq_multi_matmultiPiPS_S_S_
	.type	_Z36__device_stub__gpu_seq_multi_matmultiPiPS_S_S_,@function
_Z36__device_stub__gpu_seq_multi_matmultiPiPS_S_S_: # @_Z36__device_stub__gpu_seq_multi_matmultiPiPS_S_S_
	.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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	56(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	48(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	32(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z21gpu_seq_multi_matmultiPiPS_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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_end8:
	.size	_Z36__device_stub__gpu_seq_multi_matmultiPiPS_S_S_, .Lfunc_end8-_Z36__device_stub__gpu_seq_multi_matmultiPiPS_S_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z36__device_stub__gpu_par_multi_matmultiiPiPS_S_S_ # -- Begin function _Z36__device_stub__gpu_par_multi_matmultiiPiPS_S_S_
	.type	_Z36__device_stub__gpu_par_multi_matmultiiPiPS_S_S_,@function
_Z36__device_stub__gpu_par_multi_matmultiiPiPS_S_S_: # @_Z36__device_stub__gpu_par_multi_matmultiiPiPS_S_S_
	.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	$144, %rsp
	.cfi_def_cfa_offset 192
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	12(%rsp), %rax
	movl	%edi, (%rax)
	leaq	8(%rsp), %rdi
	movl	%esi, (%rdi)
	leaq	56(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	48(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	40(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	32(%rsp), %r8
	movq	%r9, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	24(%rsp), %r12
	leaq	16(%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	$_Z21gpu_par_multi_matmultiiPiPS_S_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$160, %rsp
	.cfi_adjust_cfa_offset -160
	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_end9:
	.size	_Z36__device_stub__gpu_par_multi_matmultiiPiPS_S_S_, .Lfunc_end9-_Z36__device_stub__gpu_par_multi_matmultiiPiPS_S_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z8cpu_algoiPiPS_               # -- Begin function _Z8cpu_algoiPiPS_
	.type	_Z8cpu_algoiPiPS_,@function
_Z8cpu_algoiPiPS_:                      # @_Z8cpu_algoiPiPS_
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %rbx
	movl	%edi, %ebp
	xorl	%ecx, %ecx
	callq	_Z18cpu_multi_mat_multiPiPS_i
	movq	%rax, %r14
	movl	(%rbx), %esi
	movslq	%ebp, %r15
	movl	-4(%rbx,%r15,4), %edx
	movl	$.L.str.4, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	(%rbx), %esi
	movl	-4(%rbx,%r15,4), %edx
	movq	%r14, %rdi
	callq	_Z9print_matPiii
	movl	$10, %edi
	callq	putchar@PLT
	movq	%r14, %rax
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end10:
	.size	_Z8cpu_algoiPiPS_, .Lfunc_end10-_Z8cpu_algoiPiPS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z18gpu_one_block_algoiPiPS_    # -- Begin function _Z18gpu_one_block_algoiPiPS_
	.type	_Z18gpu_one_block_algoiPiPS_,@function
_Z18gpu_one_block_algoiPiPS_:           # @_Z18gpu_one_block_algoiPiPS_
	.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	$72, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdx, %r13
	movq	%rsi, %r14
	movl	%edi, %ebx
	leal	-1(%rbx), %eax
	movl	%eax, 4(%rsp)                   # 4-byte Spill
	testl	%edi, %edi
	jle	.LBB11_1
# %bb.2:                                # %.lr.ph.preheader.i
	movl	%ebx, %eax
	xorl	%ecx, %ecx
	xorl	%r15d, %r15d
.LBB11_3:                               # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	movl	(%r14,%rcx,4), %edx
	cmpl	%r15d, %edx
	cmovgl	%edx, %r15d
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB11_3
	jmp	.LBB11_4
.LBB11_1:
	xorl	%r15d, %r15d
.LBB11_4:                               # %_Z8find_maxPii.exit
	movl	$.Lstr.8, %edi
	callq	puts@PLT
	movq	%rbx, 24(%rsp)                  # 8-byte Spill
	movslq	%ebx, %rbx
	shlq	$2, %rbx
	leaq	40(%rsp), %r12
	movq	%r12, %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	movq	(%r12), %rdi
	movq	%r14, %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.Lstr.9, %edi
	callq	puts@PLT
	imull	%r15d, %r15d
	movl	$4, %esi
	movq	%r15, %rdi
	callq	calloc
	movq	%rax, %rbx
	shlq	$2, %r15
	leaq	16(%rsp), %r12
	movq	%r12, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	leaq	8(%rsp), %rbp
	movq	%rbp, %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	movq	(%r12), %rdi
	movq	%rbx, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbp), %rdi
	movq	%rbx, 64(%rsp)                  # 8-byte Spill
	movq	%rbx, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.Lstr.10, %edi
	callq	puts@PLT
	movslq	4(%rsp), %rax                   # 4-byte Folded Reload
	movq	%rax, 56(%rsp)                  # 8-byte Spill
	leaq	(,%rax,8), %rdi
	movq	%rdi, 48(%rsp)                  # 8-byte Spill
	callq	malloc
	movq	%rax, %r15
	cmpl	$2, 24(%rsp)                    # 4-byte Folded Reload
	jl	.LBB11_7
# %bb.5:                                # %.lr.ph.preheader
	movl	4(%rsp), %r12d                  # 4-byte Reload
	leaq	4(%r14), %rbp
	shlq	$3, %r12
	xorl	%ebx, %ebx
.LBB11_6:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	leaq	(%r15,%rbx), %rdi
	movslq	-4(%rbp), %rax
	movslq	(%rbp), %rsi
	imulq	%rax, %rsi
	shlq	$2, %rsi
	callq	hipMalloc
	movq	(%r15,%rbx), %rdi
	movq	(%r13,%rbx), %rsi
	movslq	-4(%rbp), %rax
	movslq	(%rbp), %rdx
	imulq	%rax, %rdx
	shlq	$2, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	addq	$8, %rbx
	addq	$4, %rbp
	cmpq	%rbx, %r12
	jne	.LBB11_6
.LBB11_7:                               # %._crit_edge
	movl	$.Lstr.11, %edi
	callq	puts@PLT
	leaq	32(%rsp), %r13
	movq	%r13, %rdi
	movq	48(%rsp), %rbx                  # 8-byte Reload
	movq	%rbx, %rsi
	callq	hipMalloc
	movq	(%r13), %rdi
	movq	%r15, %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	1023(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movq	24(%rsp), %r12                  # 8-byte Reload
	jne	.LBB11_9
# %bb.8:
	movq	40(%rsp), %rsi
	movq	32(%rsp), %rdx
	movq	16(%rsp), %rcx
	movq	8(%rsp), %r8
	movl	%r12d, %edi
	callq	_Z36__device_stub__gpu_seq_multi_matmultiPiPS_S_S_
.LBB11_9:
	callq	hipDeviceSynchronize
	movl	%r12d, %eax
	andl	$-2147483647, %eax              # imm = 0x80000001
	movslq	(%r14), %rcx
	movq	56(%rsp), %rbx                  # 8-byte Reload
	movslq	(%r14,%rbx,4), %rdx
	imulq	%rcx, %rdx
	shlq	$2, %rdx
	cmpl	$1, %eax
	leaq	16(%rsp), %rax
	leaq	8(%rsp), %rcx
	cmoveq	%rax, %rcx
	movq	(%rcx), %rsi
	movq	64(%rsp), %r13                  # 8-byte Reload
	movq	%r13, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movl	(%r14), %esi
	movl	(%r14,%rbx,4), %edx
	movl	$.L.str.4, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr.7, %edi
	callq	puts@PLT
	movl	(%r14), %esi
	movl	(%r14,%rbx,4), %edx
	movq	%r13, %rdi
	callq	_Z9print_matPiii
	movl	$10, %edi
	callq	putchar@PLT
	movslq	(%r14), %rax
	movslq	(%r14,%rbx,4), %r14
	imulq	%rax, %r14
	shlq	$2, %r14
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, %rbx
	movq	%rax, %rdi
	movq	%r13, %rsi
	movq	%r14, %rdx
	callq	memcpy@PLT
	movq	%r13, %rdi
	callq	free
	movq	40(%rsp), %rdi
	callq	hipFree
	movq	%r15, %rdi
	callq	hipFree
	movq	32(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	cmpl	$2, %r12d
	jl	.LBB11_12
# %bb.10:                               # %.lr.ph86.preheader
	movl	4(%rsp), %r14d                  # 4-byte Reload
	xorl	%r12d, %r12d
.LBB11_11:                              # %.lr.ph86
                                        # =>This Inner Loop Header: Depth=1
	movq	(%r15,%r12,8), %rdi
	callq	hipFree
	incq	%r12
	cmpq	%r12, %r14
	jne	.LBB11_11
.LBB11_12:                              # %._crit_edge87
	movq	%r15, %rdi
	callq	free
	movq	%rbx, %rax
	addq	$72, %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_end11:
	.size	_Z18gpu_one_block_algoiPiPS_, .Lfunc_end11-_Z18gpu_one_block_algoiPiPS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z20gpu_multi_block_algoiPiPS_  # -- Begin function _Z20gpu_multi_block_algoiPiPS_
	.type	_Z20gpu_multi_block_algoiPiPS_,@function
_Z20gpu_multi_block_algoiPiPS_:         # @_Z20gpu_multi_block_algoiPiPS_
	.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	$200, %rsp
	.cfi_def_cfa_offset 256
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edi, %r13d
	leal	-1(%r13), %r14d
	cvtsi2sd	%r14d, %xmm1
	movq	%rsi, %rbp
	xorpd	%xmm0, %xmm0
	ucomisd	%xmm0, %xmm1
	movq	%rdx, 16(%rsp)                  # 8-byte Spill
	jb	.LBB12_2
# %bb.1:
	xorps	%xmm0, %xmm0
	sqrtsd	%xmm1, %xmm0
	jmp	.LBB12_3
.LBB12_2:                               # %call.sqrt
	movapd	%xmm1, %xmm0
	movsd	%xmm1, 8(%rsp)                  # 8-byte Spill
	callq	sqrt
	movsd	8(%rsp), %xmm1                  # 8-byte Reload
                                        # xmm1 = mem[0],zero
.LBB12_3:                               # %.split
	divsd	%xmm0, %xmm1
	cvttsd2si	%xmm1, %esi
	testl	%r13d, %r13d
	jle	.LBB12_4
# %bb.5:                                # %.lr.ph.preheader.i
	movl	%r13d, %eax
	xorl	%ecx, %ecx
	xorl	%ebx, %ebx
.LBB12_6:                               # %.lr.ph.i
                                        # =>This Inner Loop Header: Depth=1
	movl	(%rbp,%rcx,4), %edx
	cmpl	%ebx, %edx
	cmovgl	%edx, %ebx
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB12_6
	jmp	.LBB12_7
.LBB12_4:
	xorl	%ebx, %ebx
.LBB12_7:                               # %_Z8find_maxPii.exit
	movl	$.L.str.11, %edi
	movq	%rsi, 40(%rsp)                  # 8-byte Spill
                                        # kill: def $esi killed $esi killed $rsi
	movl	%r14d, %edx
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr.8, %edi
	callq	puts@PLT
	movslq	%r13d, %r12
	shlq	$2, %r12
	movq	%r13, 72(%rsp)                  # 8-byte Spill
	leaq	136(%rsp), %r13
	movq	%r13, %rdi
	movq	%r12, %rsi
	callq	hipMalloc
	movq	(%r13), %rdi
	movq	%rbp, 64(%rsp)                  # 8-byte Spill
	movq	%rbp, %rsi
	movq	%r12, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.Lstr.9, %edi
	callq	puts@PLT
	imull	%ebx, %ebx
	movl	$4, %esi
	movq	%rbx, %rdi
	callq	calloc
	movq	%rax, %r13
	shlq	$2, %rbx
	leaq	192(%rsp), %r12
	movq	%r12, %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	184(%rsp), %rbp
	movq	%rbp, %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	movq	(%r12), %rdi
	movq	%r13, %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbp), %rdi
	movq	%r13, 96(%rsp)                  # 8-byte Spill
	movq	%r13, %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.Lstr.10, %edi
	callq	puts@PLT
	movl	%r14d, 28(%rsp)                 # 4-byte Spill
	movslq	%r14d, %rax
	movq	%rax, 152(%rsp)                 # 8-byte Spill
	leaq	(,%rax,8), %rdi
	movq	%rdi, 8(%rsp)                   # 8-byte Spill
	callq	malloc
	movq	%rax, %r13
	cmpl	$2, 72(%rsp)                    # 4-byte Folded Reload
	movq	16(%rsp), %r15                  # 8-byte Reload
	jl	.LBB12_10
# %bb.8:                                # %.lr.ph.preheader
	movl	28(%rsp), %r12d                 # 4-byte Reload
	movq	64(%rsp), %rax                  # 8-byte Reload
	leaq	4(%rax), %rbp
	shlq	$3, %r12
	xorl	%r14d, %r14d
.LBB12_9:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	leaq	(%r14,%r13), %rdi
	movslq	-4(%rbp), %rax
	movslq	(%rbp), %rsi
	imulq	%rax, %rsi
	shlq	$2, %rsi
	callq	hipMalloc
	movq	(%r13,%r14), %rdi
	movq	(%r15,%r14), %rsi
	movslq	-4(%rbp), %rax
	movslq	(%rbp), %rdx
	imulq	%rax, %rdx
	shlq	$2, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	addq	$8, %r14
	addq	$4, %rbp
	cmpq	%r14, %r12
	jne	.LBB12_9
.LBB12_10:                              # %._crit_edge
	movl	$.Lstr.11, %edi
	callq	puts@PLT
	leaq	128(%rsp), %r14
	movq	%r14, %rdi
	movq	8(%rsp), %r15                   # 8-byte Reload
	movq	%r15, %rsi
	callq	hipMalloc
	movq	(%r14), %rdi
	movq	%r13, 160(%rsp)                 # 8-byte Spill
	movq	%r13, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.Lstr.12, %edi
	callq	puts@PLT
	movq	40(%rsp), %r15                  # 8-byte Reload
	movslq	%r15d, %rax
	movq	%rax, 144(%rsp)                 # 8-byte Spill
	leaq	(,%rax,8), %r14
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movq	%r14, 88(%rsp)                  # 8-byte Spill
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	movl	%r15d, %eax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	testl	%r15d, %r15d
	movq	96(%rsp), %r13                  # 8-byte Reload
	jle	.LBB12_13
# %bb.11:                               # %.lr.ph215.preheader
	movq	16(%rsp), %r12                  # 8-byte Reload
	movq	32(%rsp), %r14                  # 8-byte Reload
	movq	48(%rsp), %rbp                  # 8-byte Reload
.LBB12_12:                              # %.lr.ph215
                                        # =>This Inner Loop Header: Depth=1
	movq	%r14, %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	movq	%rbp, %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	movq	(%r14), %rdi
	movq	%r13, %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbp), %rdi
	movq	%r13, %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	addq	$8, %rbp
	addq	$8, %r14
	decq	%r12
	jne	.LBB12_12
.LBB12_13:                              # %._crit_edge216
	movl	$.Lstr.13, %edi
	callq	puts@PLT
	movq	88(%rsp), %r14                  # 8-byte Reload
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, 80(%rsp)                  # 8-byte Spill
	leaq	120(%rsp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	movl	$.Lstr.14, %edi
	callq	puts@PLT
	movq	72(%rsp), %rdi                  # 8-byte Reload
                                        # kill: def $edi killed $edi killed $rdi
	movl	%r15d, %esi
	callq	_Z13equipartitionii
	movq	%rax, %rbp
	movl	$.Lstr.15, %edi
	callq	puts@PLT
	leal	1(%r15), %eax
	movl	%eax, 24(%rsp)                  # 4-byte Spill
	movslq	%eax, %rdi
	shlq	$2, %rdi
	movq	%rdi, 8(%rsp)                   # 8-byte Spill
	callq	malloc
	movq	%rax, %r14
	testl	%r15d, %r15d
	movq	64(%rsp), %r12                  # 8-byte Reload
	js	.LBB12_16
# %bb.14:                               # %.lr.ph220.preheader
	movl	24(%rsp), %eax                  # 4-byte Reload
	xorl	%ecx, %ecx
.LBB12_15:                              # %.lr.ph220
                                        # =>This Inner Loop Header: Depth=1
	movslq	(%rbp,%rcx,4), %rdx
	movl	(%r12,%rdx,4), %edx
	movl	%edx, (%r14,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB12_15
.LBB12_16:                              # %._crit_edge221
	movq	%rbp, 56(%rsp)                  # 8-byte Spill
	movl	$.Lstr.16, %edi
	callq	puts@PLT
	leaq	176(%rsp), %r15
	movq	%r15, %rdi
	movq	8(%rsp), %rbp                   # 8-byte Reload
	movq	%rbp, %rsi
	callq	hipMalloc
	movq	(%r15), %rdi
	movq	%r14, %rsi
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.Lstr.17, %edi
	callq	puts@PLT
	leaq	112(%rsp), %r14
	movq	%r14, %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	leaq	104(%rsp), %rbp
	movq	%rbp, %rdi
	movq	%rbx, %rsi
	callq	hipMalloc
	movq	(%r14), %rdi
	movq	%r13, %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	(%rbp), %rdi
	movq	%r13, %rsi
	movq	%rbx, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.Lstr.18, %edi
	callq	puts@PLT
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movq	40(%rsp), %r15                  # 8-byte Reload
	testl	%r15d, %r15d
	movq	32(%rsp), %rbp                  # 8-byte Reload
	jle	.LBB12_24
# %bb.17:                               # %.lr.ph224.preheader
	xorl	%ebx, %ebx
	leaq	1023(%rdi), %r14
.LBB12_18:                              # %.lr.ph224
                                        # =>This Inner Loop Header: Depth=1
	movl	$1, %esi
	movq	%r14, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB12_20
# %bb.19:                               #   in Loop: Header=BB12_18 Depth=1
	movq	56(%rsp), %rax                  # 8-byte Reload
	movl	(%rax,%rbx,4), %edi
	movl	4(%rax,%rbx,4), %esi
	incl	%esi
	movq	136(%rsp), %rdx
	movq	128(%rsp), %rcx
	movq	(%rbp,%rbx,8), %r8
	movq	48(%rsp), %rax                  # 8-byte Reload
	movq	(%rax,%rbx,8), %r9
	callq	_Z36__device_stub__gpu_par_multi_matmultiiPiPS_S_S_
.LBB12_20:                              #   in Loop: Header=BB12_18 Depth=1
	incq	%rbx
	cmpq	%rbx, 16(%rsp)                  # 8-byte Folded Reload
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	jne	.LBB12_18
# %bb.21:                               # %._crit_edge225
	callq	hipDeviceSynchronize
	movl	$.Lstr.19, %edi
	callq	puts@PLT
	testl	%r15d, %r15d
	movq	80(%rsp), %rsi                  # 8-byte Reload
	jle	.LBB12_25
# %bb.22:                               # %.lr.ph228.preheader
	movq	%r12, %r15
	movq	%r13, %rbp
	movq	56(%rsp), %rdx                  # 8-byte Reload
	movl	(%rdx), %r14d
	xorl	%r13d, %r13d
	movq	32(%rsp), %rbx                  # 8-byte Reload
	movq	48(%rsp), %r12                  # 8-byte Reload
.LBB12_23:                              # %.lr.ph228
                                        # =>This Inner Loop Header: Depth=1
	movq	(%rbx), %rsi
	movslq	%r14d, %rax
	movq	%rax, 168(%rsp)                 # 8-byte Spill
	movslq	(%r15,%rax,4), %rax
	leaq	1(%r13), %rdx
	movq	%rdx, 8(%rsp)                   # 8-byte Spill
	movq	56(%rsp), %rdx                  # 8-byte Reload
	movslq	4(%rdx,%r13,4), %r14
	movslq	(%r15,%r14,4), %rdx
	imulq	%rax, %rdx
	shlq	$2, %rdx
	movq	%rbp, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	80(%rsp), %rsi                  # 8-byte Reload
	movl	%r14d, %eax
	subl	168(%rsp), %eax                 # 4-byte Folded Reload
	andl	$-2147483647, %eax              # imm = 0x80000001
	cmpl	$1, %eax
	movq	%r12, %rax
	cmoveq	%rbx, %rax
	movq	(%rax), %rax
	movq	%rax, (%rsi,%r13,8)
	movq	16(%rsp), %rax                  # 8-byte Reload
	addq	$8, %r12
	addq	$8, %rbx
	movq	8(%rsp), %rdx                   # 8-byte Reload
	movq	%rdx, %r13
	cmpq	%rdx, %rax
	jne	.LBB12_23
	jmp	.LBB12_25
.LBB12_24:                              # %._crit_edge229.critedge
	callq	hipDeviceSynchronize
	movl	$.Lstr.19, %edi
	callq	puts@PLT
	movq	80(%rsp), %rsi                  # 8-byte Reload
.LBB12_25:                              # %._crit_edge229
	movq	120(%rsp), %rdi
	movq	88(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$.Lstr.20, %edi
	callq	puts@PLT
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	1023(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB12_27
# %bb.26:
	movq	176(%rsp), %rdx
	movq	120(%rsp), %rcx
	movq	112(%rsp), %r8
	movq	104(%rsp), %r9
	xorl	%edi, %edi
	movl	24(%rsp), %esi                  # 4-byte Reload
	callq	_Z36__device_stub__gpu_par_multi_matmultiiPiPS_S_S_
.LBB12_27:
	callq	hipDeviceSynchronize
	movq	40(%rsp), %rbx                  # 8-byte Reload
	movl	%ebx, %eax
	andl	$-2147483647, %eax              # imm = 0x80000001
	movq	64(%rsp), %r14                  # 8-byte Reload
	movslq	(%r14), %rcx
	movq	152(%rsp), %r12                 # 8-byte Reload
	movslq	(%r14,%r12,4), %rdx
	imulq	%rcx, %rdx
	shlq	$2, %rdx
	cmpl	$1, %eax
	leaq	112(%rsp), %rax
	leaq	104(%rsp), %rcx
	cmoveq	%rax, %rcx
	movq	(%rcx), %rsi
	movq	96(%rsp), %r15                  # 8-byte Reload
	movq	%r15, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	56(%rsp), %rcx                  # 8-byte Reload
	movslq	(%rcx), %rax
	movl	(%r14,%rax,4), %esi
	movq	144(%rsp), %rax                 # 8-byte Reload
	movslq	(%rcx,%rax,4), %rax
	movl	(%r14,%rax,4), %edx
	movl	$.L.str.21, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.22, %edi
	movl	%ebx, %esi
	xorl	%eax, %eax
	callq	printf
	movl	(%r14), %esi
	movl	(%r14,%r12,4), %edx
	movq	%r15, %rdi
	callq	_Z9print_matPiii
	movl	$10, %edi
	callq	putchar@PLT
	movslq	(%r14), %rax
	movslq	(%r14,%r12,4), %r14
	imulq	%rax, %r14
	shlq	$2, %r14
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, %rbx
	movq	%rax, %rdi
	movq	%r15, %rsi
	movq	%r14, %rdx
	callq	memcpy@PLT
	movq	%r15, %rdi
	callq	free
	movq	136(%rsp), %rdi
	callq	hipFree
	movq	160(%rsp), %r12                 # 8-byte Reload
	movq	%r12, %rdi
	callq	hipFree
	movq	128(%rsp), %rdi
	callq	hipFree
	movq	192(%rsp), %rdi
	callq	hipFree
	movq	184(%rsp), %rdi
	callq	hipFree
	cmpl	$2, 72(%rsp)                    # 4-byte Folded Reload
	movq	16(%rsp), %r13                  # 8-byte Reload
	jl	.LBB12_30
# %bb.28:                               # %.lr.ph232.preheader
	movl	28(%rsp), %r14d                 # 4-byte Reload
	xorl	%r15d, %r15d
.LBB12_29:                              # %.lr.ph232
                                        # =>This Inner Loop Header: Depth=1
	movq	(%r12,%r15,8), %rdi
	callq	hipFree
	incq	%r15
	cmpq	%r15, %r14
	jne	.LBB12_29
.LBB12_30:                              # %._crit_edge233
	movq	%r12, %rdi
	callq	free
	cmpl	$0, 40(%rsp)                    # 4-byte Folded Reload
	movq	32(%rsp), %r15                  # 8-byte Reload
	movq	48(%rsp), %r12                  # 8-byte Reload
	jle	.LBB12_33
# %bb.31:                               # %.lr.ph236.preheader
	xorl	%r14d, %r14d
.LBB12_32:                              # %.lr.ph236
                                        # =>This Inner Loop Header: Depth=1
	movq	(%r15,%r14,8), %rdi
	callq	hipFree
	movq	(%r12,%r14,8), %rdi
	callq	hipFree
	incq	%r14
	cmpq	%r14, %r13
	jne	.LBB12_32
.LBB12_33:                              # %._crit_edge237
	movq	%r15, %rdi
	callq	free
	movq	%r12, %rdi
	callq	free
	movq	%rbx, %rax
	addq	$200, %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_end12:
	.size	_Z20gpu_multi_block_algoiPiPS_, .Lfunc_end12-_Z20gpu_multi_block_algoiPiPS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z9read_filePc                  # -- Begin function _Z9read_filePc
	.type	_Z9read_filePc,@function
_Z9read_filePc:                         # @_Z9read_filePc
	.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
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$.L.str.23, %esi
	callq	fopen
	movq	%rax, %rbx
	movl	$.L.str.24, %esi
	movl	$NUM_DIM, %edx
	movq	%rax, %rdi
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	movslq	NUM_DIM(%rip), %r14
	leaq	(,%r14,4), %rdi
	callq	malloc
	movq	%rax, MAT_DIM(%rip)
	testq	%r14, %r14
	jle	.LBB13_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%r15d, %r15d
	xorl	%r12d, %r12d
.LBB13_2:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	MAT_DIM(%rip), %rdx
	addq	%r15, %rdx
	movl	$.L.str.24, %esi
	movq	%rbx, %rdi
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	incq	%r12
	movslq	NUM_DIM(%rip), %r14
	addq	$4, %r15
	cmpq	%r14, %r12
	jl	.LBB13_2
.LBB13_3:                               # %._crit_edge
	movslq	%r14d, %r15
	leaq	-8(,%r15,8), %rdi
	callq	malloc
	movq	%rax, MAT_LIST(%rip)
	cmpl	$2, %r15d
	jl	.LBB13_10
# %bb.4:                                # %.lr.ph26.preheader
	movq	MAT_DIM(%rip), %rax
	xorl	%r15d, %r15d
.LBB13_7:                               # %.lr.ph26
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB13_9 Depth 2
	movq	%r15, %r12
	movslq	(%rax,%r15,4), %rcx
	incq	%r15
	movslq	4(%rax,%r12,4), %rdi
	imulq	%rcx, %rdi
	shlq	$2, %rdi
	callq	malloc
	movq	MAT_LIST(%rip), %rcx
	movq	%rax, (%rcx,%r12,8)
	movq	MAT_DIM(%rip), %rax
	movl	4(%rax,%r12,4), %ecx
	imull	(%rax,%r12,4), %ecx
	testl	%ecx, %ecx
	jle	.LBB13_6
# %bb.8:                                # %.lr.ph23.preheader
                                        #   in Loop: Header=BB13_7 Depth=1
	xorl	%r14d, %r14d
	xorl	%r13d, %r13d
.LBB13_9:                               # %.lr.ph23
                                        #   Parent Loop BB13_7 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	MAT_LIST(%rip), %rax
	movq	(%rax,%r12,8), %rdx
	addq	%r14, %rdx
	movl	$.L.str.24, %esi
	movq	%rbx, %rdi
	xorl	%eax, %eax
	callq	__isoc23_fscanf
	incq	%r13
	movq	MAT_DIM(%rip), %rax
	movslq	(%rax,%r12,4), %rcx
	movslq	(%rax,%r15,4), %rdx
	imulq	%rcx, %rdx
	addq	$4, %r14
	cmpq	%rdx, %r13
	jl	.LBB13_9
# %bb.5:                                # %.loopexit.loopexit
                                        #   in Loop: Header=BB13_7 Depth=1
	movl	NUM_DIM(%rip), %r14d
.LBB13_6:                               # %.loopexit
                                        #   in Loop: Header=BB13_7 Depth=1
	leal	-1(%r14), %ecx
	movslq	%ecx, %rcx
	cmpq	%rcx, %r15
	jl	.LBB13_7
.LBB13_10:                              # %._crit_edge27
	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_end13:
	.size	_Z9read_filePc, .Lfunc_end13-_Z9read_filePc
	.cfi_endproc
                                        # -- End function
	.globl	_Z5usagev                       # -- Begin function _Z5usagev
	.type	_Z5usagev,@function
_Z5usagev:                              # @_Z5usagev
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	$.Lstr.21, %edi
	callq	puts@PLT
	movl	$.Lstr.22, %edi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	puts@PLT                        # TAILCALL
.Lfunc_end14:
	.size	_Z5usagev, .Lfunc_end14-_Z5usagev
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$256, %rsp                      # imm = 0x100
	.cfi_def_cfa_offset 272
	.cfi_offset %rbx, -16
	cmpl	$3, %edi
	jne	.LBB15_5
# %bb.1:
	movq	%rsi, %rbx
	movq	8(%rsi), %rsi
	movq	%rsp, %rdi
	callq	strcpy
	movq	16(%rbx), %rdi
	callq	atoi
	cmpl	$2, %eax
	ja	.LBB15_5
# %bb.2:
	movl	%eax, %ebx
	movq	%rsp, %rdi
	callq	_Z9read_filePc
	movl	NUM_DIM(%rip), %edi
	movq	MAT_DIM(%rip), %rsi
	movq	MAT_LIST(%rip), %rdx
	testl	%ebx, %ebx
	je	.LBB15_6
# %bb.3:
	cmpl	$1, %ebx
	jne	.LBB15_7
# %bb.4:
	callq	_Z18gpu_one_block_algoiPiPS_
	jmp	.LBB15_8
.LBB15_5:
	movl	$.Lstr.21, %edi
	callq	puts@PLT
	movl	$.Lstr.22, %edi
	callq	puts@PLT
.LBB15_8:
	xorl	%eax, %eax
	addq	$256, %rsp                      # imm = 0x100
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.LBB15_6:
	.cfi_def_cfa_offset 272
	callq	_Z8cpu_algoiPiPS_
	jmp	.LBB15_8
.LBB15_7:
	callq	_Z20gpu_multi_block_algoiPiPS_
	jmp	.LBB15_8
.Lfunc_end15:
	.size	main, .Lfunc_end15-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB16_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB16_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z21gpu_seq_multi_matmultiPiPS_S_S_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z21gpu_par_multi_matmultiiPiPS_S_S_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_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	NUM_DIM,@object                 # @NUM_DIM
	.bss
	.globl	NUM_DIM
	.p2align	2, 0x0
NUM_DIM:
	.long	0                               # 0x0
	.size	NUM_DIM, 4

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

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

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

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

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

	.type	.L.str.4,@object                # @.str.4
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.4:
	.asciz	"row: %d  col: %d\n"
	.size	.L.str.4, 18

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"using %d blocks for %d matrices\n"
	.size	.L.str.11, 33

	.type	.L.str.21,@object               # @.str.21
.L.str.21:
	.asciz	"rows:%d  cols: %d\n"
	.size	.L.str.21, 19

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	"printing gpu (using %d blocks) results\n"
	.size	.L.str.22, 40

	.type	.L.str.23,@object               # @.str.23
.L.str.23:
	.asciz	"r"
	.size	.L.str.23, 2

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

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z21gpu_seq_multi_matmultiPiPS_S_S_"
	.size	.L__unnamed_1, 36

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z21gpu_par_multi_matmultiiPiPS_S_S_"
	.size	.L__unnamed_2, 37

	.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

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"Matrices are not equal"
	.size	.Lstr, 23

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"Matrices are equal"
	.size	.Lstr.1, 19

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"printing cpu sequential result"
	.size	.Lstr.2, 31

	.type	.Lstr.7,@object                 # @str.7
.Lstr.7:
	.asciz	"printing gpu (one thread block) result"
	.size	.Lstr.7, 39

	.type	.Lstr.8,@object                 # @str.8
.Lstr.8:
	.asciz	"Copying matrix dimensions to device"
	.size	.Lstr.8, 36

	.type	.Lstr.9,@object                 # @str.9
.Lstr.9:
	.asciz	"Allocating space to store output matrix"
	.size	.Lstr.9, 40

	.type	.Lstr.10,@object                # @str.10
.Lstr.10:
	.asciz	"Allocating space for each matrix, and storing pointer address of matrices on the host"
	.size	.Lstr.10, 86

	.type	.Lstr.11,@object                # @str.11
.Lstr.11:
	.asciz	"Copying pointer addresses of matrices from host to device"
	.size	.Lstr.11, 58

	.type	.Lstr.12,@object                # @str.12
.Lstr.12:
	.asciz	"Allocating a set of intermediate arrays to store partial results"
	.size	.Lstr.12, 65

	.type	.Lstr.13,@object                # @str.13
.Lstr.13:
	.asciz	"Copying pointer addresses of intermediate products from host to device"
	.size	.Lstr.13, 71

	.type	.Lstr.14,@object                # @str.14
.Lstr.14:
	.asciz	"Defining a partition to split work evenly"
	.size	.Lstr.14, 42

	.type	.Lstr.15,@object                # @str.15
.Lstr.15:
	.asciz	"Getting the dimensions of the partial products matrices"
	.size	.Lstr.15, 56

	.type	.Lstr.16,@object                # @str.16
.Lstr.16:
	.asciz	"Moving the dimensions of partial products from host to device"
	.size	.Lstr.16, 62

	.type	.Lstr.17,@object                # @str.17
.Lstr.17:
	.asciz	"Defining two output arrays which store so that one array storing partial results can be used as input of another output, which is store in the new output array"
	.size	.Lstr.17, 160

	.type	.Lstr.18,@object                # @str.18
.Lstr.18:
	.asciz	"Calling kernels to calculate a subset of input matrices asynchronously"
	.size	.Lstr.18, 71

	.type	.Lstr.19,@object                # @str.19
.Lstr.19:
	.asciz	"output arrays accounted for, will be used to aggregate results"
	.size	.Lstr.19, 63

	.type	.Lstr.20,@object                # @str.20
.Lstr.20:
	.asciz	"Calling last kernel to aggregate partial results"
	.size	.Lstr.20, 49

	.type	.Lstr.21,@object                # @str.21
.Lstr.21:
	.asciz	"usage\n ./multi_matmult <filename> <algorithm>"
	.size	.Lstr.21, 46

	.type	.Lstr.22,@object                # @str.22
.Lstr.22:
	.asciz	"algorithm:\n\t 0 - cpu (sequential), 1 - gpu (parallel, one thread block), 2 - (parallel, multiple thread blocks)"
	.size	.Lstr.22, 112

	.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 _Z36__device_stub__gpu_seq_multi_matmultiPiPS_S_S_
	.addrsig_sym _Z36__device_stub__gpu_par_multi_matmultiiPiPS_S_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym NUM_DIM
	.addrsig_sym _Z21gpu_seq_multi_matmultiPiPS_S_S_
	.addrsig_sym _Z21gpu_par_multi_matmultiiPiPS_S_S_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z21gpu_seq_multi_matmultiPiPS_S_S_ ; -- Begin function _Z21gpu_seq_multi_matmultiPiPS_S_S_
	.globl	_Z21gpu_seq_multi_matmultiPiPS_S_S_
	.p2align	8
	.type	_Z21gpu_seq_multi_matmultiPiPS_S_S_,@function
_Z21gpu_seq_multi_matmultiPiPS_S_S_:    ; @_Z21gpu_seq_multi_matmultiPiPS_S_S_
; %bb.0:
	s_load_b256 s[4:11], s[0:1], 0x8
	s_mov_b32 s21, exec_lo
	s_waitcnt lgkmcnt(0)
	s_clause 0x1
	s_load_b64 s[2:3], s[4:5], 0x0
	s_load_b32 s16, s[4:5], 0x8
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s20, s16, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	v_cmpx_gt_i32_e64 s20, v0
	s_cbranch_execz .LBB0_6
; %bb.1:                                ; %.lr.ph34.i
	s_cmp_gt_i32 s3, 0
	s_load_b128 s[12:15], s[6:7], 0x0
	s_cselect_b32 s22, -1, 0
	s_ashr_i32 s17, s16, 31
	s_mov_b32 s24, 0
	s_add_i32 s2, s16, s17
	s_lshl_b64 s[18:19], s[16:17], 2
	s_xor_b32 s23, s2, s17
	v_mov_b32_e32 v7, v0
	v_cvt_f32_u32_e32 v1, s23
	s_sub_i32 s2, 0, s23
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, s2, v1
	v_mul_hi_u32 v2, v1, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v1, v2
.LBB0_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v1, v7, v6
	v_add_nc_u32_e32 v3, s16, v7
	v_mov_b32_e32 v9, 0
	v_mul_lo_u32 v1, v1, s23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v1, v7, v1
	v_subrev_nc_u32_e32 v2, s23, v1
	v_cmp_le_u32_e32 vcc_lo, s23, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_subrev_nc_u32_e32 v2, s23, v1
	v_cmp_le_u32_e32 vcc_lo, s23, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_sub_nc_u32_e32 v2, v3, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_add_nc_u32_e32 v2, v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v2, v2, v3
	v_xor_b32_e32 v3, s17, v3
	v_mul_hi_u32 v4, v2, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v5, v4, s23
	v_sub_nc_u32_e32 v2, v2, v5
	v_add_nc_u32_e32 v5, 1, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v8, s23, v2
	v_cmp_le_u32_e32 vcc_lo, s23, v2
	v_cndmask_b32_e32 v2, v2, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s23, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, 1, v4
	v_cndmask_b32_e32 v2, v4, v5, vcc_lo
	s_and_not1_b32 vcc_lo, exec_lo, s22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v3
	v_xad_u32 v8, v3, -1, v2
	s_cbranch_vccnz .LBB0_5
; %bb.3:                                ; %.lr.ph.i
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, v8, s3
	v_ashrrev_i32_e32 v2, 31, v1
	s_mov_b32 s2, s3
	v_lshlrev_b64 v[9:10], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v4, 31, v3
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s14, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[4:5], 2, v[3:4]
	v_add_co_ci_u32_e32 v3, vcc_lo, s15, v10, vcc_lo
	v_mov_b32_e32 v9, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v4, vcc_lo, s12, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s13, v5, vcc_lo
.LBB0_4:                                ;   Parent Loop BB0_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v10, v[4:5], off
	global_load_b32 v11, v[2:3], off
	v_add_co_u32 v2, vcc_lo, v2, s18
	v_add_co_ci_u32_e32 v3, vcc_lo, s19, v3, vcc_lo
	v_add_co_u32 v4, vcc_lo, v4, 4
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_add_i32 s2, s2, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_cmp_eq_u32 s2, 0
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v10, v11, v10
	v_mul_hi_i32 v11, 0x68c8c4ad, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v12, 31, v11
	v_ashrrev_i32_e32 v11, 12, v11
	v_add_nc_u32_e32 v11, v11, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_i32_i24_e32 v11, 0x2717, v11
	v_sub_nc_u32_e32 v10, v10, v11
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v9, v10, v9
	s_cbranch_scc0 .LBB0_4
.LBB0_5:                                ; %._crit_edge.i
                                        ;   in Loop: Header=BB0_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_i32 v4, 0x68c8c4ad, v9
	v_mad_u64_u32 v[2:3], null, v8, s16, v[1:2]
	v_add_nc_u32_e32 v7, 0x400, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_le_i32_e32 vcc_lo, s20, v7
	v_lshrrev_b32_e32 v1, 31, v4
	v_ashrrev_i32_e32 v4, 12, v4
	v_ashrrev_i32_e32 v3, 31, v2
	s_or_b32 s24, vcc_lo, s24
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v4, v1
	v_lshlrev_b64 v[1:2], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_i32_i24_e32 v3, 0x2717, v4
	v_add_co_u32 v1, s2, s8, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v2, s2, s9, v2, s2
	v_sub_nc_u32_e32 v3, v9, v3
	global_store_b32 v[1:2], v3, off
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_and_not1_b32 exec_lo, exec_lo, s24
	s_cbranch_execnz .LBB0_2
.LBB0_6:                                ; %Flow141
	s_or_b32 exec_lo, exec_lo, s21
	s_load_b32 s19, s[0:1], 0x0
	s_waitcnt lgkmcnt(0)
	s_barrier
	buffer_gl0_inv
	s_barrier
	buffer_gl0_inv
	s_cmp_lt_i32 s19, 4
	s_cbranch_scc1 .LBB0_27
; %bb.7:                                ; %.lr.ph.preheader
	v_mov_b32_e32 v6, 0
	s_add_u32 s1, s6, 8
	s_addc_u32 s16, s7, 0
	s_add_u32 s17, s4, 8
	s_addc_u32 s18, s5, 0
	s_clause 0x1
	global_load_b32 v1, v6, s[4:5] offset:8
	global_load_b32 v7, v6, s[4:5]
	s_add_i32 s19, s19, -3
	s_mov_b32 s2, 1
	s_mov_b32 s3, 0
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s20, v1
.LBB0_8:                                ; %.lr.ph
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_11 Depth 2
                                        ;       Child Loop BB0_13 Depth 3
                                        ;     Child Loop BB0_20 Depth 2
                                        ;       Child Loop BB0_22 Depth 3
	s_bitcmp0_b32 s2, 0
	s_mov_b32 s0, -1
	s_cbranch_scc1 .LBB0_17
; %bb.9:                                ;   in Loop: Header=BB0_8 Depth=1
	s_lshl_b64 s[4:5], s[2:3], 2
	s_mov_b32 s21, exec_lo
	s_add_u32 s4, s17, s4
	s_addc_u32 s5, s18, s5
	global_load_b32 v1, v6, s[4:5]
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v8, v1, v7
	v_readfirstlane_b32 s4, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_lt_i32_e64 v0, v8
	s_cbranch_execz .LBB0_16
; %bb.10:                               ; %.lr.ph34.i39
                                        ;   in Loop: Header=BB0_8 Depth=1
	s_lshl_b64 s[12:13], s[2:3], 3
	s_mov_b32 s24, 0
	s_add_u32 s12, s1, s12
	s_addc_u32 s13, s16, s13
	s_cmp_gt_i32 s20, 0
	s_load_b64 s[12:13], s[12:13], 0x0
	s_cselect_b32 s22, -1, 0
	s_ashr_i32 s5, s4, 31
	v_mov_b32_e32 v10, v0
	s_add_i32 s0, s4, s5
	s_lshl_b64 s[14:15], s[4:5], 2
	s_xor_b32 s23, s0, s5
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cvt_f32_u32_e32 v1, s23
	s_sub_i32 s0, 0, s23
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v1, v1
	v_mul_lo_u32 v2, s0, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v1, v2
	v_add_nc_u32_e32 v9, v1, v2
.LBB0_11:                               ;   Parent Loop BB0_8 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_13 Depth 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v1, v10, v9
	v_add_nc_u32_e32 v3, s4, v10
	v_mul_lo_u32 v1, v1, s23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v1, v10, v1
	v_subrev_nc_u32_e32 v2, s23, v1
	v_cmp_le_u32_e32 vcc_lo, s23, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_subrev_nc_u32_e32 v2, s23, v1
	v_cmp_le_u32_e32 vcc_lo, s23, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_sub_nc_u32_e32 v2, v3, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_add_nc_u32_e32 v2, v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v2, v2, v3
	v_xor_b32_e32 v3, s5, v3
	v_mul_hi_u32 v4, v2, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v5, v4, s23
	v_sub_nc_u32_e32 v2, v2, v5
	v_add_nc_u32_e32 v5, 1, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v11, s23, v2
	v_cmp_le_u32_e32 vcc_lo, s23, v2
	v_cndmask_b32_e32 v2, v2, v11, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s23, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, 1, v4
	v_cndmask_b32_e32 v2, v4, v5, vcc_lo
	s_and_not1_b32 vcc_lo, exec_lo, s22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v3
	v_xad_u32 v11, v3, -1, v2
	s_cbranch_vccnz .LBB0_14
; %bb.12:                               ; %.lr.ph.i43
                                        ;   in Loop: Header=BB0_11 Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, v11, s20
	v_ashrrev_i32_e32 v2, 31, v1
	s_mov_b32 s0, s20
	v_lshlrev_b64 v[12:13], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v4, 31, v3
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s12, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[4:5], 2, v[3:4]
	v_add_co_ci_u32_e32 v3, vcc_lo, s13, v13, vcc_lo
	v_mov_b32_e32 v12, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v4, vcc_lo, s8, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s9, v5, vcc_lo
.LBB0_13:                               ;   Parent Loop BB0_8 Depth=1
                                        ;     Parent Loop BB0_11 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	global_load_b32 v13, v[4:5], off
	global_load_b32 v14, v[2:3], off
	v_add_co_u32 v2, vcc_lo, v2, s14
	v_add_co_ci_u32_e32 v3, vcc_lo, s15, v3, vcc_lo
	v_add_co_u32 v4, vcc_lo, v4, 4
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_add_i32 s0, s0, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_cmp_eq_u32 s0, 0
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v13, v14, v13
	v_mul_hi_i32 v14, 0x68c8c4ad, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v15, 31, v14
	v_ashrrev_i32_e32 v14, 12, v14
	v_add_nc_u32_e32 v14, v14, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_i32_i24_e32 v14, 0x2717, v14
	v_sub_nc_u32_e32 v13, v13, v14
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v12, v13, v12
	s_cbranch_scc0 .LBB0_13
	s_branch .LBB0_15
.LBB0_14:                               ;   in Loop: Header=BB0_11 Depth=2
	v_mov_b32_e32 v12, 0
.LBB0_15:                               ; %._crit_edge.i41
                                        ;   in Loop: Header=BB0_11 Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_i32 v4, 0x68c8c4ad, v12
	v_mad_u64_u32 v[2:3], null, v11, s4, v[1:2]
	v_add_nc_u32_e32 v10, 0x400, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_ge_i32_e32 vcc_lo, v10, v8
	v_lshrrev_b32_e32 v1, 31, v4
	v_ashrrev_i32_e32 v4, 12, v4
	v_ashrrev_i32_e32 v3, 31, v2
	s_or_b32 s24, vcc_lo, s24
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v4, v1
	v_lshlrev_b64 v[1:2], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_i32_i24_e32 v3, 0x2717, v4
	v_add_co_u32 v1, s0, s10, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v2, s0, s11, v2, s0
	v_sub_nc_u32_e32 v3, v12, v3
	global_store_b32 v[1:2], v3, off
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_and_not1_b32 exec_lo, exec_lo, s24
	s_cbranch_execnz .LBB0_11
.LBB0_16:                               ; %Flow129
                                        ;   in Loop: Header=BB0_8 Depth=1
	s_or_b32 exec_lo, exec_lo, s21
	s_mov_b32 s0, 0
	s_barrier
.LBB0_17:                               ; %Flow134
                                        ;   in Loop: Header=BB0_8 Depth=1
	s_and_b32 vcc_lo, exec_lo, s0
	s_cbranch_vccz .LBB0_26
; %bb.18:                               ;   in Loop: Header=BB0_8 Depth=1
	s_lshl_b64 s[4:5], s[2:3], 2
	s_mov_b32 s21, exec_lo
	s_add_u32 s4, s17, s4
	s_addc_u32 s5, s18, s5
	global_load_b32 v1, v6, s[4:5]
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v8, v1, v7
	v_readfirstlane_b32 s4, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_lt_i32_e64 v0, v8
	s_cbranch_execz .LBB0_25
; %bb.19:                               ; %.lr.ph34.i48
                                        ;   in Loop: Header=BB0_8 Depth=1
	s_or_b32 s12, s2, 1
	s_mov_b32 s13, s3
	s_mov_b32 s24, 0
	s_lshl_b64 s[12:13], s[12:13], 3
	v_mov_b32_e32 v10, v0
	s_add_u32 s12, s6, s12
	s_addc_u32 s13, s7, s13
	s_cmp_gt_i32 s20, 0
	s_load_b64 s[12:13], s[12:13], 0x0
	s_cselect_b32 s22, -1, 0
	s_ashr_i32 s5, s4, 31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	s_add_i32 s0, s4, s5
	s_lshl_b64 s[14:15], s[4:5], 2
	s_xor_b32 s23, s0, s5
	v_cvt_f32_u32_e32 v1, s23
	s_sub_i32 s0, 0, s23
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, s0, v1
	v_mul_hi_u32 v2, v1, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v9, v1, v2
.LBB0_20:                               ;   Parent Loop BB0_8 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_22 Depth 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v1, v10, v9
	v_add_nc_u32_e32 v3, s4, v10
	v_mul_lo_u32 v1, v1, s23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v1, v10, v1
	v_subrev_nc_u32_e32 v2, s23, v1
	v_cmp_le_u32_e32 vcc_lo, s23, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_subrev_nc_u32_e32 v2, s23, v1
	v_cmp_le_u32_e32 vcc_lo, s23, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_sub_nc_u32_e32 v2, v3, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_add_nc_u32_e32 v2, v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v2, v2, v3
	v_xor_b32_e32 v3, s5, v3
	v_mul_hi_u32 v4, v2, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v5, v4, s23
	v_sub_nc_u32_e32 v2, v2, v5
	v_add_nc_u32_e32 v5, 1, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v11, s23, v2
	v_cmp_le_u32_e32 vcc_lo, s23, v2
	v_cndmask_b32_e32 v2, v2, v11, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s23, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, 1, v4
	v_cndmask_b32_e32 v2, v4, v5, vcc_lo
	s_and_not1_b32 vcc_lo, exec_lo, s22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v3
	v_xad_u32 v11, v3, -1, v2
	s_cbranch_vccnz .LBB0_23
; %bb.21:                               ; %.lr.ph.i52
                                        ;   in Loop: Header=BB0_20 Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, v11, s20
	v_ashrrev_i32_e32 v2, 31, v1
	s_mov_b32 s0, s20
	v_lshlrev_b64 v[12:13], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v4, 31, v3
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s12, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[4:5], 2, v[3:4]
	v_add_co_ci_u32_e32 v3, vcc_lo, s13, v13, vcc_lo
	v_mov_b32_e32 v12, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v4, vcc_lo, s10, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v5, vcc_lo
.LBB0_22:                               ;   Parent Loop BB0_8 Depth=1
                                        ;     Parent Loop BB0_20 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	global_load_b32 v13, v[4:5], off
	global_load_b32 v14, v[2:3], off
	v_add_co_u32 v2, vcc_lo, v2, s14
	v_add_co_ci_u32_e32 v3, vcc_lo, s15, v3, vcc_lo
	v_add_co_u32 v4, vcc_lo, v4, 4
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_add_i32 s0, s0, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_cmp_eq_u32 s0, 0
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v13, v14, v13
	v_mul_hi_i32 v14, 0x68c8c4ad, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v15, 31, v14
	v_ashrrev_i32_e32 v14, 12, v14
	v_add_nc_u32_e32 v14, v14, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_i32_i24_e32 v14, 0x2717, v14
	v_sub_nc_u32_e32 v13, v13, v14
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v12, v13, v12
	s_cbranch_scc0 .LBB0_22
	s_branch .LBB0_24
.LBB0_23:                               ;   in Loop: Header=BB0_20 Depth=2
	v_mov_b32_e32 v12, 0
.LBB0_24:                               ; %._crit_edge.i50
                                        ;   in Loop: Header=BB0_20 Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_i32 v4, 0x68c8c4ad, v12
	v_mad_u64_u32 v[2:3], null, v11, s4, v[1:2]
	v_add_nc_u32_e32 v10, 0x400, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_ge_i32_e32 vcc_lo, v10, v8
	v_lshrrev_b32_e32 v1, 31, v4
	v_ashrrev_i32_e32 v4, 12, v4
	v_ashrrev_i32_e32 v3, 31, v2
	s_or_b32 s24, vcc_lo, s24
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v4, v1
	v_lshlrev_b64 v[1:2], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_i32_i24_e32 v3, 0x2717, v4
	v_add_co_u32 v1, s0, s8, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v2, s0, s9, v2, s0
	v_sub_nc_u32_e32 v3, v12, v3
	global_store_b32 v[1:2], v3, off
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_and_not1_b32 exec_lo, exec_lo, s24
	s_cbranch_execnz .LBB0_20
.LBB0_25:                               ; %Flow133
                                        ;   in Loop: Header=BB0_8 Depth=1
	s_or_b32 exec_lo, exec_lo, s21
	s_barrier
.LBB0_26:                               ;   in Loop: Header=BB0_8 Depth=1
	s_lshl_b64 s[4:5], s[2:3], 2
	s_waitcnt vmcnt(0)
	buffer_gl0_inv
	s_add_u32 s4, s17, s4
	s_addc_u32 s5, s18, s5
	s_add_i32 s0, s2, 1
	global_load_b32 v1, v6, s[4:5]
	s_cmp_lg_u32 s2, s19
	s_mov_b32 s2, s0
	s_waitcnt vmcnt(0)
	s_barrier
	buffer_gl0_inv
	v_readfirstlane_b32 s20, v1
	s_cbranch_scc1 .LBB0_8
.LBB0_27:                               ; %._crit_edge
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z21gpu_seq_multi_matmultiPiPS_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 40
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 16
		.amdhsa_next_free_sgpr 25
		.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	_Z21gpu_seq_multi_matmultiPiPS_S_S_, .Lfunc_end0-_Z21gpu_seq_multi_matmultiPiPS_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 2344
; NumSgprs: 27
; NumVgprs: 16
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 27
; NumVGPRsForWavesPerEU: 16
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z21gpu_par_multi_matmultiiPiPS_S_S_ ; -- Begin function _Z21gpu_par_multi_matmultiiPiPS_S_S_
	.globl	_Z21gpu_par_multi_matmultiiPiPS_S_S_
	.p2align	8
	.type	_Z21gpu_par_multi_matmultiiPiPS_S_S_,@function
_Z21gpu_par_multi_matmultiiPiPS_S_S_:   ; @_Z21gpu_par_multi_matmultiiPiPS_S_S_
; %bb.0:
	s_clause 0x1
	s_load_b64 s[14:15], s[0:1], 0x0
	s_load_b256 s[0:7], s[0:1], 0x8
	s_mov_b32 s21, exec_lo
	s_waitcnt lgkmcnt(0)
	s_ashr_i32 s9, s14, 31
	s_mov_b32 s8, s14
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_lshl_b64 s[10:11], s[8:9], 3
	s_add_u32 s2, s2, s10
	s_addc_u32 s3, s3, s11
	s_lshl_b64 s[8:9], s[8:9], 2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s12, s0, s8
	s_addc_u32 s13, s1, s9
	s_clause 0x1
	s_load_b64 s[0:1], s[12:13], 0x0
	s_load_b32 s16, s[12:13], 0x8
	s_waitcnt lgkmcnt(0)
	s_mul_i32 s20, s16, s0
	s_delay_alu instid0(SALU_CYCLE_1)
	v_cmpx_gt_i32_e64 s20, v0
	s_cbranch_execz .LBB1_6
; %bb.1:                                ; %.lr.ph34.i
	s_cmp_gt_i32 s1, 0
	s_load_b128 s[8:11], s[2:3], 0x0
	s_cselect_b32 s22, -1, 0
	s_ashr_i32 s17, s16, 31
	s_mov_b32 s24, 0
	s_add_i32 s0, s16, s17
	s_lshl_b64 s[18:19], s[16:17], 2
	s_xor_b32 s23, s0, s17
	v_mov_b32_e32 v7, v0
	v_cvt_f32_u32_e32 v1, s23
	s_sub_i32 s0, 0, s23
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	v_cvt_u32_f32_e32 v1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v2, s0, v1
	v_mul_hi_u32 v2, v1, v2
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v1, v2
.LBB1_2:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_4 Depth 2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_mul_hi_u32 v1, v7, v6
	v_add_nc_u32_e32 v3, s16, v7
	v_mov_b32_e32 v9, 0
	v_mul_lo_u32 v1, v1, s23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v1, v7, v1
	v_subrev_nc_u32_e32 v2, s23, v1
	v_cmp_le_u32_e32 vcc_lo, s23, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_subrev_nc_u32_e32 v2, s23, v1
	v_cmp_le_u32_e32 vcc_lo, s23, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_sub_nc_u32_e32 v2, v3, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_add_nc_u32_e32 v2, v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v2, v2, v3
	v_xor_b32_e32 v3, s17, v3
	v_mul_hi_u32 v4, v2, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v5, v4, s23
	v_sub_nc_u32_e32 v2, v2, v5
	v_add_nc_u32_e32 v5, 1, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v8, s23, v2
	v_cmp_le_u32_e32 vcc_lo, s23, v2
	v_cndmask_b32_e32 v2, v2, v8, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s23, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, 1, v4
	v_cndmask_b32_e32 v2, v4, v5, vcc_lo
	s_and_not1_b32 vcc_lo, exec_lo, s22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v3
	v_xad_u32 v8, v3, -1, v2
	s_cbranch_vccnz .LBB1_5
; %bb.3:                                ; %.lr.ph.i
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, v8, s1
	v_ashrrev_i32_e32 v2, 31, v1
	s_mov_b32 s0, s1
	v_lshlrev_b64 v[9:10], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v4, 31, v3
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s10, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[4:5], 2, v[3:4]
	v_add_co_ci_u32_e32 v3, vcc_lo, s11, v10, vcc_lo
	v_mov_b32_e32 v9, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v4, vcc_lo, s8, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s9, v5, vcc_lo
.LBB1_4:                                ;   Parent Loop BB1_2 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	global_load_b32 v10, v[4:5], off
	global_load_b32 v11, v[2:3], off
	v_add_co_u32 v2, vcc_lo, v2, s18
	v_add_co_ci_u32_e32 v3, vcc_lo, s19, v3, vcc_lo
	v_add_co_u32 v4, vcc_lo, v4, 4
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_add_i32 s0, s0, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_cmp_eq_u32 s0, 0
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v10, v11, v10
	v_mul_hi_i32 v11, 0x68c8c4ad, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v12, 31, v11
	v_ashrrev_i32_e32 v11, 12, v11
	v_add_nc_u32_e32 v11, v11, v12
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_i32_i24_e32 v11, 0x2717, v11
	v_sub_nc_u32_e32 v10, v10, v11
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v9, v10, v9
	s_cbranch_scc0 .LBB1_4
.LBB1_5:                                ; %._crit_edge.i
                                        ;   in Loop: Header=BB1_2 Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_hi_i32 v4, 0x68c8c4ad, v9
	v_mad_u64_u32 v[2:3], null, v8, s16, v[1:2]
	v_add_nc_u32_e32 v7, 0x400, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_le_i32_e32 vcc_lo, s20, v7
	v_lshrrev_b32_e32 v1, 31, v4
	v_ashrrev_i32_e32 v4, 12, v4
	v_ashrrev_i32_e32 v3, 31, v2
	s_or_b32 s24, vcc_lo, s24
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v4, v1
	v_lshlrev_b64 v[1:2], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_i32_i24_e32 v3, 0x2717, v4
	v_add_co_u32 v1, s0, s4, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v2, s0, s5, v2, s0
	v_sub_nc_u32_e32 v3, v9, v3
	global_store_b32 v[1:2], v3, off
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_and_not1_b32 exec_lo, exec_lo, s24
	s_cbranch_execnz .LBB1_2
.LBB1_6:                                ; %Flow161
	s_or_b32 exec_lo, exec_lo, s21
	s_sub_i32 s16, s15, s14
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_i32 s16, s16, -2
	s_barrier
	s_cmp_lt_i32 s16, 2
	buffer_gl0_inv
	s_barrier
	buffer_gl0_inv
	s_cbranch_scc1 .LBB1_27
; %bb.7:                                ; %.lr.ph.preheader
	v_mov_b32_e32 v6, 0
	s_add_u32 s17, s2, 8
	s_addc_u32 s18, s3, 0
	s_add_u32 s19, s12, 8
	s_addc_u32 s20, s13, 0
	s_clause 0x1
	global_load_b32 v1, v6, s[12:13] offset:8
	global_load_b32 v7, v6, s[12:13]
	s_mov_b32 s2, 1
	s_waitcnt vmcnt(1)
	v_readfirstlane_b32 s21, v1
.LBB1_8:                                ; %.lr.ph
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB1_11 Depth 2
                                        ;       Child Loop BB1_13 Depth 3
                                        ;     Child Loop BB1_20 Depth 2
                                        ;       Child Loop BB1_22 Depth 3
	s_ashr_i32 s3, s2, 31
	s_and_b32 s13, s2, 1
	s_lshl_b64 s[0:1], s[2:3], 3
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s0, s17, s0
	s_addc_u32 s1, s18, s1
	s_lshl_b64 s[8:9], s[2:3], 2
	s_load_b64 s[10:11], s[0:1], 0x0
	s_add_u32 s8, s19, s8
	s_addc_u32 s9, s20, s9
	s_cmp_eq_u32 s13, 0
	global_load_b32 v1, v6, s[8:9]
	s_mov_b32 s1, -1
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v8, v1, v7
	v_readfirstlane_b32 s12, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_cmp_lt_i32_e64 s0, v0, v8
	s_cbranch_scc1 .LBB1_17
; %bb.9:                                ;   in Loop: Header=BB1_8 Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s3, s0
	s_cbranch_execz .LBB1_16
; %bb.10:                               ; %.lr.ph34.i55
                                        ;   in Loop: Header=BB1_8 Depth=1
	s_cmp_gt_i32 s21, 0
	s_mov_b32 s24, 0
	s_cselect_b32 s22, -1, 0
	s_ashr_i32 s13, s12, 31
	v_mov_b32_e32 v10, v0
	s_add_i32 s1, s12, s13
	s_lshl_b64 s[14:15], s[12:13], 2
	s_xor_b32 s23, s1, s13
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cvt_f32_u32_e32 v1, s23
	s_sub_i32 s1, 0, s23
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v1, v1
	v_mul_lo_u32 v2, s1, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v1, v2
	v_add_nc_u32_e32 v9, v1, v2
.LBB1_11:                               ;   Parent Loop BB1_8 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB1_13 Depth 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v1, v10, v9
	v_add_nc_u32_e32 v3, s12, v10
	v_mul_lo_u32 v1, v1, s23
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v1, v10, v1
	v_subrev_nc_u32_e32 v2, s23, v1
	v_cmp_le_u32_e32 vcc_lo, s23, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_subrev_nc_u32_e32 v2, s23, v1
	v_cmp_le_u32_e32 vcc_lo, s23, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_sub_nc_u32_e32 v2, v3, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_add_nc_u32_e32 v2, v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v2, v2, v3
	v_xor_b32_e32 v3, s13, v3
	v_mul_hi_u32 v4, v2, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v5, v4, s23
	v_sub_nc_u32_e32 v2, v2, v5
	v_add_nc_u32_e32 v5, 1, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v11, s23, v2
	v_cmp_le_u32_e32 vcc_lo, s23, v2
	v_cndmask_b32_e32 v2, v2, v11, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s23, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, 1, v4
	v_cndmask_b32_e32 v2, v4, v5, vcc_lo
	s_and_not1_b32 vcc_lo, exec_lo, s22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v3
	v_xad_u32 v11, v3, -1, v2
	s_cbranch_vccnz .LBB1_14
; %bb.12:                               ; %.lr.ph.i59
                                        ;   in Loop: Header=BB1_11 Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, v11, s21
	v_ashrrev_i32_e32 v2, 31, v1
	s_mov_b32 s1, s21
	v_lshlrev_b64 v[12:13], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v4, 31, v3
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s10, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[4:5], 2, v[3:4]
	v_add_co_ci_u32_e32 v3, vcc_lo, s11, v13, vcc_lo
	v_mov_b32_e32 v12, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v4, vcc_lo, s4, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s5, v5, vcc_lo
.LBB1_13:                               ;   Parent Loop BB1_8 Depth=1
                                        ;     Parent Loop BB1_11 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	global_load_b32 v13, v[4:5], off
	global_load_b32 v14, v[2:3], off
	v_add_co_u32 v2, vcc_lo, v2, s14
	v_add_co_ci_u32_e32 v3, vcc_lo, s15, v3, vcc_lo
	v_add_co_u32 v4, vcc_lo, v4, 4
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_add_i32 s1, s1, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_cmp_eq_u32 s1, 0
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v13, v14, v13
	v_mul_hi_i32 v14, 0x68c8c4ad, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v15, 31, v14
	v_ashrrev_i32_e32 v14, 12, v14
	v_add_nc_u32_e32 v14, v14, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_i32_i24_e32 v14, 0x2717, v14
	v_sub_nc_u32_e32 v13, v13, v14
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v12, v13, v12
	s_cbranch_scc0 .LBB1_13
	s_branch .LBB1_15
.LBB1_14:                               ;   in Loop: Header=BB1_11 Depth=2
	v_mov_b32_e32 v12, 0
.LBB1_15:                               ; %._crit_edge.i57
                                        ;   in Loop: Header=BB1_11 Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_i32 v4, 0x68c8c4ad, v12
	v_mad_u64_u32 v[2:3], null, v11, s12, v[1:2]
	v_add_nc_u32_e32 v10, 0x400, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_ge_i32_e32 vcc_lo, v10, v8
	v_lshrrev_b32_e32 v1, 31, v4
	v_ashrrev_i32_e32 v4, 12, v4
	v_ashrrev_i32_e32 v3, 31, v2
	s_or_b32 s24, vcc_lo, s24
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v4, v1
	v_lshlrev_b64 v[1:2], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_i32_i24_e32 v3, 0x2717, v4
	v_add_co_u32 v1, s1, s6, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v2, s1, s7, v2, s1
	v_sub_nc_u32_e32 v3, v12, v3
	global_store_b32 v[1:2], v3, off
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_and_not1_b32 exec_lo, exec_lo, s24
	s_cbranch_execnz .LBB1_11
.LBB1_16:                               ; %Flow149
                                        ;   in Loop: Header=BB1_8 Depth=1
	s_or_b32 exec_lo, exec_lo, s3
	s_mov_b32 s1, 0
	s_waitcnt lgkmcnt(0)
	s_barrier
.LBB1_17:                               ; %Flow154
                                        ;   in Loop: Header=BB1_8 Depth=1
	s_and_b32 vcc_lo, exec_lo, s1
	s_cbranch_vccz .LBB1_26
; %bb.18:                               ;   in Loop: Header=BB1_8 Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	s_and_saveexec_b32 s1, s0
	s_cbranch_execz .LBB1_25
; %bb.19:                               ; %.lr.ph34.i64
                                        ;   in Loop: Header=BB1_8 Depth=1
	s_cmp_gt_i32 s21, 0
	s_mov_b32 s23, 0
	s_cselect_b32 s3, -1, 0
	s_ashr_i32 s13, s12, 31
	v_mov_b32_e32 v10, v0
	s_add_i32 s0, s12, s13
	s_lshl_b64 s[14:15], s[12:13], 2
	s_xor_b32 s22, s0, s13
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cvt_f32_u32_e32 v1, s22
	s_sub_i32 s0, 0, s22
	v_rcp_iflag_f32_e32 v1, v1
	s_waitcnt_depctr 0xfff
	v_mul_f32_e32 v1, 0x4f7ffffe, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cvt_u32_f32_e32 v1, v1
	v_mul_lo_u32 v2, s0, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_hi_u32 v2, v1, v2
	v_add_nc_u32_e32 v9, v1, v2
.LBB1_20:                               ;   Parent Loop BB1_8 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB1_22 Depth 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_hi_u32 v1, v10, v9
	v_add_nc_u32_e32 v3, s12, v10
	v_mul_lo_u32 v1, v1, s22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_sub_nc_u32_e32 v1, v10, v1
	v_subrev_nc_u32_e32 v2, s22, v1
	v_cmp_le_u32_e32 vcc_lo, s22, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_subrev_nc_u32_e32 v2, s22, v1
	v_cmp_le_u32_e32 vcc_lo, s22, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v1, v1, v2, vcc_lo
	v_sub_nc_u32_e32 v2, v3, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_add_nc_u32_e32 v2, v2, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v2, v2, v3
	v_xor_b32_e32 v3, s13, v3
	v_mul_hi_u32 v4, v2, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v5, v4, s22
	v_sub_nc_u32_e32 v2, v2, v5
	v_add_nc_u32_e32 v5, 1, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_subrev_nc_u32_e32 v11, s22, v2
	v_cmp_le_u32_e32 vcc_lo, s22, v2
	v_cndmask_b32_e32 v2, v2, v11, vcc_lo
	s_delay_alu instid0(VALU_DEP_4) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	v_cmp_le_u32_e32 vcc_lo, s22, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, 1, v4
	v_cndmask_b32_e32 v2, v4, v5, vcc_lo
	s_and_not1_b32 vcc_lo, exec_lo, s3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v2, v2, v3
	v_xad_u32 v11, v3, -1, v2
	s_cbranch_vccnz .LBB1_23
; %bb.21:                               ; %.lr.ph.i68
                                        ;   in Loop: Header=BB1_20 Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_lo_u32 v3, v11, s21
	v_ashrrev_i32_e32 v2, 31, v1
	s_mov_b32 s0, s21
	v_lshlrev_b64 v[12:13], 2, v[1:2]
	s_delay_alu instid0(VALU_DEP_3) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v4, 31, v3
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s10, v12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_lshlrev_b64 v[4:5], 2, v[3:4]
	v_add_co_ci_u32_e32 v3, vcc_lo, s11, v13, vcc_lo
	v_mov_b32_e32 v12, 0
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v4, vcc_lo, s6, v4
	v_add_co_ci_u32_e32 v5, vcc_lo, s7, v5, vcc_lo
.LBB1_22:                               ;   Parent Loop BB1_8 Depth=1
                                        ;     Parent Loop BB1_20 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	global_load_b32 v13, v[4:5], off
	global_load_b32 v14, v[2:3], off
	v_add_co_u32 v2, vcc_lo, v2, s14
	v_add_co_ci_u32_e32 v3, vcc_lo, s15, v3, vcc_lo
	v_add_co_u32 v4, vcc_lo, v4, 4
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_add_i32 s0, s0, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	s_cmp_eq_u32 s0, 0
	s_waitcnt vmcnt(0)
	v_mul_lo_u32 v13, v14, v13
	v_mul_hi_i32 v14, 0x68c8c4ad, v13
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshrrev_b32_e32 v15, 31, v14
	v_ashrrev_i32_e32 v14, 12, v14
	v_add_nc_u32_e32 v14, v14, v15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_i32_i24_e32 v14, 0x2717, v14
	v_sub_nc_u32_e32 v13, v13, v14
	s_delay_alu instid0(VALU_DEP_1)
	v_add_nc_u32_e32 v12, v13, v12
	s_cbranch_scc0 .LBB1_22
	s_branch .LBB1_24
.LBB1_23:                               ;   in Loop: Header=BB1_20 Depth=2
	v_mov_b32_e32 v12, 0
.LBB1_24:                               ; %._crit_edge.i66
                                        ;   in Loop: Header=BB1_20 Depth=2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mul_hi_i32 v4, 0x68c8c4ad, v12
	v_mad_u64_u32 v[2:3], null, v11, s12, v[1:2]
	v_add_nc_u32_e32 v10, 0x400, v10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_cmp_ge_i32_e32 vcc_lo, v10, v8
	v_lshrrev_b32_e32 v1, 31, v4
	v_ashrrev_i32_e32 v4, 12, v4
	v_ashrrev_i32_e32 v3, 31, v2
	s_or_b32 s23, vcc_lo, s23
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v4, v4, v1
	v_lshlrev_b64 v[1:2], 2, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_i32_i24_e32 v3, 0x2717, v4
	v_add_co_u32 v1, s0, s4, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e64 v2, s0, s5, v2, s0
	v_sub_nc_u32_e32 v3, v12, v3
	global_store_b32 v[1:2], v3, off
	s_waitcnt lgkmcnt(0)
	s_waitcnt_vscnt null, 0x0
	s_barrier
	buffer_gl0_inv
	s_and_not1_b32 exec_lo, exec_lo, s23
	s_cbranch_execnz .LBB1_20
.LBB1_25:                               ; %Flow153
                                        ;   in Loop: Header=BB1_8 Depth=1
	s_or_b32 exec_lo, exec_lo, s1
	s_waitcnt lgkmcnt(0)
	s_barrier
.LBB1_26:                               ;   in Loop: Header=BB1_8 Depth=1
	s_waitcnt lgkmcnt(0)
	buffer_gl0_inv
	global_load_b32 v1, v6, s[8:9]
	s_add_i32 s2, s2, 1
	s_waitcnt vmcnt(0)
	s_cmp_ge_i32 s2, s16
	s_barrier
	buffer_gl0_inv
	v_readfirstlane_b32 s21, v1
	s_cbranch_scc0 .LBB1_8
.LBB1_27:                               ; %._crit_edge
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z21gpu_par_multi_matmultiiPiPS_S_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 40
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 16
		.amdhsa_next_free_sgpr 25
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z21gpu_par_multi_matmultiiPiPS_S_S_, .Lfunc_end1-_Z21gpu_par_multi_matmultiiPiPS_S_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 2320
; NumSgprs: 27
; NumVgprs: 16
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 3
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 27
; NumVGPRsForWavesPerEU: 16
; Occupancy: 16
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

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

	.end_amdgpu_metadata
