#include <iostream>
#include <vector>
#include <math.h>
#include <assert.h>
#include <memory>
#include <random>
#include "gMat.cuh"
#include "gpuerrchk.cuh"
__global__ void matMulKernel(float* A, float* B, float* P, int m, int  n, int s, int tile_size){
	//each thread in the block will be responsible for a different element of these arrays
	extern __shared__ float tilemem[];
	float* Atile=tilemem;
	float* Btile=&tilemem[tile_size*tile_size];

	int b1=blockIdx.x;
	int b2=blockIdx.y;
	int t1=threadIdx.x;
	int t2=threadIdx.y;

	int row=b1*tile_size+t1;// the indices of the element of P this thread is responsible for computing
	int col=b2*tile_size+t2;

	float Pvalue=0;
	//loop over the tiles necessary to compute the element of P of interest
	for (int j=0; j<ceil( n / (float) tile_size);j++){
		//collaborate with other threads to store the current tiles of A and B
		//this thread is responsible for storing  the (t_1,t_2) element of the (b_1,j) tile of A and the (j,b_2) tile of B
		if( (b1*tile_size+t1) < m && (j*tile_size+t2)<n )
			Atile[tile_size*t1 + t2]=A[ n*(b1*tile_size+t1) + j*tile_size  + t2 ]; // We access A[b1*tile_size+t1][j*tile_size+t2] using linearized indices (A has n columns)
		if( (j*tile_size+t1)<n && (b2*tile_size+t2)<s  )
			Btile[tile_size*t1 + t2]=B[ s*(j*tile_size+t1)  + b2*tile_size + t2 ]; //we access the B[j*tile_size+t1][b2*tile_size+t2] element using linearized indices (B has s columns)
		__syncthreads();

		//sum the elements of the current A and B tiles used in the computation of the element of P for which this thread is responsible
		for (int k=0; k<tile_size; k++)
			Pvalue += Atile[tile_size*t1 + k]*Btile[tile_size*k + t2];

		__syncthreads();

	}
	if(row<m && col<s) P[row*s + col]=Pvalue;

}

gMat::gMat(std::vector<float> datavector, int r, int c,int devnumber): rows{r}, cols{c} {
	assert(datavector.size() == rows*cols );
	int size=sizeof(float)*rows*cols;
	
	cudaDeviceProp dev_prop;	
	cudaGetDeviceProperties(&dev_prop,devnumber);
	assert(dev_prop.totalGlobalMem>size);

	gpuErrchk(cudaMalloc( (void**)&d_data, size));
	gpuErrchk(cudaMemcpy(d_data, datavector.data(), size, cudaMemcpyHostToDevice));
}

void gMat::cleanup()
{
	gpuErrchk(cudaFree(d_data));
}


//gMat::~gMat(){ 
	//	std::cout <<"de-allocating " << name << '\n' <<std::flush;	
	//	gpuErrchk(cudaFree(d_data));
//}
int gMat::getrows(){ return rows;}
int gMat::getcols(){ return cols;}

std::vector<float> gMat::tovector(){
	float* h_data=new float[rows*cols];
	int size=rows*cols*sizeof(float);
	cudaMemcpy(h_data,d_data,size,cudaMemcpyDeviceToHost);
	std::vector<float> v;
	for (int i=0; i<rows; ++i) for (int j=0; j<cols; ++j){
		v.push_back(h_data[i*cols+j]);
	}
	delete[] h_data;
	return v;
}	

gMat randgMat(int r, int c){
	std::vector<float> datavector;
	std::default_random_engine re;
	std::uniform_real_distribution<float> dist{0.0,1.0};
	for (int i=0; i<r*c; i++) datavector.push_back(dist(re));
	return gMat(datavector,r,c);
}

gMat eye(int r){
	std::vector<float> datavector;
	for (int i=0; i!=r; ++i) for(int j=0; j!=r; ++j) 
		if(i==j) 
			datavector.push_back(1.0);
		else
			datavector.push_back(0.0);
	return gMat(datavector,r,r);
}

void prod(const gMat& A, const gMat& B, gMat& P, int mem){
	assert(A.cols==B.rows && A.rows==P.rows && B.cols==P.cols);
	int tile_size= floor(sqrt( mem/2/sizeof(float)));
    assert(tile_size>0);	
	
	matMulKernel<<<dim3(ceil(A.rows/ (float) tile_size), ceil( B.cols/ (float) tile_size)), dim3(tile_size,tile_size), mem>>>(A.d_data, B.d_data, P.d_data,A.rows, A.cols, B.cols, tile_size );
	gpuErrchk( cudaPeekAtLastError() );	
}

