/*********************
MIT License

Copyright (c) 2020 Matzoros Christos

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in all
copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
SOFTWARE.

***********************/

/*********************
The purpose of this code is to execute multi-GPU matrix multiplication with multiple kernel 
invocations using streams. The program split the computation into 4 individual computations 
as it is shown below. The proportion of the size of the block is variable.
  
  ------------------------------------------------
    A * B = C   

    |  A1  |     |    |    |       C1 | C2
    -------- *   | B1 | B2 |   =   -------
    |  A2  |     |    |    |       C3 | C4 
  
    A1 * B1 = C1
    A1 * B2 = c2
    A2 * B1 = C3
    A2 * B2 = C4
    
    These 4 computations may take place simultaneously on 4 different GPUs.
  ------------------------------------------------

***********************/

#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#include <time.h>
// CUDA runtime
//#include <cuda_runtime.h>


//Error handling using functions of the CUDA runtime API
#define cudaCheckError() {                                                              \
    cudaError_t e=cudaGetLastError();                                                   \
    if(e!=cudaSuccess) {                                                                \
        printf("Cuda failure %s:%d: '%s'\n",__FILE__,__LINE__,cudaGetErrorString(e));   \
        cudaDeviceReset();                                                              \
        exit(EXIT_FAILURE);                                                             \
    }                                                                                   \
}

//This macro checks malloc() and cudaMalloc() return values
#define Check_Allocation_Return_Value(a){   \
    if(a==NULL) {                           \
    printf("Allocation Error\n");           \
    cudaDeviceReset();                      \
    exit(EXIT_FAILURE);                     \
    }                                       \
}


//general kernel(not used)
__global__ void matrix_multiplication(double *A,double *B,double *C,int width){
    int idy = blockIdx.y*blockDim.y+threadIdx.y;
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    
    int step;
    double prod_val = 0;
    if((idy>=width)||((idx>=width))) return;
    
    for(step=0;step<width;step++){
        prod_val += A[idy*width+step] * B[step*width+idx];
    }
    
    C[idy*width+idx] = prod_val;
}

// Kernel for the computation of C1 portion
__global__ void kernelC1(double *A,double *B,double *C,int width, double r){
    int idy = blockIdx.y*blockDim.y+threadIdx.y;
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    int step;
    double prod_val = 0;
    
    if((idy>=(int)(width*r))||(idx>=(int)(width*r))) return;
    
    for(step=0;step<width;step++){
        prod_val += A[idy*width+step] * B[step*(int)(width*r)+idx];
    }
    
    C[idy*(int)(width*r)+idx] = prod_val;
}

// Kernel for the computation of C2 portion
__global__ void kernelC2(double *A,double *B,double *C,int width, double r){
    int idy = blockIdx.y*blockDim.y+threadIdx.y;
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    int step;
    double prod_val = 0;
    
    if((idy>=(int)(width*r))||(idx>=(int)(width*(1-r)))) return;
    
    for(step=0;step<width;step++){
        prod_val += A[idy*width+step] * B[step*(int)(width*(1-r))+idx];
    }
    
    C[idy*(int)(width*(1-r))+idx] = prod_val;
}


// Kernel for the computation of C3 portion
__global__ void kernelC3(double *A,double *B,double *C,int width, double r){
    int idy = blockIdx.y*blockDim.y+threadIdx.y;
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    
    int step;
    double prod_val = 0;
    if((idy>=(int)(width*(1-r)))||(idx>=(int)(width*r))) return;
    
    for(step=0;step<width;step++){
        prod_val += A[idy*width+step] * B[step*(int)(width*r)+idx];
    }
    
    
    C[idy*(int)(width*r)+idx] = prod_val;
}

// // Kernel for the computation of C4 portion
__global__ void kernelC4(double *A,double *B,double *C,int width, double r){
    int idy = blockIdx.y*blockDim.y+threadIdx.y;
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    
    int step;
    double prod_val = 0;
    if((idy>=(int)(width*(1-r)))||(idx>=(int)(width*(1-r)))) return;
    
    for(step=0;step<width;step++){
        prod_val += A[idy*width+step] * B[step*(int)(width*(1-r))+idx];
    }
    C[idy*(int)(width*(1-r))+idx] = prod_val;
}



