#include <iostream>

/*
written by George Strauch on 4/03/2020

c++ program for matrix multiply using 1d arrays on the GPU
the GPU makes use of parallelism to make processes like this much faster

This implementation only uses square matrices as they are much
easier to debug, calculate and work with, however all functions can work with
non-square matrices too.

This program does not ((yet)) use a struct to make working with the
matrices easier as this runs into problems with CUDA however this will change

Execution syntax:
$ ./exec {int matrix_size} {int print_option}
where the print option can be:
1: Prints the whole of each matrix for debugging
and best used with smaller matrices <= 10.
2: Shows only the first and last element of the result.
other or no option: does not print anything.

Example run:
$ nvcc gpu_mm_device_mem.cu -o gpu
$ time ./gpu 10 1
$ time ./gpu 1500 2
*/



// calculate a single element of the matrix result of m1*m2
// res_x = res cols = m2 cols  max
// res_y = res rows = m1 rows  max
// common = m1_cols and m2_rows
__global__
void matmul(long long int* m1, long long int* m2, long long int* res, int common, int res_x, int res_y)
{
  // gets the id for the x and y position within the result
  int idx = blockIdx.x*blockDim.x + threadIdx.x;
  int idy = blockIdx.y*blockDim.y+ threadIdx.y;

  // there will be some invalid calls of this fuction
  // this checks to make sure this is calculating a valid element
  if(idx >= res_x || idy >= res_y) { return; }

  int id = idy*res_x + idx;
  res[id] = 0;
  for (size_t i = 0; i < common; i++) {
    res[id] += m1[res_y*idy+i] * m2[i*res_x+idx];
  }
}



// host side function to display matrix
__host__
void displayMatrix(long long int *mat, long long int rows, long long int cols)
{
  for (size_t i = 0; i < rows; i++) {
    for (size_t j = 0; j < cols; j++) {
      std::cout << mat[i*cols + j] << ' ';
    }
    std::cout << '\n';
  }
  std::cout << '\n';
}



// host side function to transpose
__host__
void transpose(long long int *&mat, long long int rows, long long int cols)
{
  long long int *new_mat = (long long int*)malloc(rows*cols*sizeof(long long int));
  for (size_t i = 0; i < rows; i++) {
    for (size_t j = 0; j < cols; j++) {
      new_mat[j*cols+i] = mat[i*rows+j];
    }
  }
  free(mat);
  mat = new_mat;
}