std::ostream& operator<<(std::ostream& os,  gMat& gm)  
{  
	std::vector<float> v=gm.tovector();
	for (int i=0; i<gm.rows; ++i){
		for (int j=0; j<gm.cols; ++j){
			os << v[i*gm.cols+j] << ",";
		}
		os << "\n";
	}
	return os;  

}  



	.file	"tmpxft_0024a9f4_00000000-6_gMat.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB5328:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE5328:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSt12_Vector_baseIfSaIfEE11_M_allocateEm.isra.0, @function
_ZNSt12_Vector_baseIfSaIfEE11_M_allocateEm.isra.0:
.LFB6193:
	.cfi_startproc
	testq	%rdi, %rdi
	je	.L3
	movabsq	$2305843009213693951, %rax
	cmpq	%rdi, %rax
	jnb	.L4
	pushq	%rax
	.cfi_def_cfa_offset 16
	movabsq	$4611686018427387903, %rax
	cmpq	%rdi, %rax
	jnb	.L5
	call	_ZSt28__throw_bad_array_new_lengthv@PLT
.L5:
	call	_ZSt17__throw_bad_allocv@PLT
.L4:
	.cfi_def_cfa_offset 8
	salq	$2, %rdi
	jmp	_Znwm@PLT
.L3:
	xorl	%eax, %eax
	ret
	.cfi_endproc
.LFE6193:
	.size	_ZNSt12_Vector_baseIfSaIfEE11_M_allocateEm.isra.0, .-_ZNSt12_Vector_baseIfSaIfEE11_M_allocateEm.isra.0
	.section	.rodata._Z9gpuAssert9cudaErrorPKcib.str1.1,"aMS",@progbits,1
.LC0:
	.string	"GPU assert: %s %s %d\n"
	.section	.text._Z9gpuAssert9cudaErrorPKcib,"axG",@progbits,_Z9gpuAssert9cudaErrorPKcib,comdat
	.weak	_Z9gpuAssert9cudaErrorPKcib
	.type	_Z9gpuAssert9cudaErrorPKcib, @function
_Z9gpuAssert9cudaErrorPKcib:
.LFB5294:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L16
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%ecx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movl	%edx, 12(%rsp)
	movq	%rsi, (%rsp)
	call	cudaGetErrorString@PLT
	movl	12(%rsp), %r9d
	movq	(%rsp), %r8
	leaq	.LC0(%rip), %rdx
	movq	stderr(%rip), %rdi
	movq	%rax, %rcx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	testb	%bpl, %bpl
	je	.L10
	movl	%ebx, %edi
	call	exit@PLT
.L10:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
.L16:
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE5294:
	.size	_Z9gpuAssert9cudaErrorPKcib, .-_Z9gpuAssert9cudaErrorPKcib
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/paiforsyth/cuda_experiments/master/ch4/mat/gMat.cu"
	.text
	.align 2
	.globl	_ZN4gMatC2ESt6vectorIfSaIfEEiii
	.type	_ZN4gMatC2ESt6vectorIfSaIfEEiii, @function
_ZN4gMatC2ESt6vectorIfSaIfEEiii:
.LFB5297:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA5297
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rsi, %r12
	movl	%r8d, %esi
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	subq	$1048, %rsp
	.cfi_def_cfa_offset 1088
	movq	%fs:40, %rax
	movq	%rax, 1032(%rsp)
	xorl	%eax, %eax
	leaq	32(%rdi), %rax
	movl	%edx, 8(%rdi)
	movq	%rax, 16(%rdi)
	xorl	%eax, %eax
	imull	%ecx, %edx
	movl	%ecx, 12(%rdi)
	movq	%rax, 24(%rdi)
	movb	$0, 32(%rdi)
	movq	%rsp, %rdi
	leal	0(,%rdx,4), %ebp
.LEHB0:
	call	cudaGetDeviceProperties_v2@PLT
	movslq	%ebp, %rbp
	movq	%rbx, %rdi
	movq	%rbp, %rsi
	call	cudaMalloc@PLT
	leaq	.LC1(%rip), %r13
	movl	%eax, %edi
	movl	$1, %ecx
	movl	$53, %edx
	movq	%r13, %rsi
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	(%r12), %rsi
	movq	(%rbx), %rdi
	movl	$1, %ecx
	movq	%rbp, %rdx
	call	cudaMemcpy@PLT
	movl	%eax, %edi
	movl	$1, %ecx
	movl	$54, %edx
	movq	%r13, %rsi
	call	_Z9gpuAssert9cudaErrorPKcib
.LEHE0:
	jmp	.L27
.L23:
	endbr64
	movq	%rax, %rbp