int main(int argc,char *argv[]){
    const int num_streams = 4;
    cudaStream_t streams[num_streams];
    int N =7000;
    double *hA,*hB,*hC;
    int id,j,i;
    int ndev;
    double r = 0.5;
    double inv_r = (1-r);
    double *hA1,*hA2,*hB1,*hB2,*hC1,*hC2,*hC3,*hC4;
    double *dA1,*dA1_2,*dA2,*dA2_2,*dB1,*dB1_2,*dB2,*dB2_2;
    double *dC1,*dC2,*dC3,*dC4;
    
    printf("\nNumber of elements of the final matrix: %d\n",N * N);
    printf("Block 1 width: %d\n",(int)(N*r));
    printf("Block 2 width: %d\n",(int)(N*inv_r));
         
    cudaGetDeviceCount(&ndev);
    if(ndev==0){
        printf("NO GPU DEVICES AVAILABLE\n\n");
        exit(-1);
            
    }else{
        printf("Number of available GPUs: %d\n\n",ndev);
    }
        
    cudaMallocHost(&hA,N*N*sizeof(double));
    Check_Allocation_Return_Value(hA)
    cudaMallocHost(&hB,N*N*sizeof(double));
    Check_Allocation_Return_Value(hB)
    cudaMallocHost(&hC,N*N*sizeof(double));
    Check_Allocation_Return_Value(hC)
    memset(hC,0,N*N*sizeof(double));
        
    srand (time(NULL));
    
    for(i=0;i<N*N;i++){
        hA[i] = rand()%10;
        hB[i] = rand()%10;
    }
    
    //Grid and block size initialization
    int grid_width = 1+N/32;
    dim3 dimGrid(grid_width,grid_width,1);
    dim3 dimBlock(32,32,1);
    
        
    ////////////////////////////////////////////////////////////////////////////////
    ////////////////////////////////////////////////////////////////////////////////
    // kernel 1
    id=0;
    cudaSetDevice((int)(id%ndev));
    //cudaStreamCreate(&streams[id]);
    cudaStreamCreateWithFlags(&streams[id],cudaStreamNonBlocking);
    
    cudaMallocHost(&hA1,(int)(N*N*r*sizeof(double)));
    Check_Allocation_Return_Value(hA1)
    cudaMallocHost(&hB1,(int)(N*N*r*sizeof(double)));
    Check_Allocation_Return_Value(hB1)
    cudaMallocHost(&hC1,(int)(N*N*r*r*sizeof(double)));
    Check_Allocation_Return_Value(hC1)
    
    for(int i=0;i<(int)(N*r);i++){
        for(int j=0;j<N;j++){
            hA1[i*N+j] =  hA[i*N+j];
        }
    }
    for(int i=0;i<N;i++){
        for(int j=0;j<(N*r);j++){
            hB1[i*(int)(N*r)+j] =  hB[i*N+j];
        }
    }

    cudaMalloc((void**)&dA1,(int)(N*N*r*sizeof(double)));
    cudaCheckError()
    cudaMalloc((void**)&dB1,(int)(N*N*r*sizeof(double)));
    cudaCheckError()
    cudaMalloc((void**)&dC1,(int)(N*N*r*r*sizeof(double)));
    cudaCheckError()
        
    // kernel 2
    id=1;
    cudaSetDevice((int)(id%ndev));
    //cudaStreamCreate(&streams[id]);
    cudaStreamCreateWithFlags(&streams[id],cudaStreamNonBlocking);
    
    cudaMallocHost(&hB2,(int)(N*N*inv_r*sizeof(double)));
    Check_Allocation_Return_Value(hB2)
    cudaMallocHost(&hC2,(int)(N*N*r*inv_r*sizeof(double)));
    Check_Allocation_Return_Value(hC2)
    
    for(int i=0;i<N;i++){
        for(int j=0;j<(N*inv_r);j++){
            hB2[i*(int)(N*inv_r)+j] =  hB[i*N+(int)(N*r)+j];
        }
    }
     
    cudaMalloc((void**)&dA1_2,(int)(N*N*r*sizeof(double)));
    cudaCheckError()
    cudaMalloc((void**)&dB2,(int)(N*N*inv_r*sizeof(double)));
    cudaCheckError()
    cudaMalloc((void**)&dC2,(int)(N*N*r*inv_r*sizeof(double)));
    cudaCheckError()
        
    // kernel 3
    id=2;
    cudaSetDevice(id%ndev);
    //cudaStreamCreate(&streams[id]);
    cudaStreamCreateWithFlags(&streams[id],cudaStreamNonBlocking);
    
    cudaMallocHost(&hA2,(int)(N*N*inv_r*sizeof(double)));
    Check_Allocation_Return_Value(hA2)
    cudaMallocHost(&hC3,(int)(N*N*inv_r*r*sizeof(double)));
    Check_Allocation_Return_Value(hC3)
    
    for(int i=0;i<(int)(N*inv_r);i++){
        for(int j=0;j<N;j++){
            hA2[i*N+j] =  hA[(i+(int)(N*r))*N+j];
        }
    }
    
    cudaMalloc((void**)&dA2,(int)(N*N*inv_r*sizeof(double)));
    cudaCheckError()
    cudaMalloc((void**)&dB1_2,(int)(N*N*r*sizeof(double)));
    cudaCheckError()
    cudaMalloc((void**)&dC3,(int)(N*N*r*inv_r*sizeof(double)));
    cudaCheckError()  
        
    // kernel 4
    id=3;
    cudaSetDevice(id%ndev);
    //cudaStreamCreate(&streams[id]);
    cudaStreamCreateWithFlags(&streams[id],cudaStreamNonBlocking);

    cudaMallocHost(&hC4,(int)(N*N*inv_r*inv_r*sizeof(double)));
    Check_Allocation_Return_Value(hC4)
    
    cudaMalloc((void**)&dA2_2,(int)(N*N*inv_r*sizeof(double)));
    cudaCheckError()
    cudaMalloc((void**)&dB2_2,(int)(N*N*inv_r*sizeof(double)));
    cudaCheckError()
    cudaMalloc((void**)&dC4,(int)(N*N*inv_r*inv_r*sizeof(double)));
    cudaCheckError()
        
    //////////////////////////////////////////////////////////////////////////////////////////////
    //////////////////////////////////////////////////////////////////////////////////////////////
   
    printf("CPU-->GPU Memory copy(A1,B1,C1) - cudaMemcpyAsync\n");
    
    id=0;
    cudaSetDevice(id%ndev);
        
    cudaMemcpyAsync(dA1,hA1,(int)(N*N*r*sizeof(double)),cudaMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    cudaMemcpyAsync(dB1,hB1,(int)(N*N*r*sizeof(double)),cudaMemcpyHostToDevice,streams[id]);
    cudaCheckError()
        
    printf("Kernel 1 Execution...\n");
    kernelC1 <<< dimGrid,dimBlock,0,streams[id]>>>(dA1,dB1,dC1,N,r);
    cudaCheckError()
    
    ///////////////////////////////////////////////////////////////////////////////  
    
    id=1;
    cudaSetDevice(id%ndev);
    
    printf("CPU-->GPU Memory copy(A1,B2,C2) - cudaMemcpyAsync\n");
    cudaMemcpyAsync(dA1_2,hA1,(int)(N*N*r*sizeof(double)),cudaMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    cudaMemcpyAsync(dB2,hB2,(int)(N*N*inv_r*sizeof(double)),cudaMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    
    printf("Kernel 2 Execution...\n");
    kernelC2 <<< dimGrid,dimBlock,0,streams[id]>>>(dA1_2,dB2,dC2,N,r);
    cudaCheckError()
    
    ///////////////////////////////////////////////////////////////////////////////
    
    id=2;
    cudaSetDevice(id%ndev);
    
    printf("CPU-->GPU Memory copy(A2,B1,C3) - cudaMemcpyAsync\n");
    cudaMemcpyAsync(dA2,hA2,(int)(N*N*inv_r*sizeof(double)),cudaMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    cudaMemcpyAsync(dB1_2,hB1,(int)(N*N*r*sizeof(double)),cudaMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    
    printf("Kernel 3 Execution...\n");
    kernelC3 <<< dimGrid,dimBlock,0,streams[id]>>>(dA2,dB1_2,dC3,N,r);
    cudaCheckError()

    ///////////////////////////////////////////////////////////////////////////////
    
    id=3;
    cudaSetDevice(id%ndev);
    
    printf("CPU-->GPU Memory copy(A2,B2,C4) - cudaMemcpyAsync\n");
    cudaMemcpyAsync(dA2_2,hA2,(int)(N*N*inv_r*sizeof(double)),cudaMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    cudaMemcpyAsync(dB2_2,hB2,(int)(N*N*inv_r*sizeof(double)),cudaMemcpyHostToDevice,streams[id]);
    cudaCheckError()

    printf("Kernel 4 Execution...\n");
    kernelC4 <<< dimGrid,dimBlock,0,streams[id]>>>(dA2_2,dB2_2,dC4,N,r);
    cudaCheckError()

    
    ///////////////////////////////////////////////////////////////////////////////
    printf("GPU-->CPU Memory copy (dC1) - cudaMemcpyAsync\n");
    cudaMemcpyAsync(hC1,dC1,(int)(N*N*r*r*sizeof(double)),cudaMemcpyDeviceToHost,streams[id]);
    cudaCheckError()
    
    printf("GPU-->CPU Memory copy (dC2) - cudaMemcpyAsync\n");
    cudaMemcpyAsync(hC2,dC2,(int)(N*N*r*inv_r*sizeof(double)),cudaMemcpyDeviceToHost,streams[id]);
    cudaCheckError()
    
    printf("GPU-->CPU Memory copy (dC3) - cudaMemcpyAsync\n");
    cudaMemcpyAsync(hC3,dC3,(int)(N*N*r*inv_r*sizeof(double)),cudaMemcpyDeviceToHost,streams[id]);
    cudaCheckError()
    
    printf("GPU-->CPU Memory copy (dC4) - cudaMemcpyAsync\n");
    cudaMemcpyAsync(hC4,dC4,(int)(N*N*inv_r*inv_r*sizeof(double)),cudaMemcpyDeviceToHost,streams[id]);
    cudaCheckError()
    
    
    //Synchronize in order to process the results of every invocation

    id=0;
    cudaSetDevice(id%ndev);
    cudaStreamSynchronize(streams[id]);
    
    id=1;
    cudaSetDevice(id%ndev);
    cudaStreamSynchronize(streams[id]);

    id=2;
    cudaSetDevice(id%ndev);
    cudaStreamSynchronize(streams[id]);

    id=3;
    cudaSetDevice(id%ndev);
    cudaStreamSynchronize(streams[id]);

    //create the final Matrix
    for(i=0;i<(int)N*r;i++){
        for(j=0;j<(int)N*r;j++){
              hC[i*N+j] = hC1[i*(int)(N*r)+j];
              //printf("hC[%d]:%f ",i*N+j,hC[i*N+j]);
        }
        //printf("\n");
    }
    //printf("\n");
    
    
    for(i=0;i<(int)N*r;i++){
        for(j=0;j<(int)(N*inv_r);j++){
             hC[i*N+j+(int)(N*r)] = hC2[i*(int)(N*inv_r)+j];
             //printf("hC[%d]:%f",i*N+j+(int)(N*r),hC[i*N+j+(int)(N*r)]);
        }
        //printf("\n");
    }
    //printf("\n");
    
    for(i=0;i<(int)(N*inv_r);i++){
        for(j=0;j<(int)(N*r);j++){
             hC[(i+(int)(N*r))*N+j] = hC3[i*(int)(N*r)+j];
             //printf("hC[%d]:%f",(i+(int)(N*r))*N+j,hC[(i+(int)(N*r))*N+j]);
        }
        //printf("\n");
    }
    //printf("\n"); 
    
  
    for(i=0;i<(int)(N*inv_r);i++){
        for(j=0;j<(int)(N*inv_r);j++){
            hC[(i+(int)(N*r))*N+j+(int)(N*r)] = hC4[i*(int)(N*inv_r)+j];
          //  printf("hC[%d]:%f",(i+(int)(N*r))*N+j+(int)(N*r),hC[(i+(int)(N*r))*N+j+(int)(N*r)]);
        }
       // printf("\n");
    }
  //  printf("\n"); 
    
    
    /*
    //Compare the GPU result with CPU computation(for validation)
    printf("Check results...\n");
    int k;
    double res; 
    for(i=0;i<N;i++){
        for(j=0;j<N;j++){
            res=0;
            for(k=0;k<N;k++){
                res+=hA[i*N+k]*hB[k*N+j];
            }
            
           //printf("%8.3f ",res);
           if(res != hC[i*N+j]){
                printf("NOT OK i:%d, j:%d\n",i,j);
                printf("true value:%f - computed value:%f\n\n",res,hC[i*N+j]);
           }
        }
        //printf("\n");
    }
    */
    
    
    
    printf("Free Host and Device Memory\n");
    cudaFreeHost(hA);
    cudaFreeHost(hB);
    cudaFreeHost(hC);
    cudaFreeHost(hA1);
    cudaFreeHost(hA2);
    cudaFreeHost(hB1);
    cudaFreeHost(hB2);
    cudaFreeHost(hC1);
    cudaFreeHost(hC2);
    cudaFreeHost(hC3);
    cudaFreeHost(hC4);
    
    id=0;
    cudaSetDevice(id%ndev);
    cudaFree(dA1);
    cudaCheckError()
    cudaFree(dB1);
    cudaCheckError()
    cudaFree(dC1);
    cudaCheckError()
    
    id=1;
    cudaSetDevice(id%ndev);
    cudaFree(dA1_2);
    cudaCheckError()
    cudaFree(dB2);
    cudaCheckError()
    cudaFree(dC2);
    cudaCheckError()
    
    id=2;
    cudaSetDevice(id%ndev);
    cudaFree(dA2);
    cudaCheckError()
    cudaFree(dB1_2);
    cudaCheckError()
    cudaFree(dC3);
    cudaCheckError()
    
    id=3;
    cudaSetDevice(id%ndev);
    cudaFree(dA2_2);
    cudaCheckError()
    cudaFree(dB2_2);
    cudaCheckError()
    cudaFree(dC4);
    cudaCheckError()
    
    cudaStreamDestroy(streams[0]);
    cudaStreamDestroy(streams[1]);
    cudaStreamDestroy(streams[2]);
    cudaStreamDestroy(streams[3]);
    
    return(0);
}

	.file	"tmpxft_003b929e_00000000-6_matrixMulMultiGPU.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2043:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2043:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0, @function
_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0:
.LFB2194:
	.cfi_startproc
	xorl	%edx, %edx
	jmp	cudaHostAlloc@PLT
	.cfi_endproc
.LFE2194:
	.size	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0, .-_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	.globl	_Z46__device_stub__Z21matrix_multiplicationPdS_S_iPdS_S_i
	.type	_Z46__device_stub__Z21matrix_multiplicationPdS_S_iPdS_S_i, @function
_Z46__device_stub__Z21matrix_multiplicationPdS_S_iPdS_S_i:
.LFB2065:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L3
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z21matrix_multiplicationPdS_S_i(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L3:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L5
	call	__stack_chk_fail@PLT
.L5:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2065:
	.size	_Z46__device_stub__Z21matrix_multiplicationPdS_S_iPdS_S_i, .-_Z46__device_stub__Z21matrix_multiplicationPdS_S_iPdS_S_i
	.globl	_Z21matrix_multiplicationPdS_S_i
	.type	_Z21matrix_multiplicationPdS_S_i, @function
_Z21matrix_multiplicationPdS_S_i:
.LFB2066:
	.cfi_startproc
	endbr64
	jmp	_Z46__device_stub__Z21matrix_multiplicationPdS_S_iPdS_S_i
	.cfi_endproc
.LFE2066:
	.size	_Z21matrix_multiplicationPdS_S_i, .-_Z21matrix_multiplicationPdS_S_i
	.globl	_Z33__device_stub__Z8kernelC1PdS_S_idPdS_S_id
	.type	_Z33__device_stub__Z8kernelC1PdS_S_idPdS_S_id, @function
_Z33__device_stub__Z8kernelC1PdS_S_idPdS_S_id:
.LFB2067:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	56(%rsp), %rcx
	movsd	%xmm0, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	40(%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	20(%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	.L9
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z8kernelC1PdS_S_id(%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
.L9:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2067:
	.size	_Z33__device_stub__Z8kernelC1PdS_S_idPdS_S_id, .-_Z33__device_stub__Z8kernelC1PdS_S_idPdS_S_id
	.globl	_Z8kernelC1PdS_S_id
	.type	_Z8kernelC1PdS_S_id, @function
_Z8kernelC1PdS_S_id:
.LFB2068:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z8kernelC1PdS_S_idPdS_S_id
	.cfi_endproc
.LFE2068:
	.size	_Z8kernelC1PdS_S_id, .-_Z8kernelC1PdS_S_id
	.globl	_Z33__device_stub__Z8kernelC2PdS_S_idPdS_S_id
	.type	_Z33__device_stub__Z8kernelC2PdS_S_idPdS_S_id, @function
_Z33__device_stub__Z8kernelC2PdS_S_idPdS_S_id:
.LFB2069:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	56(%rsp), %rcx
	movsd	%xmm0, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	40(%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	20(%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	.L14
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z8kernelC2PdS_S_id(%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
.L14:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L16
	call	__stack_chk_fail@PLT
.L16:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2069:
	.size	_Z33__device_stub__Z8kernelC2PdS_S_idPdS_S_id, .-_Z33__device_stub__Z8kernelC2PdS_S_idPdS_S_id
	.globl	_Z8kernelC2PdS_S_id
	.type	_Z8kernelC2PdS_S_id, @function
_Z8kernelC2PdS_S_id:
.LFB2070:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z8kernelC2PdS_S_idPdS_S_id
	.cfi_endproc
.LFE2070:
	.size	_Z8kernelC2PdS_S_id, .-_Z8kernelC2PdS_S_id
	.globl	_Z33__device_stub__Z8kernelC3PdS_S_idPdS_S_id
	.type	_Z33__device_stub__Z8kernelC3PdS_S_idPdS_S_id, @function
_Z33__device_stub__Z8kernelC3PdS_S_idPdS_S_id:
.LFB2071:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	56(%rsp), %rcx
	movsd	%xmm0, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	40(%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	20(%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	.L19
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z8kernelC3PdS_S_id(%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
.L19:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L21
	call	__stack_chk_fail@PLT
.L21:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2071:
	.size	_Z33__device_stub__Z8kernelC3PdS_S_idPdS_S_id, .-_Z33__device_stub__Z8kernelC3PdS_S_idPdS_S_id
	.globl	_Z8kernelC3PdS_S_id
	.type	_Z8kernelC3PdS_S_id, @function
_Z8kernelC3PdS_S_id:
.LFB2072:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z8kernelC3PdS_S_idPdS_S_id
	.cfi_endproc
.LFE2072:
	.size	_Z8kernelC3PdS_S_id, .-_Z8kernelC3PdS_S_id
	.globl	_Z33__device_stub__Z8kernelC4PdS_S_idPdS_S_id
	.type	_Z33__device_stub__Z8kernelC4PdS_S_idPdS_S_id, @function
_Z33__device_stub__Z8kernelC4PdS_S_idPdS_S_id:
.LFB2073:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 40(%rsp)
	leaq	64(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	48(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	56(%rsp), %rcx
	movsd	%xmm0, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	40(%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	20(%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	.L24
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z8kernelC4PdS_S_id(%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
.L24:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L26
	call	__stack_chk_fail@PLT
.L26:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2073:
	.size	_Z33__device_stub__Z8kernelC4PdS_S_idPdS_S_id, .-_Z33__device_stub__Z8kernelC4PdS_S_idPdS_S_id
	.globl	_Z8kernelC4PdS_S_id
	.type	_Z8kernelC4PdS_S_id, @function
_Z8kernelC4PdS_S_id:
.LFB2074:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z8kernelC4PdS_S_idPdS_S_id
	.cfi_endproc
.LFE2074:
	.size	_Z8kernelC4PdS_S_id, .-_Z8kernelC4PdS_S_id
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"\nNumber of elements of the final matrix: %d\n"
.LC1:
	.string	"Block 1 width: %d\n"
.LC2:
	.string	"Block 2 width: %d\n"
.LC3:
	.string	"NO GPU DEVICES AVAILABLE\n\n"
.LC4:
	.string	"Number of available GPUs: %d\n\n"
.LC5:
	.string	"Allocation Error\n"
.LC6:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/ChristosMatzoros/CUDA-MultiGPU-Tiled-Matrix-Multiplication-using-CUDA-Streams/main/matrixMulMultiGPU.cu"
.LC7:
	.string	"Cuda failure %s:%d: '%s'\n"
.LC8:
	.string	"CPU-->GPU Memory copy(A1,B1,C1) - cudaMemcpyAsync\n"
.LC9:
	.string	"Kernel 1 Execution...\n"
.LC11:
	.string	"CPU-->GPU Memory copy(A1,B2,C2) - cudaMemcpyAsync\n"
.LC12:
	.string	"Kernel 2 Execution...\n"
.LC13:
	.string	"CPU-->GPU Memory copy(A2,B1,C3) - cudaMemcpyAsync\n"
.LC14:
	.string	"Kernel 3 Execution...\n"
.LC15:
	.string	"CPU-->GPU Memory copy(A2,B2,C4) - cudaMemcpyAsync\n"
.LC16:
	.string	"Kernel 4 Execution...\n"
.LC17:
	.string	"GPU-->CPU Memory copy (dC1) - cudaMemcpyAsync\n"
.LC18:
	.string	"GPU-->CPU Memory copy (dC2) - cudaMemcpyAsync\n"
.LC19:
	.string	"GPU-->CPU Memory copy (dC3) - cudaMemcpyAsync\n"
.LC20:
	.string	"GPU-->CPU Memory copy (dC4) - cudaMemcpyAsync\n"
.LC21:
	.string	"Free Host and Device Memory\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2040:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	$49000000, %edx
	movl	$2, %edi
	leaq	.LC0(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	subq	$264, %rsp
	.cfi_def_cfa_offset 288
	movq	%fs:40, %rax
	movq	%rax, 248(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$3500, %edx
	leaq	.LC1(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$3500, %edx
	leaq	.LC2(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	leaq	4(%rsp), %rdi
	call	cudaGetDeviceCount@PLT
	movl	4(%rsp), %edx
	testl	%edx, %edx
	jne	.L30
	movl	$2, %edi
	leaq	.LC3(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	orl	$-1, %edi
	jmp	.L118
.L30:
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	8(%rsp), %rdi
	movl	$392000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	cmpq	$0, 8(%rsp)
	jne	.L31
.L32:
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L119:
	call	cudaDeviceReset@PLT
	movl	$1, %edi
.L118:
	call	exit@PLT
.L31:
	leaq	16(%rsp), %rdi
	movl	$392000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	cmpq	$0, 16(%rsp)
	je	.L32
	leaq	24(%rsp), %rdi
	movl	$392000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	24(%rsp), %rdx
	testq	%rdx, %rdx
	je	.L32
	movq	%rdx, %rdi
	xorl	%eax, %eax
	movl	$98000000, %ecx
	xorl	%ebx, %ebx
	rep stosl
	xorl	%edi, %edi
	movl	$10, %ebp
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
.L33:
	call	rand@PLT
	cltd
	idivl	%ebp
	movq	8(%rsp), %rax
	cvtsi2sdl	%edx, %xmm0
	movsd	%xmm0, (%rax,%rbx)
	call	rand@PLT
	cltd
	idivl	%ebp
	movq	16(%rsp), %rax
	cvtsi2sdl	%edx, %xmm0
	movsd	%xmm0, (%rax,%rbx)
	addq	$8, %rbx
	cmpq	$392000000, %rbx
	jne	.L33
	movabsq	$940597838043, %rax
	xorl	%edi, %edi
	movl	$1, 200(%rsp)
	movq	%rax, 192(%rsp)
	movabsq	$137438953504, %rax
	movq	%rax, 204(%rsp)
	movl	$1, 212(%rsp)
	call	cudaSetDevice@PLT
	leaq	216(%rsp), %rdi
	movl	$1, %esi
	call	cudaStreamCreateWithFlags@PLT
	leaq	32(%rsp), %rdi
	movl	$196000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	cmpq	$0, 32(%rsp)
	je	.L32
	leaq	48(%rsp), %rdi
	movl	$196000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	cmpq	$0, 48(%rsp)
	je	.L32
	leaq	64(%rsp), %rdi
	movl	$98000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	cmpq	$0, 64(%rsp)
	je	.L32
	movq	8(%rsp), %rsi
	movq	32(%rsp), %rdi
	xorl	%edx, %edx
.L36:
	movq	%rdx, %rax
	xorl	%ecx, %ecx
.L35:
	movsd	(%rsi,%rax), %xmm0
	incl	%ecx
	movsd	%xmm0, (%rdi,%rax)
	addq	$8, %rax
	cmpl	$7000, %ecx
	jne	.L35
	addq	$56000, %rdx
	cmpq	$196000000, %rdx
	jne	.L36
	movq	16(%rsp), %rcx
	movq	48(%rsp), %rsi
	xorl	%edx, %edx
.L37:
	xorl	%eax, %eax
.L38:
	movsd	(%rcx,%rax), %xmm0
	movsd	%xmm0, (%rsi,%rax)
	addq	$8, %rax
	cmpq	$28000, %rax
	jne	.L38
	addl	$3500, %edx
	addq	$56000, %rcx
	addq	$28000, %rsi
	cmpl	$24500000, %edx
	jne	.L37
	leaq	96(%rsp), %rdi
	movl	$196000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L40
	call	cudaGetErrorString@PLT
	movl	$233, %ecx
	movq	%rax, %r8
	jmp	.L120
.L40:
	leaq	128(%rsp), %rdi
	movl	$196000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L41
	call	cudaGetErrorString@PLT
	movl	$235, %ecx
	movq	%rax, %r8
	jmp	.L120
.L41:
	leaq	160(%rsp), %rdi
	movl	$98000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L42
	call	cudaGetErrorString@PLT
	movl	$237, %ecx
	movq	%rax, %r8
	jmp	.L120
.L42:
	movl	$1, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	leaq	224(%rsp), %rdi
	movl	$1, %esi
	call	cudaStreamCreateWithFlags@PLT
	leaq	56(%rsp), %rdi
	movl	$196000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	cmpq	$0, 56(%rsp)
	je	.L32
	leaq	72(%rsp), %rdi
	movl	$98000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	cmpq	$0, 72(%rsp)
	je	.L32
	movq	16(%rsp), %rcx
	movq	56(%rsp), %rdx
	movl	$3500, %esi
.L43:
	xorl	%eax, %eax
.L44:
	movsd	28000(%rcx,%rax), %xmm0
	movsd	%xmm0, (%rdx,%rax)
	addq	$8, %rax
	cmpq	$28000, %rax
	jne	.L44
	addl	$7000, %esi
	addq	$56000, %rcx
	addq	$28000, %rdx
	cmpl	$49003500, %esi
	jne	.L43
	leaq	104(%rsp), %rdi
	movl	$196000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L46
	call	cudaGetErrorString@PLT
	movl	$257, %ecx
	movq	%rax, %r8
	jmp	.L120
.L46:
	leaq	144(%rsp), %rdi
	movl	$196000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L47
	call	cudaGetErrorString@PLT
	movl	$259, %ecx
	movq	%rax, %r8
	jmp	.L120
.L47:
	leaq	168(%rsp), %rdi
	movl	$98000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L48
	call	cudaGetErrorString@PLT
	movl	$261, %ecx
	movq	%rax, %r8
	jmp	.L120
.L48:
	movl	$2, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	leaq	232(%rsp), %rdi
	movl	$1, %esi
	call	cudaStreamCreateWithFlags@PLT
	leaq	40(%rsp), %rdi
	movl	$196000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	cmpq	$0, 40(%rsp)
	je	.L32
	leaq	80(%rsp), %rdi
	movl	$98000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	cmpq	$0, 80(%rsp)
	je	.L32
	movq	8(%rsp), %r8
	movq	40(%rsp), %rsi
	xorl	%eax, %eax
.L49:
	xorl	%edx, %edx
	leaq	(%r8,%rax), %rdi
	leaq	(%rsi,%rax), %rcx
.L50:
	movsd	196000000(%rdi,%rdx), %xmm0
	movsd	%xmm0, (%rcx,%rdx)
	addq	$8, %rdx
	cmpq	$56000, %rdx
	jne	.L50
	addq	$56000, %rax
	cmpq	$196000000, %rax
	jne	.L49
	leaq	112(%rsp), %rdi
	movl	$196000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L52
	call	cudaGetErrorString@PLT
	movl	$281, %ecx
	movq	%rax, %r8
	jmp	.L120
.L52:
	leaq	136(%rsp), %rdi
	movl	$196000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L53
	call	cudaGetErrorString@PLT
	movl	$283, %ecx
	movq	%rax, %r8
	jmp	.L120
.L53:
	leaq	176(%rsp), %rdi
	movl	$98000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L54
	call	cudaGetErrorString@PLT
	movl	$285, %ecx
	movq	%rax, %r8
	jmp	.L120
.L54:
	movl	$3, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	leaq	240(%rsp), %rdi
	movl	$1, %esi
	call	cudaStreamCreateWithFlags@PLT
	leaq	88(%rsp), %rdi
	movl	$98000000, %esi
	call	_Z14cudaMallocHostIdE9cudaErrorPPT_mj.constprop.0.isra.0
	cmpq	$0, 88(%rsp)
	je	.L32
	leaq	120(%rsp), %rdi
	movl	$196000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L55
	call	cudaGetErrorString@PLT
	movl	$297, %ecx
	movq	%rax, %r8
	jmp	.L120
.L55:
	leaq	152(%rsp), %rdi
	movl	$196000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L56
	call	cudaGetErrorString@PLT
	movl	$299, %ecx
	movq	%rax, %r8
	jmp	.L120
.L56:
	leaq	184(%rsp), %rdi
	movl	$98000000, %esi
	call	cudaMalloc@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L57
	call	cudaGetErrorString@PLT
	movl	$301, %ecx
	movq	%rax, %r8
	jmp	.L120
.L57:
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%edi, %edi
	call	cudaSetDevice@PLT
	movq	96(%rsp), %rdi
	movq	32(%rsp), %rsi
	movl	$1, %ecx
	movq	216(%rsp), %r8
	movl	$196000000, %edx
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L58
	call	cudaGetErrorString@PLT
	movl	$312, %ecx
	movq	%rax, %r8
	jmp	.L120
.L58:
	movq	48(%rsp), %rsi
	movl	$1, %ecx
	movl	$196000000, %edx
	movq	128(%rsp), %rdi
	movq	216(%rsp), %r8
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L59
	call	cudaGetErrorString@PLT
	movl	$314, %ecx
	movq	%rax, %r8
	jmp	.L120
.L59:
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	212(%rsp), %ecx
	xorl	%r8d, %r8d
	movq	216(%rsp), %r9
	movq	204(%rsp), %rdx
	movq	192(%rsp), %rdi
	movl	200(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L60
	movsd	.LC10(%rip), %xmm0
	movq	96(%rsp), %rdi
	movl	$7000, %ecx
	movq	160(%rsp), %rdx
	movq	128(%rsp), %rsi
	call	_Z33__device_stub__Z8kernelC1PdS_S_idPdS_S_id
.L60:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L61
	call	cudaGetErrorString@PLT
	movl	$318, %ecx
	movq	%rax, %r8
	jmp	.L120
.L61:
	movl	$1, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	104(%rsp), %rdi
	movq	32(%rsp), %rsi
	movl	$1, %ecx
	movq	224(%rsp), %r8
	movl	$196000000, %edx
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L62
	call	cudaGetErrorString@PLT
	movl	$327, %ecx
	movq	%rax, %r8
	jmp	.L120
.L62:
	movq	56(%rsp), %rsi
	movl	$1, %ecx
	movl	$196000000, %edx
	movq	144(%rsp), %rdi
	movq	224(%rsp), %r8
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L63
	call	cudaGetErrorString@PLT
	movl	$329, %ecx
	movq	%rax, %r8
	jmp	.L120
.L63:
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	212(%rsp), %ecx
	xorl	%r8d, %r8d
	movq	224(%rsp), %r9
	movq	204(%rsp), %rdx
	movq	192(%rsp), %rdi
	movl	200(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L64
	movsd	.LC10(%rip), %xmm0
	movq	104(%rsp), %rdi
	movl	$7000, %ecx
	movq	168(%rsp), %rdx
	movq	144(%rsp), %rsi
	call	_Z33__device_stub__Z8kernelC2PdS_S_idPdS_S_id
.L64:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L65
	call	cudaGetErrorString@PLT
	movl	$333, %ecx
	movq	%rax, %r8
	jmp	.L120
.L65:
	movl	$2, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	112(%rsp), %rdi
	movq	40(%rsp), %rsi
	movl	$1, %ecx
	movq	232(%rsp), %r8
	movl	$196000000, %edx
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L66
	call	cudaGetErrorString@PLT
	movl	$342, %ecx
	movq	%rax, %r8
	jmp	.L120
.L66:
	movq	48(%rsp), %rsi
	movl	$1, %ecx
	movl	$196000000, %edx
	movq	136(%rsp), %rdi
	movq	232(%rsp), %r8
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L67
	call	cudaGetErrorString@PLT
	movl	$344, %ecx
	movq	%rax, %r8
	jmp	.L120
.L67:
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	212(%rsp), %ecx
	xorl	%r8d, %r8d
	movq	232(%rsp), %r9
	movq	204(%rsp), %rdx
	movq	192(%rsp), %rdi
	movl	200(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L68
	movsd	.LC10(%rip), %xmm0
	movq	112(%rsp), %rdi
	movl	$7000, %ecx
	movq	176(%rsp), %rdx
	movq	136(%rsp), %rsi
	call	_Z33__device_stub__Z8kernelC3PdS_S_idPdS_S_id
.L68:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L69
	call	cudaGetErrorString@PLT
	movl	$348, %ecx
	movq	%rax, %r8
	jmp	.L120
.L69:
	movl	$3, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	120(%rsp), %rdi
	movq	40(%rsp), %rsi
	movl	$1, %ecx
	movq	240(%rsp), %r8
	movl	$196000000, %edx
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L70
	call	cudaGetErrorString@PLT
	movl	$357, %ecx
	movq	%rax, %r8
	jmp	.L120
.L70:
	movq	56(%rsp), %rsi
	movl	$1, %ecx
	movl	$196000000, %edx
	movq	152(%rsp), %rdi
	movq	240(%rsp), %r8
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L71
	call	cudaGetErrorString@PLT
	movl	$359, %ecx
	movq	%rax, %r8
	jmp	.L120
.L71:
	leaq	.LC16(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	212(%rsp), %ecx
	xorl	%r8d, %r8d
	movq	240(%rsp), %r9
	movq	204(%rsp), %rdx
	movq	192(%rsp), %rdi
	movl	200(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L72
	movsd	.LC10(%rip), %xmm0
	movq	120(%rsp), %rdi
	movl	$7000, %ecx
	movq	184(%rsp), %rdx
	movq	152(%rsp), %rsi
	call	_Z33__device_stub__Z8kernelC4PdS_S_idPdS_S_id
.L72:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L73
	call	cudaGetErrorString@PLT
	movl	$363, %ecx
	movq	%rax, %r8
	jmp	.L120
.L73:
	leaq	.LC17(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	64(%rsp), %rdi
	movl	$2, %ecx
	movq	240(%rsp), %r8
	movq	160(%rsp), %rsi
	movl	$98000000, %edx
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L74
	call	cudaGetErrorString@PLT
	movl	$369, %ecx
	movq	%rax, %r8
	jmp	.L120
.L74:
	leaq	.LC18(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	72(%rsp), %rdi
	movl	$2, %ecx
	movq	240(%rsp), %r8
	movq	168(%rsp), %rsi
	movl	$98000000, %edx
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L75
	call	cudaGetErrorString@PLT
	movl	$373, %ecx
	movq	%rax, %r8
	jmp	.L120
.L75:
	leaq	.LC19(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	80(%rsp), %rdi
	movl	$2, %ecx
	movq	240(%rsp), %r8
	movq	176(%rsp), %rsi
	movl	$98000000, %edx
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L76
	call	cudaGetErrorString@PLT
	movl	$377, %ecx
	movq	%rax, %r8
	jmp	.L120
.L76:
	leaq	.LC20(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	88(%rsp), %rdi
	movl	$2, %ecx
	movq	240(%rsp), %r8
	movq	184(%rsp), %rsi
	movl	$98000000, %edx
	call	cudaMemcpyAsync@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %ebx
	testl	%eax, %eax
	je	.L77
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	movl	$381, %ecx
	movq	%rax, %r8
	jmp	.L120
.L77:
	xorl	%edi, %edi
	call	cudaSetDevice@PLT
	movq	216(%rsp), %rdi
	call	cudaStreamSynchronize@PLT
	movl	$1, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	movq	224(%rsp), %rdi
	call	cudaStreamSynchronize@PLT
	movl	$2, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	movq	232(%rsp), %rdi
	call	cudaStreamSynchronize@PLT
	movl	$3, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	movq	240(%rsp), %rdi
	call	cudaStreamSynchronize@PLT
	movq	24(%rsp), %rax
	movq	64(%rsp), %rsi
	xorl	%edi, %edi
	movq	%rax, %rcx
.L78:
	xorl	%edx, %edx
.L79:
	movsd	(%rsi,%rdx), %xmm0
	movsd	%xmm0, (%rcx,%rdx)
	addq	$8, %rdx
	cmpq	$28000, %rdx
	jne	.L79
	addl	$3500, %edi
	addq	$28000, %rsi
	addq	$56000, %rcx
	cmpl	$12250000, %edi
	jne	.L78
	movq	72(%rsp), %rsi
	movq	%rax, %rcx
	xorl	%edi, %edi
.L80:
	xorl	%edx, %edx
.L81:
	movsd	(%rsi,%rdx), %xmm0
	movsd	%xmm0, 28000(%rcx,%rdx)
	addq	$8, %rdx
	cmpq	$28000, %rdx
	jne	.L81
	addl	$3500, %edi
	addq	$28000, %rsi
	addq	$56000, %rcx
	cmpl	$12250000, %edi
	jne	.L80
	movq	80(%rsp), %rsi
	movq	%rax, %rcx
	xorl	%edi, %edi
.L84:
	xorl	%edx, %edx
.L83:
	movsd	(%rsi,%rdx), %xmm0
	movsd	%xmm0, 196000000(%rcx,%rdx)
	addq	$8, %rdx
	cmpq	$28000, %rdx
	jne	.L83
	addl	$3500, %edi
	addq	$28000, %rsi
	addq	$56000, %rcx
	cmpl	$12250000, %edi
	jne	.L84
	movq	88(%rsp), %rcx
.L85:
	xorl	%edx, %edx
.L86:
	movsd	(%rcx,%rdx), %xmm0
	movsd	%xmm0, 196028000(%rax,%rdx)
	addq	$8, %rdx
	cmpq	$28000, %rdx
	jne	.L86
	addl	$3500, %ebx
	addq	$28000, %rcx
	addq	$56000, %rax
	cmpl	$12250000, %ebx
	jne	.L85
	leaq	.LC21(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	8(%rsp), %rdi
	call	cudaFreeHost@PLT
	movq	16(%rsp), %rdi
	call	cudaFreeHost@PLT
	movq	24(%rsp), %rdi
	call	cudaFreeHost@PLT
	movq	32(%rsp), %rdi
	call	cudaFreeHost@PLT
	movq	40(%rsp), %rdi
	call	cudaFreeHost@PLT
	movq	48(%rsp), %rdi
	call	cudaFreeHost@PLT
	movq	56(%rsp), %rdi
	call	cudaFreeHost@PLT
	movq	64(%rsp), %rdi
	call	cudaFreeHost@PLT
	movq	72(%rsp), %rdi
	call	cudaFreeHost@PLT
	movq	80(%rsp), %rdi
	call	cudaFreeHost@PLT
	movq	88(%rsp), %rdi
	call	cudaFreeHost@PLT
	xorl	%edi, %edi
	call	cudaSetDevice@PLT
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L88
	call	cudaGetErrorString@PLT
	movl	$482, %ecx
	movq	%rax, %r8
	jmp	.L120
.L88:
	movq	128(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L89
	call	cudaGetErrorString@PLT
	movl	$484, %ecx
	movq	%rax, %r8
	jmp	.L120
.L89:
	movq	160(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L90
	call	cudaGetErrorString@PLT
	movl	$486, %ecx
	movq	%rax, %r8
	jmp	.L120
.L90:
	movl	$1, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L91
	call	cudaGetErrorString@PLT
	movl	$491, %ecx
	movq	%rax, %r8
	jmp	.L120
.L91:
	movq	144(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L92
	call	cudaGetErrorString@PLT
	movl	$493, %ecx
	movq	%rax, %r8
	jmp	.L120
.L92:
	movq	168(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L93
	call	cudaGetErrorString@PLT
	movl	$495, %ecx
	movq	%rax, %r8
	jmp	.L120
.L93:
	movl	$2, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	movq	112(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L94
	call	cudaGetErrorString@PLT
	movl	$500, %ecx
	movq	%rax, %r8
	jmp	.L120
.L94:
	movq	136(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L95
	call	cudaGetErrorString@PLT
	movl	$502, %ecx
	movq	%rax, %r8
	jmp	.L120
.L95:
	movq	176(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L96
	call	cudaGetErrorString@PLT
	movl	$504, %ecx
	movq	%rax, %r8
	jmp	.L120
.L96:
	movl	$3, %eax
	cltd
	idivl	4(%rsp)
	movl	%edx, %edi
	call	cudaSetDevice@PLT
	movq	120(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L97
	call	cudaGetErrorString@PLT
	movl	$509, %ecx
	movq	%rax, %r8
	jmp	.L120
.L97:
	movq	152(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L98
	call	cudaGetErrorString@PLT
	movl	$511, %ecx
	movq	%rax, %r8
	jmp	.L120
.L98:
	movq	184(%rsp), %rdi
	call	cudaFree@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L99
	call	cudaGetErrorString@PLT
	movl	$513, %ecx
	movq	%rax, %r8
.L120:
	leaq	.LC6(%rip), %rdx
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L119
.L99:
	movq	216(%rsp), %rdi
	call	cudaStreamDestroy@PLT
	movq	224(%rsp), %rdi
	call	cudaStreamDestroy@PLT
	movq	232(%rsp), %rdi
	call	cudaStreamDestroy@PLT
	movq	240(%rsp), %rdi
	call	cudaStreamDestroy@PLT
	movq	248(%rsp), %rax
	subq	%fs:40, %rax
	je	.L100
	call	__stack_chk_fail@PLT
.L100:
	addq	$264, %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
.LFE2040:
	.size	main, .-main
	.section	.rodata.str1.1
.LC22:
	.string	"_Z8kernelC4PdS_S_id"
.LC23:
	.string	"_Z8kernelC3PdS_S_id"
.LC24:
	.string	"_Z8kernelC2PdS_S_id"
.LC25:
	.string	"_Z8kernelC1PdS_S_id"
.LC26:
	.string	"_Z21matrix_multiplicationPdS_S_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2076:
	.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	.LC22(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z8kernelC4PdS_S_id(%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	.LC23(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z8kernelC3PdS_S_id(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC24(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z8kernelC2PdS_S_id(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC25(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z8kernelC1PdS_S_id(%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
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC26(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z21matrix_multiplicationPdS_S_i(%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
.LFE2076:
	.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
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC10:
	.long	0
	.long	1071644672
	.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:

/*********************
MIT License

Copyright (c) 2020 Matzoros Christos

Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:

The above copyright notice and this permission notice shall be included in all
copies or substantial portions of the Software.

THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
SOFTWARE.

***********************/

/*********************
The purpose of this code is to execute multi-GPU matrix multiplication with multiple kernel 
invocations using streams. The program split the computation into 4 individual computations 
as it is shown below. The proportion of the size of the block is variable.
  
  ------------------------------------------------
    A * B = C   

    |  A1  |     |    |    |       C1 | C2
    -------- *   | B1 | B2 |   =   -------
    |  A2  |     |    |    |       C3 | C4 
  
    A1 * B1 = C1
    A1 * B2 = c2
    A2 * B1 = C3
    A2 * B2 = C4
    
    These 4 computations may take place simultaneously on 4 different GPUs.
  ------------------------------------------------

***********************/


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <unistd.h>
#include <time.h>
// CUDA runtime
//#include <cuda_runtime.h>


//Error handling using functions of the CUDA runtime API
#define cudaCheckError() {                                                              \
    hipError_t e=hipGetLastError();                                                   \
    if(e!=hipSuccess) {                                                                \
        printf("Cuda failure %s:%d: '%s'\n",__FILE__,__LINE__,hipGetErrorString(e));   \
        hipDeviceReset();                                                              \
        exit(EXIT_FAILURE);                                                             \
    }                                                                                   \
}

//This macro checks malloc() and cudaMalloc() return values
#define Check_Allocation_Return_Value(a){   \
    if(a==NULL) {                           \
    printf("Allocation Error\n");           \
    hipDeviceReset();                      \
    exit(EXIT_FAILURE);                     \
    }                                       \
}


//general kernel(not used)
__global__ void matrix_multiplication(double *A,double *B,double *C,int width){
    int idy = blockIdx.y*blockDim.y+threadIdx.y;
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    
    int step;
    double prod_val = 0;
    if((idy>=width)||((idx>=width))) return;
    
    for(step=0;step<width;step++){
        prod_val += A[idy*width+step] * B[step*width+idx];
    }
    
    C[idy*width+idx] = prod_val;
}

// Kernel for the computation of C1 portion
__global__ void kernelC1(double *A,double *B,double *C,int width, double r){
    int idy = blockIdx.y*blockDim.y+threadIdx.y;
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    int step;
    double prod_val = 0;
    
    if((idy>=(int)(width*r))||(idx>=(int)(width*r))) return;
    
    for(step=0;step<width;step++){
        prod_val += A[idy*width+step] * B[step*(int)(width*r)+idx];
    }
    
    C[idy*(int)(width*r)+idx] = prod_val;
}

// Kernel for the computation of C2 portion
__global__ void kernelC2(double *A,double *B,double *C,int width, double r){
    int idy = blockIdx.y*blockDim.y+threadIdx.y;
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    int step;
    double prod_val = 0;
    
    if((idy>=(int)(width*r))||(idx>=(int)(width*(1-r)))) return;
    
    for(step=0;step<width;step++){
        prod_val += A[idy*width+step] * B[step*(int)(width*(1-r))+idx];
    }
    
    C[idy*(int)(width*(1-r))+idx] = prod_val;
}


// Kernel for the computation of C3 portion
__global__ void kernelC3(double *A,double *B,double *C,int width, double r){
    int idy = blockIdx.y*blockDim.y+threadIdx.y;
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    
    int step;
    double prod_val = 0;
    if((idy>=(int)(width*(1-r)))||(idx>=(int)(width*r))) return;
    
    for(step=0;step<width;step++){
        prod_val += A[idy*width+step] * B[step*(int)(width*r)+idx];
    }
    
    
    C[idy*(int)(width*r)+idx] = prod_val;
}

// // Kernel for the computation of C4 portion
__global__ void kernelC4(double *A,double *B,double *C,int width, double r){
    int idy = blockIdx.y*blockDim.y+threadIdx.y;
    int idx = blockIdx.x*blockDim.x+threadIdx.x;
    
    int step;
    double prod_val = 0;
    if((idy>=(int)(width*(1-r)))||(idx>=(int)(width*(1-r)))) return;
    
    for(step=0;step<width;step++){
        prod_val += A[idy*width+step] * B[step*(int)(width*(1-r))+idx];
    }
    C[idy*(int)(width*(1-r))+idx] = prod_val;
}



int main(int argc,char *argv[]){
    const int num_streams = 4;
    hipStream_t streams[num_streams];
    int N =7000;
    double *hA,*hB,*hC;
    int id,j,i;
    int ndev;
    double r = 0.5;
    double inv_r = (1-r);
    double *hA1,*hA2,*hB1,*hB2,*hC1,*hC2,*hC3,*hC4;
    double *dA1,*dA1_2,*dA2,*dA2_2,*dB1,*dB1_2,*dB2,*dB2_2;
    double *dC1,*dC2,*dC3,*dC4;
    
    printf("\nNumber of elements of the final matrix: %d\n",N * N);
    printf("Block 1 width: %d\n",(int)(N*r));
    printf("Block 2 width: %d\n",(int)(N*inv_r));
         
    hipGetDeviceCount(&ndev);
    if(ndev==0){
        printf("NO GPU DEVICES AVAILABLE\n\n");
        exit(-1);
            
    }else{
        printf("Number of available GPUs: %d\n\n",ndev);
    }
        
    hipHostMalloc(&hA,N*N*sizeof(double), hipHostMallocDefault);
    Check_Allocation_Return_Value(hA)
    hipHostMalloc(&hB,N*N*sizeof(double), hipHostMallocDefault);
    Check_Allocation_Return_Value(hB)
    hipHostMalloc(&hC,N*N*sizeof(double), hipHostMallocDefault);
    Check_Allocation_Return_Value(hC)
    memset(hC,0,N*N*sizeof(double));
        
    srand (time(NULL));
    
    for(i=0;i<N*N;i++){
        hA[i] = rand()%10;
        hB[i] = rand()%10;
    }
    
    //Grid and block size initialization
    int grid_width = 1+N/32;
    dim3 dimGrid(grid_width,grid_width,1);
    dim3 dimBlock(32,32,1);
    
        
    ////////////////////////////////////////////////////////////////////////////////
    ////////////////////////////////////////////////////////////////////////////////
    // kernel 1
    id=0;
    hipSetDevice((int)(id%ndev));
    //cudaStreamCreate(&streams[id]);
    hipStreamCreateWithFlags(&streams[id],hipStreamNonBlocking);
    
    hipHostMalloc(&hA1,(int)(N*N*r*sizeof(double)), hipHostMallocDefault);
    Check_Allocation_Return_Value(hA1)
    hipHostMalloc(&hB1,(int)(N*N*r*sizeof(double)), hipHostMallocDefault);
    Check_Allocation_Return_Value(hB1)
    hipHostMalloc(&hC1,(int)(N*N*r*r*sizeof(double)), hipHostMallocDefault);
    Check_Allocation_Return_Value(hC1)
    
    for(int i=0;i<(int)(N*r);i++){
        for(int j=0;j<N;j++){
            hA1[i*N+j] =  hA[i*N+j];
        }
    }
    for(int i=0;i<N;i++){
        for(int j=0;j<(N*r);j++){
            hB1[i*(int)(N*r)+j] =  hB[i*N+j];
        }
    }

    hipMalloc((void**)&dA1,(int)(N*N*r*sizeof(double)));
    cudaCheckError()
    hipMalloc((void**)&dB1,(int)(N*N*r*sizeof(double)));
    cudaCheckError()
    hipMalloc((void**)&dC1,(int)(N*N*r*r*sizeof(double)));
    cudaCheckError()
        
    // kernel 2
    id=1;
    hipSetDevice((int)(id%ndev));
    //cudaStreamCreate(&streams[id]);
    hipStreamCreateWithFlags(&streams[id],hipStreamNonBlocking);
    
    hipHostMalloc(&hB2,(int)(N*N*inv_r*sizeof(double)), hipHostMallocDefault);
    Check_Allocation_Return_Value(hB2)
    hipHostMalloc(&hC2,(int)(N*N*r*inv_r*sizeof(double)), hipHostMallocDefault);
    Check_Allocation_Return_Value(hC2)
    
    for(int i=0;i<N;i++){
        for(int j=0;j<(N*inv_r);j++){
            hB2[i*(int)(N*inv_r)+j] =  hB[i*N+(int)(N*r)+j];
        }
    }
     
    hipMalloc((void**)&dA1_2,(int)(N*N*r*sizeof(double)));
    cudaCheckError()
    hipMalloc((void**)&dB2,(int)(N*N*inv_r*sizeof(double)));
    cudaCheckError()
    hipMalloc((void**)&dC2,(int)(N*N*r*inv_r*sizeof(double)));
    cudaCheckError()
        
    // kernel 3
    id=2;
    hipSetDevice(id%ndev);
    //cudaStreamCreate(&streams[id]);
    hipStreamCreateWithFlags(&streams[id],hipStreamNonBlocking);
    
    hipHostMalloc(&hA2,(int)(N*N*inv_r*sizeof(double)), hipHostMallocDefault);
    Check_Allocation_Return_Value(hA2)
    hipHostMalloc(&hC3,(int)(N*N*inv_r*r*sizeof(double)), hipHostMallocDefault);
    Check_Allocation_Return_Value(hC3)
    
    for(int i=0;i<(int)(N*inv_r);i++){
        for(int j=0;j<N;j++){
            hA2[i*N+j] =  hA[(i+(int)(N*r))*N+j];
        }
    }
    
    hipMalloc((void**)&dA2,(int)(N*N*inv_r*sizeof(double)));
    cudaCheckError()
    hipMalloc((void**)&dB1_2,(int)(N*N*r*sizeof(double)));
    cudaCheckError()
    hipMalloc((void**)&dC3,(int)(N*N*r*inv_r*sizeof(double)));
    cudaCheckError()  
        
    // kernel 4
    id=3;
    hipSetDevice(id%ndev);
    //cudaStreamCreate(&streams[id]);
    hipStreamCreateWithFlags(&streams[id],hipStreamNonBlocking);

    hipHostMalloc(&hC4,(int)(N*N*inv_r*inv_r*sizeof(double)), hipHostMallocDefault);
    Check_Allocation_Return_Value(hC4)
    
    hipMalloc((void**)&dA2_2,(int)(N*N*inv_r*sizeof(double)));
    cudaCheckError()
    hipMalloc((void**)&dB2_2,(int)(N*N*inv_r*sizeof(double)));
    cudaCheckError()
    hipMalloc((void**)&dC4,(int)(N*N*inv_r*inv_r*sizeof(double)));
    cudaCheckError()
        
    //////////////////////////////////////////////////////////////////////////////////////////////
    //////////////////////////////////////////////////////////////////////////////////////////////
   
    printf("CPU-->GPU Memory copy(A1,B1,C1) - cudaMemcpyAsync\n");
    
    id=0;
    hipSetDevice(id%ndev);
        
    hipMemcpyAsync(dA1,hA1,(int)(N*N*r*sizeof(double)),hipMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    hipMemcpyAsync(dB1,hB1,(int)(N*N*r*sizeof(double)),hipMemcpyHostToDevice,streams[id]);
    cudaCheckError()
        
    printf("Kernel 1 Execution...\n");
    kernelC1 <<< dimGrid,dimBlock,0,streams[id]>>>(dA1,dB1,dC1,N,r);
    cudaCheckError()
    
    ///////////////////////////////////////////////////////////////////////////////  
    
    id=1;
    hipSetDevice(id%ndev);
    
    printf("CPU-->GPU Memory copy(A1,B2,C2) - cudaMemcpyAsync\n");
    hipMemcpyAsync(dA1_2,hA1,(int)(N*N*r*sizeof(double)),hipMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    hipMemcpyAsync(dB2,hB2,(int)(N*N*inv_r*sizeof(double)),hipMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    
    printf("Kernel 2 Execution...\n");
    kernelC2 <<< dimGrid,dimBlock,0,streams[id]>>>(dA1_2,dB2,dC2,N,r);
    cudaCheckError()
    
    ///////////////////////////////////////////////////////////////////////////////
    
    id=2;
    hipSetDevice(id%ndev);
    
    printf("CPU-->GPU Memory copy(A2,B1,C3) - cudaMemcpyAsync\n");
    hipMemcpyAsync(dA2,hA2,(int)(N*N*inv_r*sizeof(double)),hipMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    hipMemcpyAsync(dB1_2,hB1,(int)(N*N*r*sizeof(double)),hipMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    
    printf("Kernel 3 Execution...\n");
    kernelC3 <<< dimGrid,dimBlock,0,streams[id]>>>(dA2,dB1_2,dC3,N,r);
    cudaCheckError()

    ///////////////////////////////////////////////////////////////////////////////
    
    id=3;
    hipSetDevice(id%ndev);
    
    printf("CPU-->GPU Memory copy(A2,B2,C4) - cudaMemcpyAsync\n");
    hipMemcpyAsync(dA2_2,hA2,(int)(N*N*inv_r*sizeof(double)),hipMemcpyHostToDevice,streams[id]);
    cudaCheckError()
    hipMemcpyAsync(dB2_2,hB2,(int)(N*N*inv_r*sizeof(double)),hipMemcpyHostToDevice,streams[id]);
    cudaCheckError()

    printf("Kernel 4 Execution...\n");
    kernelC4 <<< dimGrid,dimBlock,0,streams[id]>>>(dA2_2,dB2_2,dC4,N,r);
    cudaCheckError()

    
    ///////////////////////////////////////////////////////////////////////////////
    printf("GPU-->CPU Memory copy (dC1) - cudaMemcpyAsync\n");
    hipMemcpyAsync(hC1,dC1,(int)(N*N*r*r*sizeof(double)),hipMemcpyDeviceToHost,streams[id]);
    cudaCheckError()
    
    printf("GPU-->CPU Memory copy (dC2) - cudaMemcpyAsync\n");
    hipMemcpyAsync(hC2,dC2,(int)(N*N*r*inv_r*sizeof(double)),hipMemcpyDeviceToHost,streams[id]);
    cudaCheckError()
    
    printf("GPU-->CPU Memory copy (dC3) - cudaMemcpyAsync\n");
    hipMemcpyAsync(hC3,dC3,(int)(N*N*r*inv_r*sizeof(double)),hipMemcpyDeviceToHost,streams[id]);
    cudaCheckError()
    
    printf("GPU-->CPU Memory copy (dC4) - cudaMemcpyAsync\n");
    hipMemcpyAsync(hC4,dC4,(int)(N*N*inv_r*inv_r*sizeof(double)),hipMemcpyDeviceToHost,streams[id]);
    cudaCheckError()
    
    
    //Synchronize in order to process the results of every invocation

    id=0;
    hipSetDevice(id%ndev);
    hipStreamSynchronize(streams[id]);
    
    id=1;
    hipSetDevice(id%ndev);
    hipStreamSynchronize(streams[id]);

    id=2;
    hipSetDevice(id%ndev);
    hipStreamSynchronize(streams[id]);

    id=3;
    hipSetDevice(id%ndev);
    hipStreamSynchronize(streams[id]);

    //create the final Matrix
    for(i=0;i<(int)N*r;i++){
        for(j=0;j<(int)N*r;j++){
              hC[i*N+j] = hC1[i*(int)(N*r)+j];
              //printf("hC[%d]:%f ",i*N+j,hC[i*N+j]);
        }
        //printf("\n");
    }
    //printf("\n");
    
    
    for(i=0;i<(int)N*r;i++){
        for(j=0;j<(int)(N*inv_r);j++){
             hC[i*N+j+(int)(N*r)] = hC2[i*(int)(N*inv_r)+j];
             //printf("hC[%d]:%f",i*N+j+(int)(N*r),hC[i*N+j+(int)(N*r)]);
        }
        //printf("\n");
    }
    //printf("\n");
    
    for(i=0;i<(int)(N*inv_r);i++){
        for(j=0;j<(int)(N*r);j++){
             hC[(i+(int)(N*r))*N+j] = hC3[i*(int)(N*r)+j];
             //printf("hC[%d]:%f",(i+(int)(N*r))*N+j,hC[(i+(int)(N*r))*N+j]);
        }
        //printf("\n");
    }
    //printf("\n"); 
    
  
    for(i=0;i<(int)(N*inv_r);i++){
        for(j=0;j<(int)(N*inv_r);j++){
            hC[(i+(int)(N*r))*N+j+(int)(N*r)] = hC4[i*(int)(N*inv_r)+j];
          //  printf("hC[%d]:%f",(i+(int)(N*r))*N+j+(int)(N*r),hC[(i+(int)(N*r))*N+j+(int)(N*r)]);
        }
       // printf("\n");
    }
  //  printf("\n"); 
    
    
    /*
    //Compare the GPU result with CPU computation(for validation)
    printf("Check results...\n");
    int k;
    double res; 
    for(i=0;i<N;i++){
        for(j=0;j<N;j++){
            res=0;
            for(k=0;k<N;k++){
                res+=hA[i*N+k]*hB[k*N+j];
            }
            
           //printf("%8.3f ",res);
           if(res != hC[i*N+j]){
                printf("NOT OK i:%d, j:%d\n",i,j);
                printf("true value:%f - computed value:%f\n\n",res,hC[i*N+j]);
           }
        }
        //printf("\n");
    }
    */
    
    
    
    printf("Free Host and Device Memory\n");
    hipHostFree(hA);
    hipHostFree(hB);
    hipHostFree(hC);
    hipHostFree(hA1);
    hipHostFree(hA2);
    hipHostFree(hB1);
    hipHostFree(hB2);
    hipHostFree(hC1);
    hipHostFree(hC2);
    hipHostFree(hC3);
    hipHostFree(hC4);
    
    id=0;
    hipSetDevice(id%ndev);
    hipFree(dA1);
    cudaCheckError()
    hipFree(dB1);
    cudaCheckError()
    hipFree(dC1);
    cudaCheckError()
    
    id=1;
    hipSetDevice(id%ndev);
    hipFree(dA1_2);
    cudaCheckError()
    hipFree(dB2);
    cudaCheckError()
    hipFree(dC2);
    cudaCheckError()
    
    id=2;
    hipSetDevice(id%ndev);
    hipFree(dA2);
    cudaCheckError()
    hipFree(dB1_2);
    cudaCheckError()
    hipFree(dC3);
    cudaCheckError()
    
    id=3;
    hipSetDevice(id%ndev);
    hipFree(dA2_2);
    cudaCheckError()
    hipFree(dB2_2);
    cudaCheckError()
    hipFree(dC4);
    cudaCheckError()
    
    hipStreamDestroy(streams[0]);
    hipStreamDestroy(streams[1]);
    hipStreamDestroy(streams[2]);
    hipStreamDestroy(streams[3]);
    
    return(0);
}

	.text
	.file	"matrixMulMultiGPU.hip"
	.globl	_Z36__device_stub__matrix_multiplicationPdS_S_i # -- Begin function _Z36__device_stub__matrix_multiplicationPdS_S_i
	.type	_Z36__device_stub__matrix_multiplicationPdS_S_i,@function
_Z36__device_stub__matrix_multiplicationPdS_S_i: # @_Z36__device_stub__matrix_multiplicationPdS_S_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z21matrix_multiplicationPdS_S_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z36__device_stub__matrix_multiplicationPdS_S_i, .Lfunc_end0-_Z36__device_stub__matrix_multiplicationPdS_S_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z23__device_stub__kernelC1PdS_S_id # -- Begin function _Z23__device_stub__kernelC1PdS_S_id
	.type	_Z23__device_stub__kernelC1PdS_S_id,@function
_Z23__device_stub__kernelC1PdS_S_id:    # @_Z23__device_stub__kernelC1PdS_S_id
	.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	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	12(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	32(%rsp), %rcx
	movsd	%xmm0, (%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	$_Z8kernelC1PdS_S_id, %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_end1:
	.size	_Z23__device_stub__kernelC1PdS_S_id, .Lfunc_end1-_Z23__device_stub__kernelC1PdS_S_id
	.cfi_endproc
                                        # -- End function
	.globl	_Z23__device_stub__kernelC2PdS_S_id # -- Begin function _Z23__device_stub__kernelC2PdS_S_id
	.type	_Z23__device_stub__kernelC2PdS_S_id,@function
_Z23__device_stub__kernelC2PdS_S_id:    # @_Z23__device_stub__kernelC2PdS_S_id
	.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	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	12(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	32(%rsp), %rcx
	movsd	%xmm0, (%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	$_Z8kernelC2PdS_S_id, %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_end2:
	.size	_Z23__device_stub__kernelC2PdS_S_id, .Lfunc_end2-_Z23__device_stub__kernelC2PdS_S_id
	.cfi_endproc
                                        # -- End function
	.globl	_Z23__device_stub__kernelC3PdS_S_id # -- Begin function _Z23__device_stub__kernelC3PdS_S_id
	.type	_Z23__device_stub__kernelC3PdS_S_id,@function
_Z23__device_stub__kernelC3PdS_S_id:    # @_Z23__device_stub__kernelC3PdS_S_id
	.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	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	12(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	32(%rsp), %rcx
	movsd	%xmm0, (%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	$_Z8kernelC3PdS_S_id, %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_end3:
	.size	_Z23__device_stub__kernelC3PdS_S_id, .Lfunc_end3-_Z23__device_stub__kernelC3PdS_S_id
	.cfi_endproc
                                        # -- End function
	.globl	_Z23__device_stub__kernelC4PdS_S_id # -- Begin function _Z23__device_stub__kernelC4PdS_S_id
	.type	_Z23__device_stub__kernelC4PdS_S_id,@function
_Z23__device_stub__kernelC4PdS_S_id:    # @_Z23__device_stub__kernelC4PdS_S_id
	.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	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	12(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	32(%rsp), %rcx
	movsd	%xmm0, (%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	$_Z8kernelC4PdS_S_id, %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_end4:
	.size	_Z23__device_stub__kernelC4PdS_S_id, .Lfunc_end4-_Z23__device_stub__kernelC4PdS_S_id
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function main
.LCPI5_0:
	.quad	0x3fe0000000000000              # double 0.5
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$232, %rsp
	.cfi_def_cfa_offset 256
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movl	$.L.str, %edi
	movl	$49000000, %esi                 # imm = 0x2EBAE40
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.1, %edi
	movl	$3500, %esi                     # imm = 0xDAC
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.2, %edi
	movl	$3500, %esi                     # imm = 0xDAC
	xorl	%eax, %eax
	callq	printf
	leaq	12(%rsp), %rbx
	movq	%rbx, %rdi
	callq	hipGetDeviceCount
	movl	(%rbx), %esi
	testl	%esi, %esi
	je	.LBB5_138
# %bb.1:
	movl	$.L.str.4, %edi
	xorl	%eax, %eax
	callq	printf
	leaq	88(%rsp), %rbx
	movl	$392000000, %esi                # imm = 0x175D7200
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	cmpq	$0, (%rbx)
	je	.LBB5_2
# %bb.4:
	leaq	80(%rsp), %rbx
	movl	$392000000, %esi                # imm = 0x175D7200
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	cmpq	$0, (%rbx)
	je	.LBB5_2
# %bb.5:
	leaq	224(%rsp), %rbx
	movl	$392000000, %esi                # imm = 0x175D7200
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	movq	(%rbx), %rdi
	testq	%rdi, %rdi
	je	.LBB5_2
# %bb.6:
	xorl	%ebx, %ebx
	movl	$392000000, %edx                # imm = 0x175D7200
	xorl	%esi, %esi
	callq	memset@PLT
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
.LBB5_7:                                # =>This Inner Loop Header: Depth=1
	callq	rand
	cltq
	imulq	$1717986919, %rax, %rcx         # imm = 0x66666667
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$34, %rcx
	addl	%edx, %ecx
	addl	%ecx, %ecx
	leal	(%rcx,%rcx,4), %ecx
	subl	%ecx, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	movq	88(%rsp), %rax
	movsd	%xmm0, (%rax,%rbx,8)
	callq	rand
	cltq
	imulq	$1717986919, %rax, %rcx         # imm = 0x66666667
	movq	%rcx, %rdx
	shrq	$63, %rdx
	sarq	$34, %rcx
	addl	%edx, %ecx
	addl	%ecx, %ecx
	leal	(%rcx,%rcx,4), %ecx
	subl	%ecx, %eax
	xorps	%xmm0, %xmm0
	cvtsi2sd	%eax, %xmm0
	movq	80(%rsp), %rax
	movsd	%xmm0, (%rax,%rbx,8)
	incq	%rbx
	cmpq	$49000000, %rbx                 # imm = 0x2EBAE40
	jne	.LBB5_7
# %bb.8:
	xorl	%edi, %edi
	callq	hipSetDevice
	leaq	16(%rsp), %rdi
	movl	$1, %esi
	callq	hipStreamCreateWithFlags
	leaq	72(%rsp), %rbx
	movl	$196000000, %esi                # imm = 0xBAEB900
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	cmpq	$0, (%rbx)
	je	.LBB5_2
# %bb.9:
	leaq	56(%rsp), %rbx
	movl	$196000000, %esi                # imm = 0xBAEB900
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	cmpq	$0, (%rbx)
	je	.LBB5_2
# %bb.10:
	leaq	216(%rsp), %rbx
	movl	$98000000, %esi                 # imm = 0x5D75C80
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	cmpq	$0, (%rbx)
	je	.LBB5_2
# %bb.11:                               # %.preheader564
	movq	88(%rsp), %rax
	movq	72(%rsp), %rcx
	xorl	%edx, %edx
	movl	$56000, %esi                    # imm = 0xDAC0
.LBB5_12:                               # %.preheader563
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_13 Depth 2
	xorl	%edi, %edi
.LBB5_13:                               #   Parent Loop BB5_12 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rax,%rdi,8), %xmm0            # xmm0 = mem[0],zero
	movsd	%xmm0, (%rcx,%rdi,8)
	incq	%rdi
	cmpq	$7000, %rdi                     # imm = 0x1B58
	jne	.LBB5_13
# %bb.14:                               #   in Loop: Header=BB5_12 Depth=1
	incq	%rdx
	addq	%rsi, %rcx
	addq	%rsi, %rax
	cmpq	$3500, %rdx                     # imm = 0xDAC
	jne	.LBB5_12
# %bb.15:                               # %.preheader562
	movq	80(%rsp), %rax
	movq	56(%rsp), %rcx
	xorl	%edx, %edx
.LBB5_16:                               # %.preheader561
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_17 Depth 2
	xorl	%esi, %esi
.LBB5_17:                               #   Parent Loop BB5_16 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rax,%rsi,8), %xmm0            # xmm0 = mem[0],zero
	movsd	%xmm0, (%rcx,%rsi,8)
	incq	%rsi
	cmpq	$3500, %rsi                     # imm = 0xDAC
	jne	.LBB5_17
# %bb.18:                               #   in Loop: Header=BB5_16 Depth=1
	incq	%rdx
	addq	$28000, %rcx                    # imm = 0x6D60
	addq	$56000, %rax                    # imm = 0xDAC0
	cmpq	$7000, %rdx                     # imm = 0x1B58
	jne	.LBB5_16
# %bb.19:
	leaq	184(%rsp), %rdi
	movl	$196000000, %esi                # imm = 0xBAEB900
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_20
# %bb.22:
	leaq	152(%rsp), %rdi
	movl	$196000000, %esi                # imm = 0xBAEB900
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_23
# %bb.24:
	leaq	120(%rsp), %rdi
	movl	$98000000, %esi                 # imm = 0x5D75C80
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_25
# %bb.26:
	movl	$1, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	leaq	24(%rsp), %rdi
	movl	$1, %esi
	callq	hipStreamCreateWithFlags
	leaq	48(%rsp), %rbx
	movl	$196000000, %esi                # imm = 0xBAEB900
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	cmpq	$0, (%rbx)
	je	.LBB5_2
# %bb.27:
	leaq	208(%rsp), %rbx
	movl	$98000000, %esi                 # imm = 0x5D75C80
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	cmpq	$0, (%rbx)
	je	.LBB5_2
# %bb.28:                               # %.preheader560
	movq	48(%rsp), %rax
	movl	$28000, %ecx                    # imm = 0x6D60
	addq	80(%rsp), %rcx
	xorl	%edx, %edx
.LBB5_29:                               # %.preheader559
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_30 Depth 2
	xorl	%esi, %esi
.LBB5_30:                               #   Parent Loop BB5_29 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rcx,%rsi,8), %xmm0            # xmm0 = mem[0],zero
	movsd	%xmm0, (%rax,%rsi,8)
	incq	%rsi
	cmpq	$3500, %rsi                     # imm = 0xDAC
	jne	.LBB5_30
# %bb.31:                               #   in Loop: Header=BB5_29 Depth=1
	incq	%rdx
	addq	$28000, %rax                    # imm = 0x6D60
	addq	$56000, %rcx                    # imm = 0xDAC0
	cmpq	$7000, %rdx                     # imm = 0x1B58
	jne	.LBB5_29
# %bb.32:
	leaq	176(%rsp), %rdi
	movl	$196000000, %esi                # imm = 0xBAEB900
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_33
# %bb.34:
	leaq	136(%rsp), %rdi
	movl	$196000000, %esi                # imm = 0xBAEB900
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_35
# %bb.36:
	leaq	112(%rsp), %rdi
	movl	$98000000, %esi                 # imm = 0x5D75C80
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_37
# %bb.38:
	movl	$2, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	leaq	32(%rsp), %rdi
	movl	$1, %esi
	callq	hipStreamCreateWithFlags
	leaq	64(%rsp), %rbx
	movl	$196000000, %esi                # imm = 0xBAEB900
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	cmpq	$0, (%rbx)
	je	.LBB5_2
# %bb.39:
	leaq	200(%rsp), %rbx
	movl	$98000000, %esi                 # imm = 0x5D75C80
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	cmpq	$0, (%rbx)
	je	.LBB5_2
# %bb.40:                               # %.preheader558
	movq	64(%rsp), %rax
	movl	$196000000, %ecx                # imm = 0xBAEB900
	addq	88(%rsp), %rcx
	xorl	%edx, %edx
	movl	$56000, %esi                    # imm = 0xDAC0
.LBB5_41:                               # %.preheader557
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_42 Depth 2
	xorl	%edi, %edi
.LBB5_42:                               #   Parent Loop BB5_41 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rcx,%rdi,8), %xmm0            # xmm0 = mem[0],zero
	movsd	%xmm0, (%rax,%rdi,8)
	incq	%rdi
	cmpq	$7000, %rdi                     # imm = 0x1B58
	jne	.LBB5_42
# %bb.43:                               #   in Loop: Header=BB5_41 Depth=1
	incq	%rdx
	addq	%rsi, %rax
	addq	%rsi, %rcx
	cmpq	$3500, %rdx                     # imm = 0xDAC
	jne	.LBB5_41
# %bb.44:
	leaq	168(%rsp), %rdi
	movl	$196000000, %esi                # imm = 0xBAEB900
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_45
# %bb.46:
	leaq	144(%rsp), %rdi
	movl	$196000000, %esi                # imm = 0xBAEB900
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_47
# %bb.48:
	leaq	104(%rsp), %rdi
	movl	$98000000, %esi                 # imm = 0x5D75C80
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_49
# %bb.50:
	movl	$3, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	leaq	40(%rsp), %rdi
	movl	$1, %esi
	callq	hipStreamCreateWithFlags
	leaq	192(%rsp), %rbx
	movl	$98000000, %esi                 # imm = 0x5D75C80
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	cmpq	$0, (%rbx)
	je	.LBB5_2
# %bb.51:
	leaq	160(%rsp), %rdi
	movl	$196000000, %esi                # imm = 0xBAEB900
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_52
# %bb.53:
	leaq	128(%rsp), %rdi
	movl	$196000000, %esi                # imm = 0xBAEB900
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_54
# %bb.55:
	leaq	96(%rsp), %rdi
	movl	$98000000, %esi                 # imm = 0x5D75C80
	callq	hipMalloc
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_56
# %bb.57:
	movl	$.Lstr, %edi
	callq	puts@PLT
	xorl	%edi, %edi
	callq	hipSetDevice
	movq	184(%rsp), %rdi
	movq	72(%rsp), %rsi
	movq	16(%rsp), %r8
	movl	$196000000, %edx                # imm = 0xBAEB900
	movl	$1, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_58
# %bb.59:
	movq	152(%rsp), %rdi
	movq	56(%rsp), %rsi
	movq	16(%rsp), %r8
	movl	$196000000, %edx                # imm = 0xBAEB900
	movl	$1, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_60
# %bb.61:
	movabsq	$940597838043, %r14             # imm = 0xDB000000DB
	movabsq	$137438953504, %rbx             # imm = 0x2000000020
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movq	16(%rsp), %r9
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_63
# %bb.62:
	movq	184(%rsp), %rdi
	movq	152(%rsp), %rsi
	movq	120(%rsp), %rdx
	movsd	.LCPI5_0(%rip), %xmm0           # xmm0 = mem[0],zero
	movl	$7000, %ecx                     # imm = 0x1B58
	callq	_Z23__device_stub__kernelC1PdS_S_id
.LBB5_63:
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_64
# %bb.65:
	movl	$1, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movq	176(%rsp), %rdi
	movq	72(%rsp), %rsi
	movq	24(%rsp), %r8
	movl	$196000000, %edx                # imm = 0xBAEB900
	movl	$1, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_66
# %bb.67:
	movq	136(%rsp), %rdi
	movq	48(%rsp), %rsi
	movq	24(%rsp), %r8
	movl	$196000000, %edx                # imm = 0xBAEB900
	movl	$1, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_68
# %bb.69:
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movq	24(%rsp), %r9
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_71
# %bb.70:
	movq	176(%rsp), %rdi
	movq	136(%rsp), %rsi
	movq	112(%rsp), %rdx
	movsd	.LCPI5_0(%rip), %xmm0           # xmm0 = mem[0],zero
	movl	$7000, %ecx                     # imm = 0x1B58
	callq	_Z23__device_stub__kernelC2PdS_S_id
.LBB5_71:
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_72
# %bb.73:
	movl	$2, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	movq	168(%rsp), %rdi
	movq	64(%rsp), %rsi
	movq	32(%rsp), %r8
	movl	$196000000, %edx                # imm = 0xBAEB900
	movl	$1, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_74
# %bb.75:
	movq	144(%rsp), %rdi
	movq	56(%rsp), %rsi
	movq	32(%rsp), %r8
	movl	$196000000, %edx                # imm = 0xBAEB900
	movl	$1, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_76
# %bb.77:
	movl	$.Lstr.5, %edi
	callq	puts@PLT
	movq	32(%rsp), %r9
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_79
# %bb.78:
	movq	168(%rsp), %rdi
	movq	144(%rsp), %rsi
	movq	104(%rsp), %rdx
	movsd	.LCPI5_0(%rip), %xmm0           # xmm0 = mem[0],zero
	movl	$7000, %ecx                     # imm = 0x1B58
	callq	_Z23__device_stub__kernelC3PdS_S_id
.LBB5_79:
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_80
# %bb.81:
	movl	$3, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	movl	$.Lstr.6, %edi
	callq	puts@PLT
	movq	160(%rsp), %rdi
	movq	64(%rsp), %rsi
	movq	40(%rsp), %r8
	movl	$196000000, %edx                # imm = 0xBAEB900
	movl	$1, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_82
# %bb.83:
	movq	128(%rsp), %rdi
	movq	48(%rsp), %rsi
	movq	40(%rsp), %r8
	movl	$196000000, %edx                # imm = 0xBAEB900
	movl	$1, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_84
# %bb.85:
	movl	$.Lstr.7, %edi
	callq	puts@PLT
	movq	40(%rsp), %r9
	movq	%r14, %rdi
	movl	$1, %esi
	movq	%rbx, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_87
# %bb.86:
	movq	160(%rsp), %rdi
	movq	128(%rsp), %rsi
	movq	96(%rsp), %rdx
	movsd	.LCPI5_0(%rip), %xmm0           # xmm0 = mem[0],zero
	movl	$7000, %ecx                     # imm = 0x1B58
	callq	_Z23__device_stub__kernelC4PdS_S_id
.LBB5_87:
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_88
# %bb.89:
	movl	$.Lstr.8, %edi
	callq	puts@PLT
	movq	216(%rsp), %rdi
	movq	120(%rsp), %rsi
	movq	40(%rsp), %r8
	movl	$98000000, %edx                 # imm = 0x5D75C80
	movl	$2, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_90
# %bb.91:
	movl	$.Lstr.9, %edi
	callq	puts@PLT
	movq	208(%rsp), %rdi
	movq	112(%rsp), %rsi
	movq	40(%rsp), %r8
	movl	$98000000, %edx                 # imm = 0x5D75C80
	movl	$2, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_92
# %bb.93:
	movl	$.Lstr.10, %edi
	callq	puts@PLT
	movq	200(%rsp), %rdi
	movq	104(%rsp), %rsi
	movq	40(%rsp), %r8
	movl	$98000000, %edx                 # imm = 0x5D75C80
	movl	$2, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_94
# %bb.95:
	movl	$.Lstr.11, %edi
	callq	puts@PLT
	movq	192(%rsp), %rdi
	movq	96(%rsp), %rsi
	movq	40(%rsp), %r8
	movl	$98000000, %edx                 # imm = 0x5D75C80
	movl	$2, %ecx
	callq	hipMemcpyAsync
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_96
# %bb.97:
	xorl	%ebx, %ebx
	xorl	%edi, %edi
	callq	hipSetDevice
	movq	16(%rsp), %rdi
	callq	hipStreamSynchronize
	movl	$1, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	movq	24(%rsp), %rdi
	callq	hipStreamSynchronize
	movl	$2, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	movq	32(%rsp), %rdi
	callq	hipStreamSynchronize
	movl	$3, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	movq	40(%rsp), %rdi
	callq	hipStreamSynchronize
	movq	216(%rsp), %rcx
	movq	224(%rsp), %rax
	movq	%rax, %rdx
.LBB5_98:                               # %.preheader556
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_99 Depth 2
	xorl	%esi, %esi
.LBB5_99:                               #   Parent Loop BB5_98 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rcx,%rsi,8), %xmm0            # xmm0 = mem[0],zero
	movsd	%xmm0, (%rdx,%rsi,8)
	incq	%rsi
	cmpq	$3500, %rsi                     # imm = 0xDAC
	jne	.LBB5_99
# %bb.100:                              #   in Loop: Header=BB5_98 Depth=1
	incq	%rbx
	addq	$56000, %rdx                    # imm = 0xDAC0
	addq	$28000, %rcx                    # imm = 0x6D60
	cmpq	$3500, %rbx                     # imm = 0xDAC
	jne	.LBB5_98
# %bb.101:                              # %.preheader555
	movq	208(%rsp), %rcx
	leaq	28000(%rax), %rdx
	xorl	%esi, %esi
.LBB5_102:                              # %.preheader554
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_103 Depth 2
	xorl	%edi, %edi
.LBB5_103:                              #   Parent Loop BB5_102 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rcx,%rdi,8), %xmm0            # xmm0 = mem[0],zero
	movsd	%xmm0, (%rdx,%rdi,8)
	incq	%rdi
	cmpq	$3500, %rdi                     # imm = 0xDAC
	jne	.LBB5_103
# %bb.104:                              #   in Loop: Header=BB5_102 Depth=1
	incq	%rsi
	addq	$56000, %rdx                    # imm = 0xDAC0
	addq	$28000, %rcx                    # imm = 0x6D60
	cmpq	$3500, %rsi                     # imm = 0xDAC
	jne	.LBB5_102
# %bb.105:                              # %.preheader553
	movq	200(%rsp), %rcx
	leaq	196000000(%rax), %rdx
	xorl	%esi, %esi
.LBB5_106:                              # %.preheader552
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_107 Depth 2
	xorl	%edi, %edi
.LBB5_107:                              #   Parent Loop BB5_106 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rcx,%rdi,8), %xmm0            # xmm0 = mem[0],zero
	movsd	%xmm0, (%rdx,%rdi,8)
	incq	%rdi
	cmpq	$3500, %rdi                     # imm = 0xDAC
	jne	.LBB5_107
# %bb.108:                              #   in Loop: Header=BB5_106 Depth=1
	incq	%rsi
	addq	$28000, %rcx                    # imm = 0x6D60
	addq	$56000, %rdx                    # imm = 0xDAC0
	cmpq	$3500, %rsi                     # imm = 0xDAC
	jne	.LBB5_106
# %bb.109:                              # %.preheader551
	movq	192(%rsp), %rcx
	addq	$196028000, %rax                # imm = 0xBAF2660
	xorl	%edx, %edx
.LBB5_110:                              # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB5_111 Depth 2
	xorl	%esi, %esi
.LBB5_111:                              #   Parent Loop BB5_110 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movsd	(%rcx,%rsi,8), %xmm0            # xmm0 = mem[0],zero
	movsd	%xmm0, (%rax,%rsi,8)
	incq	%rsi
	cmpq	$3500, %rsi                     # imm = 0xDAC
	jne	.LBB5_111
# %bb.112:                              #   in Loop: Header=BB5_110 Depth=1
	incq	%rdx
	addq	$28000, %rcx                    # imm = 0x6D60
	addq	$56000, %rax                    # imm = 0xDAC0
	cmpq	$3500, %rdx                     # imm = 0xDAC
	jne	.LBB5_110
# %bb.113:
	movl	$.Lstr.12, %edi
	callq	puts@PLT
	movq	88(%rsp), %rdi
	callq	hipHostFree
	movq	80(%rsp), %rdi
	callq	hipHostFree
	movq	224(%rsp), %rdi
	callq	hipHostFree
	movq	72(%rsp), %rdi
	callq	hipHostFree
	movq	64(%rsp), %rdi
	callq	hipHostFree
	movq	56(%rsp), %rdi
	callq	hipHostFree
	movq	48(%rsp), %rdi
	callq	hipHostFree
	movq	216(%rsp), %rdi
	callq	hipHostFree
	movq	208(%rsp), %rdi
	callq	hipHostFree
	movq	200(%rsp), %rdi
	callq	hipHostFree
	movq	192(%rsp), %rdi
	callq	hipHostFree
	xorl	%edi, %edi
	callq	hipSetDevice
	movq	184(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_114
# %bb.115:
	movq	152(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_116
# %bb.117:
	movq	120(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_118
# %bb.119:
	movl	$1, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	movq	176(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_120
# %bb.121:
	movq	136(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_122
# %bb.123:
	movq	112(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_124
# %bb.125:
	movl	$2, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	movq	168(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_126
# %bb.127:
	movq	144(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_128
# %bb.129:
	movq	104(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_130
# %bb.131:
	movl	$3, %eax
	xorl	%edx, %edx
	idivl	12(%rsp)
	movl	%edx, %edi
	callq	hipSetDevice
	movq	160(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_132
# %bb.133:
	movq	128(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_134
# %bb.135:
	movq	96(%rsp), %rdi
	callq	hipFree
	callq	hipGetLastError
	testl	%eax, %eax
	jne	.LBB5_136
# %bb.137:
	movq	16(%rsp), %rdi
	callq	hipStreamDestroy
	movq	24(%rsp), %rdi
	callq	hipStreamDestroy
	movq	32(%rsp), %rdi
	callq	hipStreamDestroy
	movq	40(%rsp), %rdi
	callq	hipStreamDestroy
	xorl	%eax, %eax
	addq	$232, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.LBB5_2:
	.cfi_def_cfa_offset 256
	movl	$.Lstr.23, %edi
	callq	puts@PLT
	jmp	.LBB5_3
.LBB5_138:
	movl	$.Lstr.24, %edi
	callq	puts@PLT
	movl	$-1, %edi
	callq	exit
.LBB5_20:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$235, %edx
	jmp	.LBB5_21
.LBB5_23:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$237, %edx
	jmp	.LBB5_21
.LBB5_25:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$239, %edx
	jmp	.LBB5_21
.LBB5_33:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$259, %edx                      # imm = 0x103
	jmp	.LBB5_21
.LBB5_35:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$261, %edx                      # imm = 0x105
	jmp	.LBB5_21
.LBB5_37:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$263, %edx                      # imm = 0x107
	jmp	.LBB5_21
.LBB5_45:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$283, %edx                      # imm = 0x11B
	jmp	.LBB5_21
.LBB5_47:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$285, %edx                      # imm = 0x11D
	jmp	.LBB5_21
.LBB5_49:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$287, %edx                      # imm = 0x11F
	jmp	.LBB5_21
.LBB5_52:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$299, %edx                      # imm = 0x12B
	jmp	.LBB5_21
.LBB5_54:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$301, %edx                      # imm = 0x12D
	jmp	.LBB5_21
.LBB5_56:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$303, %edx                      # imm = 0x12F
	jmp	.LBB5_21
.LBB5_58:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$314, %edx                      # imm = 0x13A
	jmp	.LBB5_21
.LBB5_60:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$316, %edx                      # imm = 0x13C
	jmp	.LBB5_21
.LBB5_64:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$320, %edx                      # imm = 0x140
	jmp	.LBB5_21
.LBB5_66:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$329, %edx                      # imm = 0x149
	jmp	.LBB5_21
.LBB5_68:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$331, %edx                      # imm = 0x14B
	jmp	.LBB5_21
.LBB5_72:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$335, %edx                      # imm = 0x14F
	jmp	.LBB5_21
.LBB5_74:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$344, %edx                      # imm = 0x158
	jmp	.LBB5_21
.LBB5_76:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$346, %edx                      # imm = 0x15A
	jmp	.LBB5_21
.LBB5_80:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$350, %edx                      # imm = 0x15E
	jmp	.LBB5_21
.LBB5_82:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$359, %edx                      # imm = 0x167
	jmp	.LBB5_21
.LBB5_84:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$361, %edx                      # imm = 0x169
	jmp	.LBB5_21
.LBB5_88:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$365, %edx                      # imm = 0x16D
	jmp	.LBB5_21
.LBB5_90:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$371, %edx                      # imm = 0x173
	jmp	.LBB5_21
.LBB5_92:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$375, %edx                      # imm = 0x177
	jmp	.LBB5_21
.LBB5_94:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$379, %edx                      # imm = 0x17B
	jmp	.LBB5_21
.LBB5_96:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$383, %edx                      # imm = 0x17F
	jmp	.LBB5_21
.LBB5_114:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$484, %edx                      # imm = 0x1E4
	jmp	.LBB5_21
.LBB5_116:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$486, %edx                      # imm = 0x1E6
	jmp	.LBB5_21
.LBB5_118:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$488, %edx                      # imm = 0x1E8
	jmp	.LBB5_21
.LBB5_120:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$493, %edx                      # imm = 0x1ED
	jmp	.LBB5_21
.LBB5_122:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$495, %edx                      # imm = 0x1EF
	jmp	.LBB5_21
.LBB5_124:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$497, %edx                      # imm = 0x1F1
	jmp	.LBB5_21
.LBB5_126:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$502, %edx                      # imm = 0x1F6
	jmp	.LBB5_21
.LBB5_128:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$504, %edx                      # imm = 0x1F8
	jmp	.LBB5_21
.LBB5_130:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$506, %edx                      # imm = 0x1FA
	jmp	.LBB5_21
.LBB5_132:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$511, %edx                      # imm = 0x1FF
	jmp	.LBB5_21
.LBB5_134:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$513, %edx                      # imm = 0x201
	jmp	.LBB5_21
.LBB5_136:
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$515, %edx                      # imm = 0x203
.LBB5_21:
	movq	%rax, %rcx
	xorl	%eax, %eax
	callq	printf
.LBB5_3:
	callq	hipDeviceReset
	movl	$1, %edi
	callq	exit
.Lfunc_end5:
	.size	main, .Lfunc_end5-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	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z21matrix_multiplicationPdS_S_i, %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	$_Z8kernelC1PdS_S_id, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8kernelC2PdS_S_id, %esi
	movl	$.L__unnamed_3, %edx
	movl	$.L__unnamed_3, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8kernelC3PdS_S_id, %esi
	movl	$.L__unnamed_4, %edx
	movl	$.L__unnamed_4, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z8kernelC4PdS_S_id, %esi
	movl	$.L__unnamed_5, %edx
	movl	$.L__unnamed_5, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB7_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z21matrix_multiplicationPdS_S_i,@object # @_Z21matrix_multiplicationPdS_S_i
	.section	.rodata,"a",@progbits
	.globl	_Z21matrix_multiplicationPdS_S_i
	.p2align	3, 0x0
_Z21matrix_multiplicationPdS_S_i:
	.quad	_Z36__device_stub__matrix_multiplicationPdS_S_i
	.size	_Z21matrix_multiplicationPdS_S_i, 8

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

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

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

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

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"\nNumber of elements of the final matrix: %d\n"
	.size	.L.str, 45

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Block 1 width: %d\n"
	.size	.L.str.1, 19

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Block 2 width: %d\n"
	.size	.L.str.2, 19

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Number of available GPUs: %d\n\n"
	.size	.L.str.4, 31

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Cuda failure %s:%d: '%s'\n"
	.size	.L.str.6, 26

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/ChristosMatzoros/CUDA-MultiGPU-Tiled-Matrix-Multiplication-using-CUDA-Streams/main/matrixMulMultiGPU.hip"
	.size	.L.str.7, 167

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z21matrix_multiplicationPdS_S_i"
	.size	.L__unnamed_1, 33

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z8kernelC1PdS_S_id"
	.size	.L__unnamed_2, 20

	.type	.L__unnamed_3,@object           # @2
.L__unnamed_3:
	.asciz	"_Z8kernelC2PdS_S_id"
	.size	.L__unnamed_3, 20

	.type	.L__unnamed_4,@object           # @3
.L__unnamed_4:
	.asciz	"_Z8kernelC3PdS_S_id"
	.size	.L__unnamed_4, 20

	.type	.L__unnamed_5,@object           # @4
.L__unnamed_5:
	.asciz	"_Z8kernelC4PdS_S_id"
	.size	.L__unnamed_5, 20

	.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	"CPU-->GPU Memory copy(A1,B1,C1) - cudaMemcpyAsync"
	.size	.Lstr, 50

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"Kernel 1 Execution..."
	.size	.Lstr.1, 22

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"CPU-->GPU Memory copy(A1,B2,C2) - cudaMemcpyAsync"
	.size	.Lstr.2, 50

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"Kernel 2 Execution..."
	.size	.Lstr.3, 22

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"CPU-->GPU Memory copy(A2,B1,C3) - cudaMemcpyAsync"
	.size	.Lstr.4, 50

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"Kernel 3 Execution..."
	.size	.Lstr.5, 22

	.type	.Lstr.6,@object                 # @str.6
.Lstr.6:
	.asciz	"CPU-->GPU Memory copy(A2,B2,C4) - cudaMemcpyAsync"
	.size	.Lstr.6, 50

	.type	.Lstr.7,@object                 # @str.7
.Lstr.7:
	.asciz	"Kernel 4 Execution..."
	.size	.Lstr.7, 22

	.type	.Lstr.8,@object                 # @str.8
.Lstr.8:
	.asciz	"GPU-->CPU Memory copy (dC1) - cudaMemcpyAsync"
	.size	.Lstr.8, 46

	.type	.Lstr.9,@object                 # @str.9
.Lstr.9:
	.asciz	"GPU-->CPU Memory copy (dC2) - cudaMemcpyAsync"
	.size	.Lstr.9, 46

	.type	.Lstr.10,@object                # @str.10
.Lstr.10:
	.asciz	"GPU-->CPU Memory copy (dC3) - cudaMemcpyAsync"
	.size	.Lstr.10, 46

	.type	.Lstr.11,@object                # @str.11
.Lstr.11:
	.asciz	"GPU-->CPU Memory copy (dC4) - cudaMemcpyAsync"
	.size	.Lstr.11, 46

	.type	.Lstr.12,@object                # @str.12
.Lstr.12:
	.asciz	"Free Host and Device Memory"
	.size	.Lstr.12, 28

	.type	.Lstr.23,@object                # @str.23
.Lstr.23:
	.asciz	"Allocation Error"
	.size	.Lstr.23, 17

	.type	.Lstr.24,@object                # @str.24
.Lstr.24:
	.asciz	"NO GPU DEVICES AVAILABLE\n"
	.size	.Lstr.24, 26

	.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__matrix_multiplicationPdS_S_i
	.addrsig_sym _Z23__device_stub__kernelC1PdS_S_id
	.addrsig_sym _Z23__device_stub__kernelC2PdS_S_id
	.addrsig_sym _Z23__device_stub__kernelC3PdS_S_id
	.addrsig_sym _Z23__device_stub__kernelC4PdS_S_id
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z21matrix_multiplicationPdS_S_i
	.addrsig_sym _Z8kernelC1PdS_S_id
	.addrsig_sym _Z8kernelC2PdS_S_id
	.addrsig_sym _Z8kernelC3PdS_S_id
	.addrsig_sym _Z8kernelC4PdS_S_id
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z21matrix_multiplicationPdS_S_i ; -- Begin function _Z21matrix_multiplicationPdS_S_i
	.globl	_Z21matrix_multiplicationPdS_S_i
	.p2align	8
	.type	_Z21matrix_multiplicationPdS_S_i,@function
_Z21matrix_multiplicationPdS_S_i:       ; @_Z21matrix_multiplicationPdS_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x2c
	s_load_b32 s2, s[0:1], 0x18
	v_bfe_u32 v1, v0, 10, 10
	v_and_b32_e32 v4, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s4, s3, 16
	s_and_b32 s3, s3, 0xffff
	v_mad_u64_u32 v[2:3], null, s15, s4, v[1:2]
	v_mad_u64_u32 v[0:1], null, s14, s3, v[4:5]
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max_i32_e32 v1, v2, v0
	v_cmpx_gt_i32_e64 s2, v1
	s_cbranch_execz .LBB0_6
; %bb.1:                                ; %.preheader
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_mul_lo_u32 v1, v2, s2
	s_cmp_lt_i32 s2, 1
	s_cbranch_scc1 .LBB0_4
; %bb.2:                                ; %.lr.ph.preheader
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_2)
	v_ashrrev_i32_e32 v2, 31, v1
	s_mov_b32 s3, s2
	v_mov_b32_e32 v6, v0
	v_lshlrev_b64 v[4:5], 3, v[1:2]
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, 0
	s_waitcnt lgkmcnt(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
.LBB0_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v7, 31, v6
	s_add_i32 s3, s3, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u32 s3, 0
	v_lshlrev_b64 v[7:8], 3, v[6:7]
	v_add_nc_u32_e32 v6, s2, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, s6, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	global_load_b64 v[9:10], v[4:5], off
	global_load_b64 v[7:8], v[7:8], off
	v_add_co_u32 v4, vcc_lo, v4, 8
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_waitcnt vmcnt(0)
	v_fma_f64 v[2:3], v[9:10], v[7:8], v[2:3]
	s_cbranch_scc1 .LBB0_3
	s_branch .LBB0_5
.LBB0_4:
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, 0
.LBB0_5:                                ; %._crit_edge
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v1, v0
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b64 v[0:1], v[2:3], off
.LBB0_6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z21matrix_multiplicationPdS_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 14
		.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 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 11
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z21matrix_multiplicationPdS_S_i, .Lfunc_end0-_Z21matrix_multiplicationPdS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 344
; NumSgprs: 18
; NumVgprs: 11
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 11
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.protected	_Z8kernelC1PdS_S_id     ; -- Begin function _Z8kernelC1PdS_S_id
	.globl	_Z8kernelC1PdS_S_id
	.p2align	8
	.type	_Z8kernelC1PdS_S_id,@function
_Z8kernelC1PdS_S_id:                    ; @_Z8kernelC1PdS_S_id
; %bb.0:
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x18
	s_load_b64 s[4:5], s[0:1], 0x20
	s_load_b32 s3, s[0:1], 0x34
	v_and_b32_e32 v3, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	v_cvt_f64_i32_e32 v[1:2], s2
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mul_f64 v[1:2], v[1:2], s[4:5]
	s_lshr_b32 s4, s3, 16
	s_and_b32 s3, s3, 0xffff
	v_cvt_i32_f64_e32 v8, v[1:2]
	v_bfe_u32 v2, v0, 10, 10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[0:1], null, s15, s4, v[2:3]
	v_mad_u64_u32 v[1:2], null, s14, s3, v[3:4]
	s_mov_b32 s3, exec_lo
	v_max_i32_e32 v2, v0, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_lt_i32_e64 v2, v8
	s_cbranch_execz .LBB1_6
; %bb.1:                                ; %.preheader
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_cmp_lt_i32 s2, 1
	s_cbranch_scc1 .LBB1_4
; %bb.2:                                ; %.lr.ph
	v_mul_lo_u32 v2, v0, s2
	v_mov_b32_e32 v6, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[4:5], 3, v[2:3]
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, 0
	s_waitcnt lgkmcnt(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_3:                                ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v7, 31, v6
	s_add_i32 s2, s2, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u32 s2, 0
	v_lshlrev_b64 v[9:10], 3, v[6:7]
	v_add_nc_u32_e32 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v9, vcc_lo, s6, v9
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v10, vcc_lo
	global_load_b64 v[11:12], v[4:5], off
	global_load_b64 v[9:10], v[9:10], off
	v_add_co_u32 v4, vcc_lo, v4, 8
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_waitcnt vmcnt(0)
	v_fma_f64 v[2:3], v[11:12], v[9:10], v[2:3]
	s_cbranch_scc1 .LBB1_3
	s_branch .LBB1_5
.LBB1_4:
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, 0
.LBB1_5:                                ; %._crit_edge
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[4:5], null, v0, v8, v[1:2]
	v_ashrrev_i32_e32 v5, 31, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[4:5]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b64 v[0:1], v[2:3], off
.LBB1_6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8kernelC1PdS_S_id
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 14
		.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 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 13
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z8kernelC1PdS_S_id, .Lfunc_end1-_Z8kernelC1PdS_S_id
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 376
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.protected	_Z8kernelC2PdS_S_id     ; -- Begin function _Z8kernelC2PdS_S_id
	.globl	_Z8kernelC2PdS_S_id
	.p2align	8
	.type	_Z8kernelC2PdS_S_id,@function
_Z8kernelC2PdS_S_id:                    ; @_Z8kernelC2PdS_S_id
; %bb.0:
	s_clause 0x2
	s_load_b32 s8, s[0:1], 0x18
	s_load_b64 s[2:3], s[0:1], 0x20
	s_load_b32 s4, s[0:1], 0x34
	v_bfe_u32 v5, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	v_cvt_f64_i32_e32 v[3:4], s8
	s_lshr_b32 s5, s4, 16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[1:2], v[3:4], s[2:3]
	v_cvt_i32_f64_e32 v6, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s5, v[5:6]
	s_mov_b32 s5, exec_lo
	v_cmpx_lt_i32_e64 v1, v6
	s_cbranch_execz .LBB2_7
; %bb.1:
	v_add_f64 v[5:6], -s[2:3], 1.0
	v_and_b32_e32 v0, 0x3ff, v0
	s_and_b32 s2, s4, 0xffff
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[2:3], v[5:6], v[3:4]
	v_cvt_i32_f64_e32 v9, v[2:3]
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[2:3], null, s14, s2, v[0:1]
	v_cmp_lt_i32_e32 vcc_lo, v2, v9
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB2_7
; %bb.2:                                ; %.preheader
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_cmp_lt_i32 s8, 1
	s_cbranch_scc1 .LBB2_5
; %bb.3:                                ; %.lr.ph
	v_mul_lo_u32 v3, v1, s8
	v_mov_b32_e32 v7, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v4, 31, v3
	v_lshlrev_b64 v[5:6], 3, v[3:4]
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s4, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v6, vcc_lo
.LBB2_4:                                ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v8, 31, v7
	s_add_i32 s8, s8, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u32 s8, 0
	v_lshlrev_b64 v[10:11], 3, v[7:8]
	v_add_nc_u32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v10, vcc_lo, s6, v10
	v_add_co_ci_u32_e32 v11, vcc_lo, s7, v11, vcc_lo
	global_load_b64 v[12:13], v[5:6], off
	global_load_b64 v[10:11], v[10:11], off
	v_add_co_u32 v5, vcc_lo, v5, 8
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v6, vcc_lo
	s_waitcnt vmcnt(0)
	v_fma_f64 v[3:4], v[12:13], v[10:11], v[3:4]
	s_cbranch_scc1 .LBB2_4
	s_branch .LBB2_6
.LBB2_5:
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
.LBB2_6:                                ; %._crit_edge
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[5:6], null, v1, v9, v[2:3]
	v_ashrrev_i32_e32 v6, 31, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[5:6]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b64 v[0:1], v[3:4], off
.LBB2_7:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8kernelC2PdS_S_id
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 14
		.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 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 14
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end2:
	.size	_Z8kernelC2PdS_S_id, .Lfunc_end2-_Z8kernelC2PdS_S_id
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 408
; NumSgprs: 18
; NumVgprs: 14
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 14
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.protected	_Z8kernelC3PdS_S_id     ; -- Begin function _Z8kernelC3PdS_S_id
	.globl	_Z8kernelC3PdS_S_id
	.p2align	8
	.type	_Z8kernelC3PdS_S_id,@function
_Z8kernelC3PdS_S_id:                    ; @_Z8kernelC3PdS_S_id
; %bb.0:
	s_clause 0x2
	s_load_b32 s8, s[0:1], 0x18
	s_load_b64 s[2:3], s[0:1], 0x20
	s_load_b32 s4, s[0:1], 0x34
	v_bfe_u32 v5, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	v_cvt_f64_i32_e32 v[3:4], s8
	v_add_f64 v[1:2], -s[2:3], 1.0
	s_lshr_b32 s5, s4, 16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mul_f64 v[1:2], v[1:2], v[3:4]
	v_cvt_i32_f64_e32 v6, v[1:2]
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s5, v[5:6]
	s_mov_b32 s5, exec_lo
	v_cmpx_lt_i32_e64 v1, v6
	s_cbranch_execz .LBB3_7
; %bb.1:
	v_mul_f64 v[2:3], v[3:4], s[2:3]
	v_and_b32_e32 v0, 0x3ff, v0
	s_and_b32 s2, s4, 0xffff
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cvt_i32_f64_e32 v9, v[2:3]
	v_mad_u64_u32 v[2:3], null, s14, s2, v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_lt_i32_e32 vcc_lo, v2, v9
	s_and_b32 exec_lo, exec_lo, vcc_lo
	s_cbranch_execz .LBB3_7
; %bb.2:                                ; %.preheader
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_cmp_lt_i32 s8, 1
	s_cbranch_scc1 .LBB3_5
; %bb.3:                                ; %.lr.ph
	v_mul_lo_u32 v3, v1, s8
	v_mov_b32_e32 v7, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v4, 31, v3
	v_lshlrev_b64 v[5:6], 3, v[3:4]
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_4)
	v_add_co_u32 v5, vcc_lo, s4, v5
	v_add_co_ci_u32_e32 v6, vcc_lo, s5, v6, vcc_lo
.LBB3_4:                                ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v8, 31, v7
	s_add_i32 s8, s8, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u32 s8, 0
	v_lshlrev_b64 v[10:11], 3, v[7:8]
	v_add_nc_u32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v10, vcc_lo, s6, v10
	v_add_co_ci_u32_e32 v11, vcc_lo, s7, v11, vcc_lo
	global_load_b64 v[12:13], v[5:6], off
	global_load_b64 v[10:11], v[10:11], off
	v_add_co_u32 v5, vcc_lo, v5, 8
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v6, vcc_lo
	s_waitcnt vmcnt(0)
	v_fma_f64 v[3:4], v[12:13], v[10:11], v[3:4]
	s_cbranch_scc1 .LBB3_4
	s_branch .LBB3_6
.LBB3_5:
	v_mov_b32_e32 v3, 0
	v_mov_b32_e32 v4, 0
.LBB3_6:                                ; %._crit_edge
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[5:6], null, v1, v9, v[2:3]
	v_ashrrev_i32_e32 v6, 31, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[5:6]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b64 v[0:1], v[3:4], off
.LBB3_7:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8kernelC3PdS_S_id
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 14
		.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 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 14
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end3:
	.size	_Z8kernelC3PdS_S_id, .Lfunc_end3-_Z8kernelC3PdS_S_id
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 408
; NumSgprs: 18
; NumVgprs: 14
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 14
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.protected	_Z8kernelC4PdS_S_id     ; -- Begin function _Z8kernelC4PdS_S_id
	.globl	_Z8kernelC4PdS_S_id
	.p2align	8
	.type	_Z8kernelC4PdS_S_id,@function
_Z8kernelC4PdS_S_id:                    ; @_Z8kernelC4PdS_S_id
; %bb.0:
	s_clause 0x2
	s_load_b32 s2, s[0:1], 0x18
	s_load_b64 s[4:5], s[0:1], 0x20
	s_load_b32 s3, s[0:1], 0x34
	s_waitcnt lgkmcnt(0)
	v_cvt_f64_i32_e32 v[1:2], s2
	v_add_f64 v[3:4], -s[4:5], 1.0
	s_lshr_b32 s4, s3, 16
	s_and_b32 s3, s3, 0xffff
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mul_f64 v[1:2], v[3:4], v[1:2]
	v_and_b32_e32 v3, 0x3ff, v0
	v_cvt_i32_f64_e32 v8, v[1:2]
	v_bfe_u32 v2, v0, 10, 10
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[0:1], null, s15, s4, v[2:3]
	v_mad_u64_u32 v[1:2], null, s14, s3, v[3:4]
	s_mov_b32 s3, exec_lo
	v_max_i32_e32 v2, v0, v1
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_lt_i32_e64 v2, v8
	s_cbranch_execz .LBB4_6
; %bb.1:                                ; %.preheader
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	s_cmp_lt_i32 s2, 1
	s_cbranch_scc1 .LBB4_4
; %bb.2:                                ; %.lr.ph
	v_mul_lo_u32 v2, v0, s2
	v_mov_b32_e32 v6, v1
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[4:5], 3, v[2:3]
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, 0
	s_waitcnt lgkmcnt(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
.LBB4_3:                                ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v7, 31, v6
	s_add_i32 s2, s2, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u32 s2, 0
	v_lshlrev_b64 v[9:10], 3, v[6:7]
	v_add_nc_u32_e32 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v9, vcc_lo, s6, v9
	v_add_co_ci_u32_e32 v10, vcc_lo, s7, v10, vcc_lo
	global_load_b64 v[11:12], v[4:5], off
	global_load_b64 v[9:10], v[9:10], off
	v_add_co_u32 v4, vcc_lo, v4, 8
	v_add_co_ci_u32_e32 v5, vcc_lo, 0, v5, vcc_lo
	s_waitcnt vmcnt(0)
	v_fma_f64 v[2:3], v[11:12], v[9:10], v[2:3]
	s_cbranch_scc1 .LBB4_3
	s_branch .LBB4_5
.LBB4_4:
	v_mov_b32_e32 v2, 0
	v_mov_b32_e32 v3, 0
.LBB4_5:                                ; %._crit_edge
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[4:5], null, v0, v8, v[1:2]
	v_ashrrev_i32_e32 v5, 31, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 3, v[4:5]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b64 v[0:1], v[2:3], off
.LBB4_6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z8kernelC4PdS_S_id
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 296
		.amdhsa_user_sgpr_count 14
		.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 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 13
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end4:
	.size	_Z8kernelC4PdS_S_id, .Lfunc_end4-_Z8kernelC4PdS_S_id
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 384
; NumSgprs: 18
; NumVgprs: 13
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 13
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 1
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z21matrix_multiplicationPdS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z21matrix_multiplicationPdS_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     11
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           8
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8kernelC1PdS_S_id
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z8kernelC1PdS_S_id.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           8
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8kernelC2PdS_S_id
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z8kernelC2PdS_S_id.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     14
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           8
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8kernelC3PdS_S_id
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z8kernelC3PdS_S_id.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     14
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           8
        .value_kind:     by_value
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         44
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         52
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         54
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         56
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         58
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         60
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         62
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         104
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 296
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z8kernelC4PdS_S_id
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z8kernelC4PdS_S_id.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     13
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