int main(int argc, char const *argv[])
{
  // gets the matrix size from user, see header
  int N = atoi(argv[1]);
  std::cout << "N: " << N << '\n';

  // host side matrices
  long long int *m1 = (long long int*)malloc(N*N*sizeof(long long int));
  long long int *m2 = (long long int*)malloc(N*N*sizeof(long long int));
  long long int *rs = (long long int*)malloc(N*N*sizeof(long long int));

  // device side matricies
  long long int *d_m1;
  long long int *d_m2;
  long long int *d_rs;

  // allocate memory for the matrices on the device
  cudaMalloc(&d_m1, N*N*sizeof(long long int));
  cudaMalloc(&d_m2, N*N*sizeof(long long int));
  cudaMalloc(&d_rs, N*N*sizeof(long long int));

  // initalizes the host matrices
  for (size_t a = 0; a < N; a++) {
    for (size_t b = 0; b < N; b++) {
      m1[a*N+b] = b;
      m2[a*N+b] = b;
    }
  }

  // transpose(m2, N, N);

  // copies host matrices to devices memory
  cudaMemcpy(d_m1, m1, N*N*sizeof(long long int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_m2, m2, N*N*sizeof(long long int), cudaMemcpyHostToDevice);


  // options for building the block grid. Subject to Change.
  //------------------------------------------
  int thdX = 16;
  int thdY = 16;
  dim3 threads_in_block(thdX, thdY);
  dim3 block_grid((N/thdX)+1, (N/thdY)+1);
  //------------------------------------------

  // debug info about the block grid
  std::cout << "----------------------------" << '\n';
  std::cout << "gird x: " << block_grid.x << " grid y: " << block_grid.y << '\n';
  std::cout << "total threads: " << block_grid.x * block_grid.y * thdX * thdY << ", needed: " << N*N << '\n';
  std::cout << "----------------------------" << '\n';

  // performs matrix multiply on the GPU, each thread will handle one element
  // then copys the result to host memory
  std::cout << "\nstart" << '\n';
  matmul<<<block_grid, threads_in_block>>>(d_m1, d_m2, d_rs, N, N, N);
  cudaMemcpy(rs, d_rs, N*N*sizeof(long long int), cudaMemcpyDeviceToHost);
  std::cout << "done\n" << '\n';


// display array
// display options listed in header
  if (argc > 2) {
    if (atoi(argv[2]) == 1) {
      std::cout << "matrix 1: " << '\n';
      displayMatrix(m1, N, N);

      std::cout << "matrix 2: " << '\n';
      displayMatrix(m2, N, N);

      std::cout << "result: " << '\n';
      displayMatrix(rs, N, N);
    }
    else if (atoi(argv[2]) == 2) {
      std::cout << "first: " << rs[0] << '\n';
      std::cout << "last: " << rs[N*N-1] << '\n';
      std::cout << '\n';
    }
  }

  // frees device and host memeory
  std::cout << "freeing memory" << '\n';
  cudaFree(d_m1);
  cudaFree(d_m2);
  cudaFree(d_rs);

  free(m1);
  free(m2);
  free(rs);

  return 0;
}














//

	.file	"tmpxft_00282c15_00000000-6_gpu_mm_device_mem.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3640:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3640:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0, @function
_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0:
.LFB4292:
	.cfi_startproc
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movq	(%rdi), %rax
	movb	%sil, 12(%rsp)
	movq	-24(%rax), %rax
	cmpq	$0, 16(%rdi,%rax)
	je	.L3
	leaq	12(%rsp), %rsi
	movl	$1, %edx
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
	jmp	.L2
.L3:
	movsbl	12(%rsp), %esi
	call	_ZNSo3putEc@PLT
.L2:
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4292:
	.size	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0, .-_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	.globl	_Z13displayMatrixPxxx
	.type	_Z13displayMatrixPxxx, @function
_Z13displayMatrixPxxx:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	xorl	%r14d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	xorl	%r13d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdx, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
.L8:
	cmpq	%r15, %r14
	jnb	.L10
	leaq	(%r12,%r13,8), %rcx
	xorl	%ebp, %ebp
.L11:
	cmpq	%rbx, %rbp
	jnb	.L13
	movq	(%rcx,%rbp,8), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movq	%rcx, 8(%rsp)
	incq	%rbp
	call	_ZNSo9_M_insertIxEERSoT_@PLT
	movl	$32, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	movq	8(%rsp), %rcx
	jmp	.L11
.L13:
	movl	$10, %esi
	leaq	_ZSt4cout(%rip), %rdi
	incq	%r14
	addq	%rbx, %r13
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	jmp	.L8
.L10:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	movl	$10, %esi
	leaq	_ZSt4cout(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	.cfi_endproc
.LFE3635:
	.size	_Z13displayMatrixPxxx, .-_Z13displayMatrixPxxx
	.globl	_Z9transposeRPxxx
	.type	_Z9transposeRPxxx, @function
_Z9transposeRPxxx:
.LFB3636:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	movq	%rdi, %rbx
	movq	%rsi, %rdi
	imulq	%rdx, %rdi
	pushq	%rcx
	.cfi_def_cfa_offset 48
	salq	$3, %rdi
	call	malloc@PLT
	leaq	0(,%r12,8), %r9
	xorl	%edi, %edi
	movq	%rax, %r13
	xorl	%eax, %eax
.L15:
	cmpq	%rbp, %rax
	je	.L17
	movq	(%rbx), %r8
	leaq	0(,%rdi,8), %rsi
	leaq	0(%r13,%rax,8), %rcx
	xorl	%edx, %edx
.L18:
	cmpq	%r12, %rdx
	je	.L20
	movq	(%r8,%rsi), %r10
	incq	%rdx
	addq	$8, %rsi
	movq	%r10, (%rcx)
	addq	%r9, %rcx
	jmp	.L18
.L20:
	incq	%rax
	addq	%rbp, %rdi
	jmp	.L15
.L17:
	movq	(%rbx), %rdi
	call	free@PLT
	movq	%r13, (%rbx)
	popq	%rax
	.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
.LFE3636:
	.size	_Z9transposeRPxxx, .-_Z9transposeRPxxx
	.globl	_Z32__device_stub__Z6matmulPxS_S_iiiPxS_S_iii
	.type	_Z32__device_stub__Z6matmulPxS_S_iiiPxS_S_iii, @function
_Z32__device_stub__Z6matmulPxS_S_iiiPxS_S_iii:
.LFB3662:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	%rdi, 40(%rsp)
	leaq	72(%rsp), %rdi
	movq	%rsi, 32(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rdx, 24(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 20(%rsp)
	leaq	64(%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, 80(%rsp)
	movq	%rax, 120(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z6matmulPxS_S_iii(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	136(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L21:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L23
	call	__stack_chk_fail@PLT
.L23:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3662:
	.size	_Z32__device_stub__Z6matmulPxS_S_iiiPxS_S_iii, .-_Z32__device_stub__Z6matmulPxS_S_iiiPxS_S_iii
	.globl	_Z6matmulPxS_S_iii
	.type	_Z6matmulPxS_S_iii, @function
_Z6matmulPxS_S_iii:
.LFB3663:
	.cfi_startproc
	endbr64
	jmp	_Z32__device_stub__Z6matmulPxS_S_iiiPxS_S_iii
	.cfi_endproc
.LFE3663:
	.size	_Z6matmulPxS_S_iii, .-_Z6matmulPxS_S_iii
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"N: "
.LC1:
	.string	"----------------------------"
.LC2:
	.string	"gird x: "
.LC3:
	.string	" grid y: "
.LC4:
	.string	"total threads: "
.LC5:
	.string	", needed: "
.LC6:
	.string	"\nstart"
.LC7:
	.string	"done\n"
.LC8:
	.string	"matrix 1: "
.LC9:
	.string	"matrix 2: "
.LC10:
	.string	"result: "
.LC11:
	.string	"first: "
.LC12:
	.string	"last: "
.LC13:
	.string	"freeing memory"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3637:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %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
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movl	%edi, 12(%rsp)
	movq	8(%rsi), %rdi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	call	atoi@PLT
	leaq	.LC0(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movslq	%eax, %rbx
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	movl	%ebx, %eax
	imull	%ebx, %eax
	movl	%eax, 8(%rsp)
	cltq
	salq	$3, %rax
	movq	%rax, %rdi
	movq	%rax, (%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, %r14
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, %r13
	call	malloc@PLT
	movq	(%rsp), %rsi
	leaq	24(%rsp), %rdi
	movq	%rax, %rbp
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	leaq	32(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	(%rsp), %rsi
	leaq	40(%rsp), %rdi
	call	cudaMalloc@PLT
	movslq	%ebx, %rcx
	xorl	%edi, %edi
	xorl	%esi, %esi
.L27:
	cmpq	%rcx, %rsi
	jnb	.L36
	leaq	0(,%rdi,8), %rax
	leaq	(%r14,%rax), %r8
	leaq	0(%r13,%rax), %rdx
	xorl	%eax, %eax
.L28:
	movq	%rax, (%r8,%rax,8)
	movq	%rax, (%rdx,%rax,8)
	incq	%rax
	cmpq	%rcx, %rax
	jb	.L28
	incq	%rsi
	addq	%rcx, %rdi
	jmp	.L27
.L36:
	movq	(%rsp), %rdx
	movq	24(%rsp), %rdi
	movl	$1, %ecx
	movq	%r14, %rsi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rdx
	movq	32(%rsp), %rdi
	movq	%r13, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$16, %ecx
	movabsq	$68719476752, %rax
	leaq	.LC1(%rip), %rsi
	movq	%rax, 48(%rsp)
	movl	%ebx, %eax
	leaq	_ZSt4cout(%rip), %rdi
	cltd
	movl	$1, 56(%rsp)
	idivl	%ecx
	movl	$1, 68(%rsp)
	leal	1(%rax), %r12d
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	leaq	.LC2(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r12d, %esi
	movq	%rax, %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	leaq	.LC3(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r12d, %esi
	movq	%rax, %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	leaq	.LC4(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r12d, %esi
	imull	%r12d, %esi
	movq	%rax, %rdi
	sall	$8, %esi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	leaq	.LC5(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	8(%rsp), %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	leaq	.LC1(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	leaq	.LC6(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	movl	%r12d, 60(%rsp)
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	56(%rsp), %ecx
	movq	48(%rsp), %rdx
	movl	%r12d, 64(%rsp)
	movl	68(%rsp), %esi
	movq	60(%rsp), %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L30
	movq	40(%rsp), %rdx
	movq	32(%rsp), %rsi
	movl	%ebx, %r9d
	movl	%ebx, %r8d
	movq	24(%rsp), %rdi
	movl	%ebx, %ecx
	call	_Z32__device_stub__Z6matmulPxS_S_iiiPxS_S_iii
.L30:
	movq	(%rsp), %rdx
	movq	40(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbp, %rdi
	call	cudaMemcpy@PLT
	leaq	.LC7(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	cmpl	$2, 12(%rsp)
	jle	.L31
	movq	16(%r15), %rdi
	call	atoi@PLT
	cmpl	$1, %eax
	jne	.L32
	leaq	.LC8(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	movq	%rbx, %rdx
	movq	%rbx, %rsi
	movq	%r14, %rdi
	call	_Z13displayMatrixPxxx
	leaq	.LC9(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	movq	%rbx, %rdx
	movq	%rbx, %rsi
	movq	%r13, %rdi
	call	_Z13displayMatrixPxxx
	leaq	.LC10(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	movq	%rbx, %rdx
	movq	%rbx, %rsi
	movq	%rbp, %rdi
	call	_Z13displayMatrixPxxx
	jmp	.L31
.L32:
	cmpl	$2, %eax
	jne	.L31
	leaq	.LC11(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	0(%rbp), %rsi
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIxEERSoT_@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	leaq	.LC12(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rax, %rdi
	movq	(%rsp), %rax
	movq	-8(%rbp,%rax), %rsi
	call	_ZNSo9_M_insertIxEERSoT_@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	movl	$10, %esi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
.L31:
	leaq	.LC13(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	$10, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.isra.0
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
	movq	%r14, %rdi
	call	free@PLT
	movq	%r13, %rdi
	call	free@PLT
	movq	%rbp, %rdi
	call	free@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L33
	call	__stack_chk_fail@PLT
.L33:
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3637:
	.size	main, .-main
	.section	.rodata.str1.1
.LC14:
	.string	"_Z6matmulPxS_S_iii"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3665:
	.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	.LC14(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z6matmulPxS_S_iii(%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
.LFE3665:
	.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
	.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 <iostream>

/*
written by George Strauch on 4/03/2020

c++ program for matrix multiply using 1d arrays on the GPU
the GPU makes use of parallelism to make processes like this much faster

This implementation only uses square matrices as they are much
easier to debug, calculate and work with, however all functions can work with
non-square matrices too.

This program does not ((yet)) use a struct to make working with the
matrices easier as this runs into problems with CUDA however this will change

Execution syntax:
$ ./exec {int matrix_size} {int print_option}
where the print option can be:
1: Prints the whole of each matrix for debugging
and best used with smaller matrices <= 10.
2: Shows only the first and last element of the result.
other or no option: does not print anything.

Example run:
$ nvcc gpu_mm_device_mem.cu -o gpu
$ time ./gpu 10 1
$ time ./gpu 1500 2
*/



// calculate a single element of the matrix result of m1*m2
// res_x = res cols = m2 cols  max
// res_y = res rows = m1 rows  max
// common = m1_cols and m2_rows
__global__
void matmul(long long int* m1, long long int* m2, long long int* res, int common, int res_x, int res_y)
{
  // gets the id for the x and y position within the result
  int idx = blockIdx.x*blockDim.x + threadIdx.x;
  int idy = blockIdx.y*blockDim.y+ threadIdx.y;

  // there will be some invalid calls of this fuction
  // this checks to make sure this is calculating a valid element
  if(idx >= res_x || idy >= res_y) { return; }

  int id = idy*res_x + idx;
  res[id] = 0;
  for (size_t i = 0; i < common; i++) {
    res[id] += m1[res_y*idy+i] * m2[i*res_x+idx];
  }
}



// host side function to display matrix
__host__
void displayMatrix(long long int *mat, long long int rows, long long int cols)
{
  for (size_t i = 0; i < rows; i++) {
    for (size_t j = 0; j < cols; j++) {
      std::cout << mat[i*cols + j] << ' ';
    }
    std::cout << '\n';
  }
  std::cout << '\n';
}



// host side function to transpose
__host__
void transpose(long long int *&mat, long long int rows, long long int cols)
{
  long long int *new_mat = (long long int*)malloc(rows*cols*sizeof(long long int));
  for (size_t i = 0; i < rows; i++) {
    for (size_t j = 0; j < cols; j++) {
      new_mat[j*cols+i] = mat[i*rows+j];
    }
  }
  free(mat);
  mat = new_mat;
}



int main(int argc, char const *argv[])
{
  // gets the matrix size from user, see header
  int N = atoi(argv[1]);
  std::cout << "N: " << N << '\n';

  // host side matrices
  long long int *m1 = (long long int*)malloc(N*N*sizeof(long long int));
  long long int *m2 = (long long int*)malloc(N*N*sizeof(long long int));
  long long int *rs = (long long int*)malloc(N*N*sizeof(long long int));

  // device side matricies
  long long int *d_m1;
  long long int *d_m2;
  long long int *d_rs;

  // allocate memory for the matrices on the device
  hipMalloc(&d_m1, N*N*sizeof(long long int));
  hipMalloc(&d_m2, N*N*sizeof(long long int));
  hipMalloc(&d_rs, N*N*sizeof(long long int));

  // initalizes the host matrices
  for (size_t a = 0; a < N; a++) {
    for (size_t b = 0; b < N; b++) {
      m1[a*N+b] = b;
      m2[a*N+b] = b;
    }
  }

  // transpose(m2, N, N);

  // copies host matrices to devices memory
  hipMemcpy(d_m1, m1, N*N*sizeof(long long int), hipMemcpyHostToDevice);
  hipMemcpy(d_m2, m2, N*N*sizeof(long long int), hipMemcpyHostToDevice);


  // options for building the block grid. Subject to Change.
  //------------------------------------------
  int thdX = 16;
  int thdY = 16;
  dim3 threads_in_block(thdX, thdY);
  dim3 block_grid((N/thdX)+1, (N/thdY)+1);
  //------------------------------------------

  // debug info about the block grid
  std::cout << "----------------------------" << '\n';
  std::cout << "gird x: " << block_grid.x << " grid y: " << block_grid.y << '\n';
  std::cout << "total threads: " << block_grid.x * block_grid.y * thdX * thdY << ", needed: " << N*N << '\n';
  std::cout << "----------------------------" << '\n';

  // performs matrix multiply on the GPU, each thread will handle one element
  // then copys the result to host memory
  std::cout << "\nstart" << '\n';
  matmul<<<block_grid, threads_in_block>>>(d_m1, d_m2, d_rs, N, N, N);
  hipMemcpy(rs, d_rs, N*N*sizeof(long long int), hipMemcpyDeviceToHost);
  std::cout << "done\n" << '\n';


// display array
// display options listed in header
  if (argc > 2) {
    if (atoi(argv[2]) == 1) {
      std::cout << "matrix 1: " << '\n';
      displayMatrix(m1, N, N);

      std::cout << "matrix 2: " << '\n';
      displayMatrix(m2, N, N);

      std::cout << "result: " << '\n';
      displayMatrix(rs, N, N);
    }
    else if (atoi(argv[2]) == 2) {
      std::cout << "first: " << rs[0] << '\n';
      std::cout << "last: " << rs[N*N-1] << '\n';
      std::cout << '\n';
    }
  }

  // frees device and host memeory
  std::cout << "freeing memory" << '\n';
  hipFree(d_m1);
  hipFree(d_m2);
  hipFree(d_rs);

  free(m1);
  free(m2);
  free(rs);

  return 0;
}














//

	.text
	.file	"gpu_mm_device_mem.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z21__device_stub__matmulPxS_S_iii # -- Begin function _Z21__device_stub__matmulPxS_S_iii
	.type	_Z21__device_stub__matmulPxS_S_iii,@function
_Z21__device_stub__matmulPxS_S_iii:     # @_Z21__device_stub__matmulPxS_S_iii
	.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	20(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	16(%rsp), %rcx
	movl	%r8d, (%rcx)
	leaq	12(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	32(%rsp), %r12
	leaq	24(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z6matmulPxS_S_iii, %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_end0:
	.size	_Z21__device_stub__matmulPxS_S_iii, .Lfunc_end0-_Z21__device_stub__matmulPxS_S_iii
	.cfi_endproc
                                        # -- End function
	.globl	_Z13displayMatrixPxxx           # -- Begin function _Z13displayMatrixPxxx
	.type	_Z13displayMatrixPxxx,@function
_Z13displayMatrixPxxx:                  # @_Z13displayMatrixPxxx
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	testq	%rsi, %rsi
	je	.LBB1_6
# %bb.1:                                # %.preheader.lr.ph
	movq	%rdx, %rbx
	movq	%rsi, %r14
	movq	%rdi, %r15
	leaq	(,%rdx,8), %rax
	movq	%rax, 16(%rsp)                  # 8-byte Spill
	xorl	%r12d, %r12d
	leaq	15(%rsp), %r13
.LBB1_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_4 Depth 2
	testq	%rbx, %rbx
	je	.LBB1_10
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB1_2 Depth=1
	xorl	%ebp, %ebp
.LBB1_4:                                #   Parent Loop BB1_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	(%r15,%rbp,8), %rsi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIxEERSoT_
	movb	$32, 15(%rsp)
	movq	(%rax), %rcx
	movq	-24(%rcx), %rcx
	cmpq	$0, 16(%rax,%rcx)
	je	.LBB1_14
# %bb.5:                                #   in Loop: Header=BB1_4 Depth=2
	movl	$1, %edx
	movq	%rax, %rdi
	movq	%r13, %rsi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB1_15
.LBB1_14:                               #   in Loop: Header=BB1_4 Depth=2
	movq	%rax, %rdi
	movl	$32, %esi
	callq	_ZNSo3putEc
.LBB1_15:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit16
                                        #   in Loop: Header=BB1_4 Depth=2
	incq	%rbp
	cmpq	%rbp, %rbx
	jne	.LBB1_4
.LBB1_10:                               # %._crit_edge
                                        #   in Loop: Header=BB1_2 Depth=1
	movb	$10, 14(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	je	.LBB1_12
# %bb.11:                               #   in Loop: Header=BB1_2 Depth=1
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	leaq	14(%rsp), %rsi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB1_13
.LBB1_12:                               #   in Loop: Header=BB1_2 Depth=1
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB1_13:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit13
                                        #   in Loop: Header=BB1_2 Depth=1
	incq	%r12
	addq	16(%rsp), %r15                  # 8-byte Folded Reload
	cmpq	%r14, %r12
	jne	.LBB1_2
.LBB1_6:                                # %._crit_edge19
	movb	$10, 13(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	je	.LBB1_8
# %bb.7:
	leaq	13(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB1_9
.LBB1_8:
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB1_9:                                # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z13displayMatrixPxxx, .Lfunc_end1-_Z13displayMatrixPxxx
	.cfi_endproc
                                        # -- End function
	.globl	_Z9transposeRPxxx               # -- Begin function _Z9transposeRPxxx
	.type	_Z9transposeRPxxx,@function
_Z9transposeRPxxx:                      # @_Z9transposeRPxxx
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	pushq	%rax
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rdx, %r14
	movq	%rsi, %r15
	movq	%rdi, %rbx
	movq	%rsi, %rdi
	imulq	%rdx, %rdi
	shlq	$3, %rdi
	callq	malloc
	movq	%rax, %r12
	movq	(%rbx), %rdi
	testq	%r15, %r15
	je	.LBB2_6
# %bb.1:                                # %.preheader.lr.ph
	leaq	(,%r14,8), %rax
	leaq	(,%r15,8), %rcx
	xorl	%edx, %edx
	movq	%rdi, %rsi
	movq	%r12, %r8
.LBB2_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_4 Depth 2
	testq	%r14, %r14
	je	.LBB2_5
# %bb.3:                                # %.lr.ph
                                        #   in Loop: Header=BB2_2 Depth=1
	movq	%r8, %r9
	xorl	%r10d, %r10d
.LBB2_4:                                #   Parent Loop BB2_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	(%rsi,%r10,8), %r11
	movq	%r11, (%r9)
	incq	%r10
	addq	%rax, %r9
	cmpq	%r10, %r14
	jne	.LBB2_4
.LBB2_5:                                # %._crit_edge
                                        #   in Loop: Header=BB2_2 Depth=1
	incq	%rdx
	addq	$8, %r8
	addq	%rcx, %rsi
	cmpq	%r15, %rdx
	jne	.LBB2_2
.LBB2_6:                                # %._crit_edge23
	callq	free
	movq	%r12, (%rbx)
	addq	$8, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z9transposeRPxxx, .Lfunc_end2-_Z9transposeRPxxx
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edi, 36(%rsp)                  # 4-byte Spill
	movq	%rsi, 48(%rsp)                  # 8-byte Spill
	movq	8(%rsi), %rdi
	callq	atoi
	movl	%eax, %r13d
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$3, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%r13d, %esi
	callq	_ZNSolsEi
	movb	$10, 8(%rsp)
	movq	(%rax), %rcx
	movq	-24(%rcx), %rcx
	cmpq	$0, 16(%rax,%rcx)
	je	.LBB3_2
# %bb.1:
	leaq	8(%rsp), %rsi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_3
.LBB3_2:
	movq	%rax, %rdi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_3:                                # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit
	movl	%r13d, %r15d
	imull	%r15d, %r15d
	leaq	(,%r15,8), %rbp
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, %rbx
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, %r14
	movq	%rbp, %rdi
	callq	malloc
	movq	%rax, 56(%rsp)                  # 8-byte Spill
	leaq	8(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	24(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movq	%rbp, %rsi
	callq	hipMalloc
	movslq	%r13d, %r12
	testl	%r13d, %r13d
	je	.LBB3_8
# %bb.4:                                # %.preheader.preheader
	leaq	(,%r12,8), %rax
	xorl	%ecx, %ecx
	movq	%rbx, %rdx
	movq	%r14, %rsi
.LBB3_5:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_6 Depth 2
	xorl	%edi, %edi
.LBB3_6:                                #   Parent Loop BB3_5 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%rdi, (%rdx,%rdi,8)
	movq	%rdi, (%rsi,%rdi,8)
	incq	%rdi
	cmpq	%rdi, %r12
	jne	.LBB3_6
# %bb.7:                                #   in Loop: Header=BB3_5 Depth=1
	incq	%rcx
	addq	%rax, %rsi
	addq	%rax, %rdx
	cmpq	%r12, %rcx
	jne	.LBB3_5
.LBB3_8:                                # %._crit_edge
	movq	8(%rsp), %rdi
	movq	%rbx, 64(%rsp)                  # 8-byte Spill
	movq	%rbx, %rsi
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	24(%rsp), %rdi
	movq	%r14, %rsi
	movq	%rbp, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leal	15(%r13), %ebx
	testl	%r13d, %r13d
	cmovnsl	%r13d, %ebx
	sarl	$4, %ebx
	incl	%ebx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$28, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movb	$10, 7(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	movq	%r14, 40(%rsp)                  # 8-byte Spill
	je	.LBB3_10
# %bb.9:
	leaq	7(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_11
.LBB3_10:
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_11:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit82
	movl	$_ZSt4cout, %edi
	movl	$.L.str.2, %esi
	movl	$8, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movq	%rbx, %rsi
	callq	_ZNSo9_M_insertImEERSoT_
	movq	%rax, %r14
	movl	$.L.str.3, %esi
	movl	$9, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movq	%rbx, %rsi
	callq	_ZNSo9_M_insertImEERSoT_
	movb	$10, 7(%rsp)
	movq	(%rax), %rcx
	movq	-24(%rcx), %rcx
	cmpq	$0, 16(%rax,%rcx)
	je	.LBB3_13
# %bb.12:
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_14
.LBB3_13:
	movq	%rax, %rdi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_14:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit85
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$15, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	%ebx, %esi
	imull	%esi, %esi
	shll	$8, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertImEERSoT_
	movq	%rax, %r14
	movl	$.L.str.5, %esi
	movl	$10, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r14, %rdi
	movl	%r15d, %esi
	callq	_ZNSolsEi
	movb	$10, 7(%rsp)
	movq	(%rax), %rcx
	movq	-24(%rcx), %rcx
	cmpq	$0, 16(%rax,%rcx)
	je	.LBB3_16
# %bb.15:
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_17
.LBB3_16:
	movq	%rax, %rdi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_17:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit88
	movq	%rbx, %r14
	shlq	$32, %r14
	movl	$_ZSt4cout, %edi
	movl	$.L.str.1, %esi
	movl	$28, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movb	$10, 7(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	je	.LBB3_19
# %bb.18:
	leaq	7(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_20
.LBB3_19:
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_20:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit91
	orq	%r14, %rbx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.6, %esi
	movl	$6, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movb	$10, 7(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	je	.LBB3_22
# %bb.21:
	leaq	7(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_23
.LBB3_22:
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_23:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit94
	movq	40(%rsp), %r14                  # 8-byte Reload
	movabsq	$68719476752, %rdx              # imm = 0x1000000010
	movq	%rbx, %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_25
# %bb.24:
	movq	8(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdx
	movl	%r13d, %ecx
	movl	%r13d, %r8d
	movl	%r13d, %r9d
	callq	_Z21__device_stub__matmulPxS_S_iii
.LBB3_25:
	movq	16(%rsp), %rsi
	movq	56(%rsp), %r13                  # 8-byte Reload
	movq	%r13, %rdi
	movq	%rbp, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movl	$_ZSt4cout, %edi
	movl	$.L.str.7, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movb	$10, 7(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	movq	64(%rsp), %rbx                  # 8-byte Reload
	je	.LBB3_27
# %bb.26:
	leaq	7(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_28
.LBB3_27:
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_28:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit97
	cmpl	$3, 36(%rsp)                    # 4-byte Folded Reload
	jl	.LBB3_50
# %bb.29:
	movq	48(%rsp), %rax                  # 8-byte Reload
	movq	16(%rax), %rdi
	callq	atoi
	cmpl	$2, %eax
	je	.LBB3_41
# %bb.30:
	cmpl	$1, %eax
	jne	.LBB3_50
# %bb.31:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.8, %esi
	movl	$10, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movb	$10, 7(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	je	.LBB3_33
# %bb.32:
	leaq	7(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_34
.LBB3_41:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.11, %esi
	movl	$7, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	(%r13), %rsi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIxEERSoT_
	movb	$10, 7(%rsp)
	movq	(%rax), %rcx
	movq	-24(%rcx), %rcx
	cmpq	$0, 16(%rax,%rcx)
	je	.LBB3_43
# %bb.42:
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_44
.LBB3_33:
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_34:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit100
	movq	%rbx, %rdi
	movq	%r12, %rsi
	movq	%r12, %rdx
	callq	_Z13displayMatrixPxxx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.9, %esi
	movl	$10, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movb	$10, 7(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	je	.LBB3_36
# %bb.35:
	leaq	7(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_37
.LBB3_36:
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_37:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit103
	movq	%r14, %rdi
	movq	%r12, %rsi
	movq	%r12, %rdx
	callq	_Z13displayMatrixPxxx
	movl	$_ZSt4cout, %edi
	movl	$.L.str.10, %esi
	movl	$8, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movb	$10, 7(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	je	.LBB3_39
# %bb.38:
	leaq	7(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_40
.LBB3_39:
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_40:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit106
	movq	%r13, %rdi
	movq	%r12, %rsi
	movq	%r12, %rdx
	callq	_Z13displayMatrixPxxx
	jmp	.LBB3_50
.LBB3_43:
	movq	%rax, %rdi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_44:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit109
	movl	$_ZSt4cout, %edi
	movl	$.L.str.12, %esi
	movl	$6, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	-8(%r13,%r15,8), %rsi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIxEERSoT_
	movb	$10, 7(%rsp)
	movq	(%rax), %rcx
	movq	-24(%rcx), %rcx
	cmpq	$0, 16(%rax,%rcx)
	je	.LBB3_46
# %bb.45:
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_47
.LBB3_46:
	movq	%rax, %rdi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_47:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit112
	movb	$10, 7(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	je	.LBB3_49
# %bb.48:
	leaq	7(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_50
.LBB3_49:
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_50:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.13, %esi
	movl	$14, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movb	$10, 7(%rsp)
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	cmpq	$0, _ZSt4cout+16(%rax)
	je	.LBB3_52
# %bb.51:
	leaq	7(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_53
.LBB3_52:
	movl	$_ZSt4cout, %edi
	movl	$10, %esi
	callq	_ZNSo3putEc
.LBB3_53:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit118
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	%rbx, %rdi
	callq	free
	movq	%r14, %rdi
	callq	free
	movq	%r13, %rdi
	callq	free
	xorl	%eax, %eax
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	main, .Lfunc_end3-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:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z6matmulPxS_S_iii, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__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	.LBB5_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
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z6matmulPxS_S_iii,@object      # @_Z6matmulPxS_S_iii
	.section	.rodata,"a",@progbits
	.globl	_Z6matmulPxS_S_iii
	.p2align	3, 0x0
_Z6matmulPxS_S_iii:
	.quad	_Z21__device_stub__matmulPxS_S_iii
	.size	_Z6matmulPxS_S_iii, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"N: "
	.size	.L.str, 4

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"----------------------------"
	.size	.L.str.1, 29

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"gird x: "
	.size	.L.str.2, 9

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	" grid y: "
	.size	.L.str.3, 10

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"total threads: "
	.size	.L.str.4, 16

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	", needed: "
	.size	.L.str.5, 11

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"\nstart"
	.size	.L.str.6, 7

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"done\n"
	.size	.L.str.7, 6

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"matrix 1: "
	.size	.L.str.8, 11

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"matrix 2: "
	.size	.L.str.9, 11

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"result: "
	.size	.L.str.10, 9

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"first: "
	.size	.L.str.11, 8

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"last: "
	.size	.L.str.12, 7

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"freeing memory"
	.size	.L.str.13, 15

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z6matmulPxS_S_iii"
	.size	.L__unnamed_1, 19

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

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

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z6matmulPxS_S_iii      ; -- Begin function _Z6matmulPxS_S_iii
	.globl	_Z6matmulPxS_S_iii
	.p2align	8
	.type	_Z6matmulPxS_S_iii,@function
_Z6matmulPxS_S_iii:                     ; @_Z6matmulPxS_S_iii
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x34
	s_load_b128 s[4:7], s[0:1], 0x18
	v_and_b32_e32 v1, 0x3ff, v0
	v_bfe_u32 v0, v0, 10, 10
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s3, s2, 16
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[2:3], null, s14, s2, v[1:2]
	v_mad_u64_u32 v[3:4], null, s15, s3, v[0:1]
	v_cmp_gt_i32_e32 vcc_lo, s5, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_cmp_gt_i32_e64 s2, s6, v3
	s_and_b32 s2, vcc_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_saveexec_b32 s3, s2
	s_cbranch_execz .LBB0_4
; %bb.1:
	s_load_b64 s[2:3], s[0:1], 0x10
	v_mad_u64_u32 v[0:1], null, v3, s5, v[2:3]
	v_mov_b32_e32 v4, 0
	s_cmp_eq_u32 s4, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_mov_b32_e32 v5, v4
	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, s2, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s3, v1, vcc_lo
	global_store_b64 v[0:1], v[4:5], off
	s_cbranch_scc1 .LBB0_4
; %bb.2:                                ; %.lr.ph
	s_load_b128 s[8:11], s[0:1], 0x0
	v_mul_lo_u32 v4, v3, s6
	v_ashrrev_i32_e32 v3, 31, v2
	v_mov_b32_e32 v6, 0
	v_mov_b32_e32 v7, 0
	s_ashr_i32 s3, s5, 31
	s_mov_b32 s2, s5
	v_lshlrev_b64 v[8:9], 3, v[2:3]
	s_ashr_i32 s1, s4, 31
	v_ashrrev_i32_e32 v5, 31, v4
	s_mov_b32 s0, s4
	s_lshl_b64 s[2:3], s[2:3], 3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[4:5], 3, v[4:5]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v2, vcc_lo, s8, v4
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v5, vcc_lo
	v_add_co_u32 v4, vcc_lo, s10, v8
	v_add_co_ci_u32_e32 v5, vcc_lo, s11, v9, vcc_lo
.LBB0_3:                                ; =>This Inner Loop Header: Depth=1
	global_load_b64 v[8:9], v[2:3], off
	global_load_b64 v[10:11], v[4:5], off
	v_add_co_u32 v2, vcc_lo, v2, 8
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v3, vcc_lo
	v_add_co_u32 v4, vcc_lo, v4, s2
	s_add_u32 s0, s0, -1
	v_add_co_ci_u32_e32 v5, vcc_lo, s3, v5, vcc_lo
	s_addc_u32 s1, s1, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(VALU_DEP_1)
	s_cmp_lg_u64 s[0:1], 0
	s_waitcnt vmcnt(0)
	v_mad_u64_u32 v[12:13], null, v10, v8, v[6:7]
	v_mul_lo_u32 v6, v10, v9
	v_mul_lo_u32 v7, v11, v8
	v_add3_u32 v13, v7, v13, v6
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v6, v12 :: v_dual_mov_b32 v7, v13
	global_store_b64 v[0:1], v[12:13], off
	s_cbranch_scc1 .LBB0_3
.LBB0_4:                                ; %.loopexit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z6matmulPxS_S_iii
		.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_end0:
	.size	_Z6matmulPxS_S_iii, .Lfunc_end0-_Z6matmulPxS_S_iii
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 416
; 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
	.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:         28
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .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:           _Z6matmulPxS_S_iii
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z6matmulPxS_S_iii.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     14
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