.L20:
	leaq	16(%rbx), %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	1032(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L26
	movq	%rbp, %rdi
.LEHB1:
	call	_Unwind_Resume@PLT
.LEHE1:
.L27:
	movq	1032(%rsp), %rax
	subq	%fs:40, %rax
	je	.L22
.L26:
	call	__stack_chk_fail@PLT
.L22:
	addq	$1048, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5297:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA5297:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE5297-.LLSDACSB5297
.LLSDACSB5297:
	.uleb128 .LEHB0-.LFB5297
	.uleb128 .LEHE0-.LEHB0
	.uleb128 .L23-.LFB5297
	.uleb128 0
	.uleb128 .LEHB1-.LFB5297
	.uleb128 .LEHE1-.LEHB1
	.uleb128 0
	.uleb128 0
.LLSDACSE5297:
	.text
	.size	_ZN4gMatC2ESt6vectorIfSaIfEEiii, .-_ZN4gMatC2ESt6vectorIfSaIfEEiii
	.globl	_ZN4gMatC1ESt6vectorIfSaIfEEiii
	.set	_ZN4gMatC1ESt6vectorIfSaIfEEiii,_ZN4gMatC2ESt6vectorIfSaIfEEiii
	.align 2
	.globl	_ZN4gMat7cleanupEv
	.type	_ZN4gMat7cleanupEv, @function
_ZN4gMat7cleanupEv:
.LFB5299:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	(%rdi), %rdi
	call	cudaFree@PLT
	movl	$1, %ecx
	movl	$59, %edx
	popq	%r8
	.cfi_def_cfa_offset 8
	movl	%eax, %edi
	leaq	.LC1(%rip), %rsi
	jmp	_Z9gpuAssert9cudaErrorPKcib
	.cfi_endproc
.LFE5299:
	.size	_ZN4gMat7cleanupEv, .-_ZN4gMat7cleanupEv
	.align 2
	.globl	_ZN4gMat7getrowsEv
	.type	_ZN4gMat7getrowsEv, @function
_ZN4gMat7getrowsEv:
.LFB5300:
	.cfi_startproc
	endbr64
	movl	8(%rdi), %eax
	ret
	.cfi_endproc
.LFE5300:
	.size	_ZN4gMat7getrowsEv, .-_ZN4gMat7getrowsEv
	.align 2
	.globl	_ZN4gMat7getcolsEv
	.type	_ZN4gMat7getcolsEv, @function
_ZN4gMat7getcolsEv:
.LFB5301:
	.cfi_startproc
	endbr64
	movl	12(%rdi), %eax
	ret
	.cfi_endproc
.LFE5301:
	.size	_ZN4gMat7getcolsEv, .-_ZN4gMat7getcolsEv
	.globl	_Z40__device_stub__Z12matMulKernelPfS_S_iiiiPfS_S_iiii
	.type	_Z40__device_stub__Z12matMulKernelPfS_S_iiiiPfS_S_iiii, @function
_Z40__device_stub__Z12matMulKernelPfS_S_iiiiPfS_S_iiii:
.LFB5350:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	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
	movl	%r8d, 16(%rsp)
	movl	%r9d, 12(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%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	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L32
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z12matMulKernelPfS_S_iiii(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	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 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L32:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L34
	call	__stack_chk_fail@PLT
.L34:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5350:
	.size	_Z40__device_stub__Z12matMulKernelPfS_S_iiiiPfS_S_iiii, .-_Z40__device_stub__Z12matMulKernelPfS_S_iiiiPfS_S_iiii
	.globl	_Z12matMulKernelPfS_S_iiii
	.type	_Z12matMulKernelPfS_S_iiii, @function
_Z12matMulKernelPfS_S_iiii:
.LFB5351:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z12matMulKernelPfS_S_iiiiPfS_S_iiii
	.cfi_endproc
.LFE5351:
	.size	_Z12matMulKernelPfS_S_iiii, .-_Z12matMulKernelPfS_S_iiii
	.globl	_Z4prodRK4gMatS1_RS_i
	.type	_Z4prodRK4gMatS1_RS_i, @function
_Z4prodRK4gMatS1_RS_i:
.LFB5324:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movl	%ecx, %r12d
	movl	$2, %ecx
	movl	%r12d, %eax
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rsi, %rbp
	cltd
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
	idivl	%ecx
	subq	$48, %rsp
	.cfi_def_cfa_offset 96
	cltq
	shrq	$2, %rax
	cvtsi2sdq	%rax, %xmm0
	call	sqrt@PLT
	call	trunc@PLT
	cvttsd2sil	%xmm0, %r13d
	cvtsi2ssl	12(%rbp), %xmm0
	cvtsi2ssl	%r13d, %xmm2
	movl	%r13d, 36(%rsp)
	movl	%r13d, 40(%rsp)
	divss	%xmm2, %xmm0
	movss	%xmm2, 12(%rsp)
	call	ceilf@PLT
	movss	12(%rsp), %xmm2
	movss	%xmm0, 8(%rsp)
	cvtsi2ssl	8(%rbx), %xmm0
	divss	%xmm2, %xmm0
	call	ceilf@PLT
	movq	36(%rsp), %rdx
	xorl	%r9d, %r9d
	movslq	%r12d, %r8
	cvttss2siq	%xmm0, %rax
	movss	8(%rsp), %xmm1
	movl	$1, %ecx
	movl	$1, %esi
	movl	%eax, 24(%rsp)
	cvttss2siq	%xmm1, %rax
	movl	%eax, 28(%rsp)
	movq	24(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L38
	movl	8(%rbx), %ecx
	pushq	%rax
	.cfi_def_cfa_offset 104
	pushq	%r13
	.cfi_def_cfa_offset 112
	movq	(%r14), %rdx
	movl	12(%rbp), %r9d
	movl	12(%rbx), %r8d
	movq	0(%rbp), %rsi
	movq	(%rbx), %rdi
	call	_Z40__device_stub__Z12matMulKernelPfS_S_iiiiPfS_S_iiii
	popq	%rdx
	.cfi_def_cfa_offset 104
	popq	%rcx
	.cfi_def_cfa_offset 96
.L38:
	call	cudaPeekAtLastError@PLT
	addq	$48, %rsp
	.cfi_def_cfa_offset 48
	movl	$1, %ecx
	movl	$106, %edx
	popq	%rbx
	.cfi_def_cfa_offset 40
	movl	%eax, %edi
	leaq	.LC1(%rip), %rsi
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	_Z9gpuAssert9cudaErrorPKcib
	.cfi_endproc
.LFE5324:
	.size	_Z4prodRK4gMatS1_RS_i, .-_Z4prodRK4gMatS1_RS_i
	.section	.rodata.str1.1
.LC2:
	.string	"_Z12matMulKernelPfS_S_iiii"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB5353:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 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	.LC2(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z12matMulKernelPfS_S_iiii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	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	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE5353:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text._ZNSt6vectorIfSaIfEED2Ev,"axG",@progbits,_ZNSt6vectorIfSaIfEED5Ev,comdat
	.align 2
	.weak	_ZNSt6vectorIfSaIfEED2Ev
	.type	_ZNSt6vectorIfSaIfEED2Ev, @function
_ZNSt6vectorIfSaIfEED2Ev:
.LFB5688:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rax
	testq	%rax, %rax
	je	.L42
	movq	16(%rdi), %rsi
	movq	%rax, %rdi
	subq	%rax, %rsi
	jmp	_ZdlPvm@PLT
.L42:
	ret
	.cfi_endproc
.LFE5688:
	.size	_ZNSt6vectorIfSaIfEED2Ev, .-_ZNSt6vectorIfSaIfEED2Ev
	.weak	_ZNSt6vectorIfSaIfEED1Ev
	.set	_ZNSt6vectorIfSaIfEED1Ev,_ZNSt6vectorIfSaIfEED2Ev
	.section	.text._ZNSt6vectorIfSaIfEEC2ERKS1_,"axG",@progbits,_ZNSt6vectorIfSaIfEEC5ERKS1_,comdat
	.align 2
	.weak	_ZNSt6vectorIfSaIfEEC2ERKS1_
	.type	_ZNSt6vectorIfSaIfEEC2ERKS1_, @function
_ZNSt6vectorIfSaIfEEC2ERKS1_:
.LFB5704:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	8(%rsi), %r12
	movq	%rdi, %rbx
	subq	(%rsi), %r12
	movq	%rax, (%rdi)
	movq	%rax, 8(%rdi)
	movq	%rax, 16(%rdi)
	movq	%r12, %rdi
	sarq	$2, %rdi
	call	_ZNSt12_Vector_baseIfSaIfEE11_M_allocateEm.isra.0
	addq	%rax, %r12
	movq	%rax, (%rbx)
	movq	%rax, %rcx
	movq	%rax, 8(%rbx)
	movq	%r12, 16(%rbx)
	movq	0(%rbp), %rsi
	movq	8(%rbp), %rbp
	subq	%rsi, %rbp
	cmpq	$4, %rbp
	jle	.L45
	movq	%rbp, %rdx
	movq	%rax, %rdi
	call	memmove@PLT
	movq	%rax, %rcx
	jmp	.L46
.L45:
	jne	.L46
	movss	(%rsi), %xmm0
	movss	%xmm0, (%rax)
.L46:
	addq	%rbp, %rcx
	movq	%rcx, 8(%rbx)
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5704:
	.size	_ZNSt6vectorIfSaIfEEC2ERKS1_, .-_ZNSt6vectorIfSaIfEEC2ERKS1_
	.weak	_ZNSt6vectorIfSaIfEEC1ERKS1_
	.set	_ZNSt6vectorIfSaIfEEC1ERKS1_,_ZNSt6vectorIfSaIfEEC2ERKS1_
	.section	.text._ZNKSt6vectorIfSaIfEE12_M_check_lenEmPKc,"axG",@progbits,_ZNKSt6vectorIfSaIfEE12_M_check_lenEmPKc,comdat
	.align 2
	.weak	_ZNKSt6vectorIfSaIfEE12_M_check_lenEmPKc
	.type	_ZNKSt6vectorIfSaIfEE12_M_check_lenEmPKc, @function
_ZNKSt6vectorIfSaIfEE12_M_check_lenEmPKc:
.LFB6025:
	.cfi_startproc
	endbr64
	movq	%rdi, %rax
	movq	%rdx, %rdi
	movq	8(%rax), %rcx
	subq	(%rax), %rcx
	movabsq	$2305843009213693951, %rax
	sarq	$2, %rcx
	movq	%rax, %rdx
	subq	%rcx, %rdx
	cmpq	%rsi, %rdx
	jnb	.L49
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZSt20__throw_length_errorPKc@PLT
.L49:
	.cfi_def_cfa_offset 8
	cmpq	%rcx, %rsi
	cmovb	%rcx, %rsi
	addq	%rsi, %rcx
	jc	.L48
	movabsq	$2305843009213693951, %rax
	cmpq	%rax, %rcx
	cmovbe	%rcx, %rax
.L48:
	ret
	.cfi_endproc
.LFE6025:
	.size	_ZNKSt6vectorIfSaIfEE12_M_check_lenEmPKc, .-_ZNKSt6vectorIfSaIfEE12_M_check_lenEmPKc
	.section	.rodata.str1.1
.LC3:
	.string	"vector::_M_realloc_insert"
	.text
	.align 2
	.globl	_ZN4gMat8tovectorEv
	.type	_ZN4gMat8tovectorEv, @function
_ZN4gMat8tovectorEv:
.LFB5302:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA5302
	endbr64
	movabsq	$2305843009213693950, %rax
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdi, %rbx
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movl	8(%rsi), %edi
	imull	12(%rsi), %edi
	movslq	%edi, %rdi
	cmpq	%rdi, %rax
	jb	.L58
	movq	%rsi, %rbp
	salq	$2, %rdi
.LEHB2:
	call	_Znam@PLT
	movl	8(%rbp), %edx
	imull	12(%rbp), %edx
	movl	$2, %ecx
	movq	0(%rbp), %rsi
	movq	%rax, %rdi
	movq	%rax, 24(%rsp)
	sall	$2, %edx
	movslq	%edx, %rdx
	call	cudaMemcpy@PLT
	xorl	%edx, %edx
	xorl	%ecx, %ecx
	movq	%rdx, (%rbx)
	movq	%rdx, 8(%rbx)
	movq	%rdx, 16(%rbx)
	movl	%ecx, 16(%rsp)
	jmp	.L59
.L58:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE2:
.L65:
	movl	16(%rsp), %ecx
	movq	8(%rbx), %r12
	imull	%ecx, %eax
	movl	20(%rsp), %ecx
	addl	%ecx, %eax
	movq	24(%rsp), %rcx
	cltq
	leaq	(%rcx,%rax,4), %r15
	cmpq	16(%rbx), %r12
	je	.L60
	movss	(%r15), %xmm0
	addq	$4, %r12
	movss	%xmm0, -4(%r12)
	movq	%r12, 8(%rbx)
.L61:
	incl	20(%rsp)
.L66:
	movl	12(%rbp), %eax
	movl	20(%rsp), %ecx
	cmpl	%ecx, %eax
	jg	.L65
	incl	16(%rsp)
.L59:
	movl	16(%rsp), %eax
	cmpl	%eax, 8(%rbp)
	jle	.L76
	xorl	%eax, %eax
	movl	%eax, 20(%rsp)
	jmp	.L66
.L60:
	leaq	.LC3(%rip), %rdx
	movl	$1, %esi
	movq	%rbx, %rdi
.LEHB3:
	call	_ZNKSt6vectorIfSaIfEE12_M_check_lenEmPKc
	movq	%rax, 32(%rsp)
	movq	8(%rbx), %rax
	movq	(%rbx), %r14
	movq	32(%rsp), %rdi
	movq	%rax, 40(%rsp)
	movq	%r12, %rax
	subq	%r14, %rax
	movq	%rax, 8(%rsp)
	call	_ZNSt12_Vector_baseIfSaIfEE11_M_allocateEm.isra.0
.LEHE3:
	movq	%rax, %r13
	movss	(%r15), %xmm0
	movq	8(%rsp), %rax
	movss	%xmm0, 0(%r13,%rax)
	testq	%rax, %rax
	jle	.L62
	movq	%rax, %rdx
	movq	%r14, %rsi
	movq	%r13, %rdi
	call	memmove@PLT
.L62:
	movq	40(%rsp), %r15
	movq	8(%rsp), %rax
	subq	%r12, %r15
	leaq	4(%r13,%rax), %rcx
	testq	%r15, %r15
	jle	.L63
	movq	%rcx, %rdi
	movq	%r15, %rdx
	movq	%r12, %rsi
	call	memmove@PLT
	movq	%rax, %rcx
.L63:
	addq	%rcx, %r15
	testq	%r14, %r14
	je	.L64
	movq	16(%rbx), %rsi
	movq	%r14, %rdi
	subq	%r14, %rsi
	call	_ZdlPvm@PLT
.L64:
	movq	32(%rsp), %rax
	movq	%r13, (%rbx)
	movq	%r15, 8(%rbx)
	leaq	0(%r13,%rax,4), %rax
	movq	%rax, 16(%rbx)
	jmp	.L61
.L69:
	endbr64
	movq	%rax, %rbp
.L67:
	movq	%rbx, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	%rbp, %rdi
.LEHB4:
	call	_Unwind_Resume@PLT
.LEHE4:
.L76:
	movq	24(%rsp), %rdi
	call	_ZdaPv@PLT
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5302:
	.section	.gcc_except_table
.LLSDA5302:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE5302-.LLSDACSB5302
.LLSDACSB5302:
	.uleb128 .LEHB2-.LFB5302
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB3-.LFB5302
	.uleb128 .LEHE3-.LEHB3
	.uleb128 .L69-.LFB5302
	.uleb128 0
	.uleb128 .LEHB4-.LFB5302
	.uleb128 .LEHE4-.LEHB4
	.uleb128 0
	.uleb128 0
.LLSDACSE5302:
	.text
	.size	_ZN4gMat8tovectorEv, .-_ZN4gMat8tovectorEv
	.section	.rodata.str1.1
.LC4:
	.string	","
.LC5:
	.string	"\n"
	.text
	.globl	_ZlsRSoR4gMat
	.type	_ZlsRSoR4gMat, @function
_ZlsRSoR4gMat:
.LFB5325:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA5325
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC4(%rip), %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdi, %rbx
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %r14
	movq	%r14, %rdi
.LEHB5:
	call	_ZN4gMat8tovectorEv
.LEHE5:
.L78:
	cmpl	%r12d, 8(%rbp)
	jle	.L88
	xorl	%r13d, %r13d
.L80:
	movl	12(%rbp), %eax
	cmpl	%r13d, %eax
	jle	.L89
	imull	%r12d, %eax
	movq	(%rsp), %rdx
	movq	%rbx, %rdi
	addl	%r13d, %eax
	cltq
	cvtss2sd	(%rdx,%rax,4), %xmm0
.LEHB6:
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	movl	$1, %edx
	movq	%r15, %rsi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
	incl	%r13d
	jmp	.L80
.L89:
	movl	$1, %edx
	leaq	.LC5(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
.LEHE6:
	incl	%r12d
	jmp	.L78
.L88:
	movq	%r14, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L83
	jmp	.L86
.L85:
	endbr64
	movq	%rax, %rbx
.L81:
	movq	%r14, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L86
	movq	%rbx, %rdi
.LEHB7:
	call	_Unwind_Resume@PLT
.LEHE7:
.L86:
	call	__stack_chk_fail@PLT
.L83:
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5325:
	.section	.gcc_except_table
.LLSDA5325:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE5325-.LLSDACSB5325
.LLSDACSB5325:
	.uleb128 .LEHB5-.LFB5325
	.uleb128 .LEHE5-.LEHB5
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB6-.LFB5325
	.uleb128 .LEHE6-.LEHB6
	.uleb128 .L85-.LFB5325
	.uleb128 0
	.uleb128 .LEHB7-.LFB5325
	.uleb128 .LEHE7-.LEHB7
	.uleb128 0
	.uleb128 0
.LLSDACSE5325:
	.text
	.size	_ZlsRSoR4gMat, .-_ZlsRSoR4gMat
	.align 2
	.type	_ZNSt6vectorIfSaIfEE12emplace_backIJfEEERfDpOT_.isra.0, @function
_ZNSt6vectorIfSaIfEE12emplace_backIJfEEERfDpOT_.isra.0:
.LFB6203:
	.cfi_startproc
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movq	8(%rdi), %rbx
	cmpq	16(%rdi), %rbx
	je	.L91
	movss	(%rsi), %xmm0
	addq	$4, %rbx
	movss	%xmm0, -4(%rbx)
	movq	%rbx, 8(%rdi)
	jmp	.L90
.L91:
	movq	%rsi, 8(%rsp)
	leaq	.LC3(%rip), %rdx
	movl	$1, %esi
	movq	%rbx, %r15
	call	_ZNKSt6vectorIfSaIfEE12_M_check_lenEmPKc
	movq	(%r14), %r12
	movq	8(%r14), %r13
	movq	%rax, %rdi
	movq	%rax, (%rsp)
	call	_ZNSt12_Vector_baseIfSaIfEE11_M_allocateEm.isra.0
	movq	8(%rsp), %r8
	subq	%r12, %r15
	movq	%rax, %rbp
	movss	(%r8), %xmm0
	movss	%xmm0, (%rax,%r15)
	testq	%r15, %r15
	jle	.L93
	movq	%r15, %rdx
	movq	%r12, %rsi
	movq	%rax, %rdi
	call	memmove@PLT
.L93:
	subq	%rbx, %r13
	leaq	4(%rbp,%r15), %rdi
	testq	%r13, %r13
	jle	.L94
	movq	%r13, %rdx
	movq	%rbx, %rsi
	call	memmove@PLT
	movq	%rax, %rdi
.L94:
	addq	%rdi, %r13
	testq	%r12, %r12
	je	.L95
	movq	16(%r14), %rsi
	movq	%r12, %rdi
	subq	%r12, %rsi
	call	_ZdlPvm@PLT
.L95:
	movq	(%rsp), %rax
	movq	%rbp, (%r14)
	movq	%r13, 8(%r14)
	leaq	0(%rbp,%rax,4), %rax
	movq	%rax, 16(%r14)
.L90:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE6203:
	.size	_ZNSt6vectorIfSaIfEE12emplace_backIJfEEERfDpOT_.isra.0, .-_ZNSt6vectorIfSaIfEE12emplace_backIJfEEERfDpOT_.isra.0
	.globl	_Z3eyei
	.type	_Z3eyei, @function
_Z3eyei:
.LFB5323:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA5323
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	%esi, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%r14, %rbx
	subq	$64, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	movq	%rax, 8(%rsp)
	movq	%rax, 16(%rsp)
	movq	%rax, 24(%rsp)
.L101:
	cmpq	%r14, %r12
	je	.L117
	xorl	%r13d, %r13d
.L105:
	cmpq	%r14, %r13
	je	.L118
	cmpl	%r12d, %r13d
	leaq	32(%rsp), %rsi
	leaq	8(%rsp), %rdi
	jne	.L102
	movl	$0x3f800000, 32(%rsp)
	jmp	.L116
.L102:
	movl	$0x00000000, 32(%rsp)
.L116:
.LEHB8:
	call	_ZNSt6vectorIfSaIfEE12emplace_backIJfEEERfDpOT_.isra.0
	incq	%r13
	jmp	.L105
.L118:
	incq	%r12
	jmp	.L101
.L117:
	leaq	8(%rsp), %r13
	leaq	32(%rsp), %r12
	movq	%r13, %rsi
	movq	%r12, %rdi
	call	_ZNSt6vectorIfSaIfEEC1ERKS1_
.LEHE8:
	xorl	%r8d, %r8d
	movl	%ebx, %ecx
	movl	%ebx, %edx
	movq	%r12, %rsi
	movq	%rbp, %rdi
.LEHB9:
	call	_ZN4gMatC1ESt6vectorIfSaIfEEiii
.LEHE9:
	movq	%r12, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	%r13, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L109
	jmp	.L113
.L112:
	endbr64
	movq	%rax, %rbx
.L106:
	movq	%r12, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	jmp	.L107
.L111:
	endbr64
	movq	%rax, %rbx
.L107:
	leaq	8(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L113
	movq	%rbx, %rdi
.LEHB10:
	call	_Unwind_Resume@PLT
.LEHE10:
.L113:
	call	__stack_chk_fail@PLT
.L109:
	addq	$64, %rsp
	.cfi_def_cfa_offset 48
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5323:
	.section	.gcc_except_table
.LLSDA5323:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE5323-.LLSDACSB5323
.LLSDACSB5323:
	.uleb128 .LEHB8-.LFB5323
	.uleb128 .LEHE8-.LEHB8
	.uleb128 .L111-.LFB5323
	.uleb128 0
	.uleb128 .LEHB9-.LFB5323
	.uleb128 .LEHE9-.LEHB9
	.uleb128 .L112-.LFB5323
	.uleb128 0
	.uleb128 .LEHB10-.LFB5323
	.uleb128 .LEHE10-.LEHB10
	.uleb128 0
	.uleb128 0
.LLSDACSE5323:
	.text
	.size	_Z3eyei, .-_Z3eyei
	.globl	_Z8randgMatii
	.type	_Z8randgMatii, @function
_Z8randgMatii:
.LFB5319:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA5319
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	%esi, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	imull	%edx, %r15d
	xorl	%r14d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movl	$1, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%edx, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	movq	%rax, 24(%rsp)
	movq	%rax, 32(%rsp)
	movq	%rax, 40(%rsp)
.L120:
	leaq	24(%rsp), %rax
	cmpl	%r14d, %r15d
	leaq	48(%rsp), %rbx
	movq	%rax, (%rsp)
	jle	.L134
	imulq	$16807, %r13, %rax
	movl	$2147483647, %ecx
	xorl	%edx, %edx
	divq	%rcx
	movq	%rdx, %rax
	movq	%rdx, %r13
	decq	%rax
	js	.L121
	cvtsi2ssq	%rax, %xmm0
	jmp	.L122
.L121:
	movq	%rax, %rdx
	andl	$1, %eax
	shrq	%rdx
	orq	%rax, %rdx
	cvtsi2ssq	%rdx, %xmm0
	addss	%xmm0, %xmm0
.L122:
	xorps	%xmm1, %xmm1
	addss	%xmm1, %xmm0
	mulss	.LC9(%rip), %xmm0
	comiss	.LC6(%rip), %xmm0
	jb	.L123
	movss	.LC8(%rip), %xmm0
.L123:
	xorps	%xmm1, %xmm1
	movq	(%rsp), %rdi
	movq	%rbx, %rsi
	addss	%xmm1, %xmm0
	movss	%xmm0, 48(%rsp)
.LEHB11:
	call	_ZNSt6vectorIfSaIfEE12emplace_backIJfEEERfDpOT_.isra.0
	incl	%r14d
	jmp	.L120
.L134:
	movq	%rax, %rsi
	movq	%rbx, %rdi
	call	_ZNSt6vectorIfSaIfEEC1ERKS1_
.LEHE11:
	movq	8(%rsp), %rdi
	xorl	%r8d, %r8d
	movl	%r12d, %ecx
	movl	%ebp, %edx
	movq	%rbx, %rsi
.LEHB12:
	call	_ZN4gMatC1ESt6vectorIfSaIfEEiii
.LEHE12:
	movq	%rbx, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L128
	jmp	.L132
.L131:
	endbr64
	movq	%rax, %rbp
.L125:
	movq	%rbx, %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	jmp	.L126
.L130:
	endbr64
	movq	%rax, %rbp
.L126:
	movq	(%rsp), %rdi
	call	_ZNSt6vectorIfSaIfEED1Ev
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L132
	movq	%rbp, %rdi
.LEHB13:
	call	_Unwind_Resume@PLT
.LEHE13:
.L132:
	call	__stack_chk_fail@PLT
.L128:
	movq	8(%rsp), %rax
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE5319:
	.section	.gcc_except_table
.LLSDA5319:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE5319-.LLSDACSB5319
.LLSDACSB5319:
	.uleb128 .LEHB11-.LFB5319
	.uleb128 .LEHE11-.LEHB11
	.uleb128 .L130-.LFB5319
	.uleb128 0
	.uleb128 .LEHB12-.LFB5319
	.uleb128 .LEHE12-.LEHB12
	.uleb128 .L131-.LFB5319
	.uleb128 0
	.uleb128 .LEHB13-.LFB5319
	.uleb128 .LEHE13-.LEHB13
	.uleb128 0
	.uleb128 0
.LLSDACSE5319:
	.text
	.size	_Z8randgMatii, .-_Z8randgMatii
	.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.cst4,"aM",@progbits,4
	.align 4
.LC6:
	.long	1065353216
	.align 4
.LC8:
	.long	1065353215
	.align 4
.LC9:
	.long	805306368
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
#include <vector>
#include <memory>
#include <string>
__global__ void matMulKernel(float* A, float* B, float* P, int m, int  n, int s, int tile_size);
class gMat {
	private:
		float* d_data;
		int rows;
		int cols;
	public:
		std::string name;
		gMat(std::vector<float> data, int r, int c,int dev=0); 	
		void cleanup();
		//~gMat();
	 	int getrows();
	    int getcols();
		std::vector<float> tovector();					
		friend std::ostream& operator<<(std::ostream& os,  gMat& gm);  
		friend void prod(const gMat& A, const gMat& B, gMat& C, int mem);
};

gMat randgMat(int r, int c);
gMat eye(int r);

	.text
	.file	"gMat.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

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

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

	.end_amdgpu_metadata
