#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <cuda_runtime.h>

#ifndef __CUDACC__ 
#define __CUDACC__
#endif
#include "device_launch_parameters.h"
#include <cuda.h>
#include <device_functions.h>
#include <cuda_runtime_api.h>

// Matrices are stored in row-major order: 
// M(row, col) = *(M.elements + row * M.width + col) 
typedef struct { 
	int width; 
	int height; 
	float* elements; 
} Matrix; 

// Thread block size 
#define BLOCK_SIZE 16 
// Forward declaration of the matrix multiplication kernel 
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix); 
void MatMul(const Matrix A, const Matrix B, Matrix C);

int main() {
	float A[16] = { 1, 2, 3, 4, 5, 6, 1, 2, 3, 4, 5, 6, 1, 2, 3, 4 };
	float B[32] = { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 1, 2, 3, 4, 5, 6, 7, 8 };
	int rowsA, colsA, rowsB, colsB;
	rowsA = 4;
	colsA = 4;
	rowsB = 4;
	colsB = 8;
	float *C = (float*)calloc(rowsA * colsB, sizeof(float));

	Matrix Am, Bm, Cm;
	Am.height = rowsA;
	Am.width = colsA;
	Am.elements = A;
	Bm.height = rowsB;
	Bm.width = colsB;
	Bm.elements = B;
	Cm.height = rowsA;
	Cm.width = colsB;
	Cm.elements = C;

	MatMul(Am, Bm, Cm);

	for (int i = 0; i < Cm.height; ++i){
		for (int j = 0; j < Cm.width; ++j){
			printf("%f ", Cm.elements[i * Cm.height + j]);
		}
		printf("\n");
	}
	printf("\n");
	//print_matrix(Cm, Cm.height, Cm.width);
	//C = matmul(A, rowsA, colsA, B, rowsB, colsB);

}

// Matrix multiplication - Host code 
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE 
void MatMul(const Matrix A, const Matrix B, Matrix C) { 
	
	// Load A and B to device memory 
	Matrix d_A; d_A.width = A.width; d_A.height = A.height; 
	size_t size = A.width * A.height * sizeof(float); 
	cudaMalloc(&d_A.elements, size); 
	cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice); 
	
	Matrix d_B; d_B.width = B.width; d_B.height = B.height; 
	size = B.width * B.height * sizeof(float); 
	cudaMalloc(&d_B.elements, size); 
	cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice); 
	
	// Allocate C in device memory 
	Matrix d_C; d_C.width = C.width; d_C.height = C.height; 
	size = C.width * C.height * sizeof(float); 
	cudaMalloc(&d_C.elements, size); 
	
	// Invoke kernel 
	dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); 
	dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y); 
	MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C); 
	
	// Read C from device memory 
	cudaMemcpy(C.elements, d_C.elements, size, cudaMemcpyDeviceToHost); 
	
	// Free device memory 
	cudaFree(d_A.elements); 
	cudaFree(d_B.elements); 
	cudaFree(d_C.elements); } 

// Matrix multiplication kernel called by MatMul() 

__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) { 
	// Each thread computes one element of C 
	// by accumulating results into Cvalue 
	float Cvalue = 0; 
	int row = blockIdx.y * blockDim.y + threadIdx.y; 
	int col = blockIdx.x * blockDim.x + threadIdx.x; 
	
	for (int e = 0; e < A.width; ++e) 
		Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col]; 
		C.elements[row * C.width + col] = Cvalue; 
}

	.file	"tmpxft_002d7b82_00000000-6_kernel.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z41__device_stub__Z12MatMulKernel6MatrixS_S_R6MatrixS0_S0_
	.type	_Z41__device_stub__Z12MatMulKernel6MatrixS_S_R6MatrixS0_S0_, @function
_Z41__device_stub__Z12MatMulKernel6MatrixS_S_R6MatrixS0_S0_:
.LFB2053:
	.cfi_startproc
	endbr64
	subq	$104, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rcx
	movq	%rdi, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rsi, 72(%rsp)
	leaq	16(%rsp), %rdi
	leaq	28(%rsp), %rsi
	movq	%rdx, 80(%rsp)
	movq	%rsp, %rdx
	movq	%rax, 16(%rsp)
	movl	$1, 24(%rsp)
	movq	%rax, 28(%rsp)
	movl	$1, 36(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	8(%rsp)
	.cfi_def_cfa_offset 120
	leaq	_Z12MatMulKernel6MatrixS_S_(%rip), %rdi
	pushq	8(%rsp)
	.cfi_def_cfa_offset 128
	movq	44(%rsp), %rcx
	movl	52(%rsp), %r8d
	movq	32(%rsp), %rsi
	movl	40(%rsp), %edx
	leaq	80(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 120
	popq	%rdx
	.cfi_def_cfa_offset 112
.L2:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$104, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2053:
	.size	_Z41__device_stub__Z12MatMulKernel6MatrixS_S_R6MatrixS0_S0_, .-_Z41__device_stub__Z12MatMulKernel6MatrixS_S_R6MatrixS0_S0_
	.globl	_Z12MatMulKernel6MatrixS_S_
	.type	_Z12MatMulKernel6MatrixS_S_, @function
_Z12MatMulKernel6MatrixS_S_:
.LFB2054:
	.cfi_startproc
	endbr64
	subq	$56, %rsp
	.cfi_def_cfa_offset 64
	movq	%rdi, 32(%rsp)
	leaq	32(%rsp), %rdi
	movq	%rsi, 40(%rsp)
	leaq	16(%rsp), %rsi
	movq	%rdx, 16(%rsp)
	movq	%rsp, %rdx
	movq	%rcx, 24(%rsp)
	movq	%r8, (%rsp)
	movq	%r9, 8(%rsp)
	call	_Z41__device_stub__Z12MatMulKernel6MatrixS_S_R6MatrixS0_S0_
	addq	$56, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2054:
	.size	_Z12MatMulKernel6MatrixS_S_, .-_Z12MatMulKernel6MatrixS_S_
	.globl	_Z6MatMul6MatrixS_S_
	.type	_Z6MatMul6MatrixS_S_, @function
_Z6MatMul6MatrixS_S_:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rdi, %rax
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rsi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rcx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%r12, %r15
	movq	%r8, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	sarq	$32, %r15
	sarq	$32, %rbp
	subq	$168, %rsp
	.cfi_def_cfa_offset 224
	movq	%r9, 8(%rsp)
	movq	%rdi, %r9
	sarq	$32, %r9
	imull	%r9d, %eax
	movq	%fs:40, %rdx
	movq	%rdx, 152(%rsp)
	xorl	%edx, %edx
	movl	%edi, 56(%rsp)
	leaq	64(%rsp), %rdi
	movl	%r9d, 60(%rsp)
	movslq	%eax, %rbx
	movq	%r9, 16(%rsp)
	salq	$2, %rbx
	movq	%r8, 24(%rsp)
	movq	%rbx, %rsi
	call	cudaMalloc@PLT
	movq	64(%rsp), %rdi
	movq	%rbx, %rdx
	movq	%r14, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	%r15d, 76(%rsp)
	imull	%r12d, %r15d
	leaq	80(%rsp), %rdi
	movl	%r12d, 72(%rsp)
	movslq	%r15d, %r15
	salq	$2, %r15
	movq	%r15, %rsi
	call	cudaMalloc@PLT
	movq	80(%rsp), %rdi
	movq	%r15, %rdx
	movq	%r13, %rsi
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	24(%rsp), %r8
	movl	%ebp, 92(%rsp)
	leaq	96(%rsp), %rdi
	imull	%r8d, %ebp
	movl	%r8d, 88(%rsp)
	movslq	%ebp, %rbp
	salq	$2, %rbp
	movq	%rbp, %rsi
	call	cudaMalloc@PLT
	movq	16(%rsp), %r9
	movl	%r12d, %r10d
	xorl	%r8d, %r8d
	shrl	$4, %r10d
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$68719476752, %rdx
	shrl	$4, %r9d
	movl	%r10d, 44(%rsp)
	movl	%r9d, 48(%rsp)
	movq	44(%rsp), %rdi
	xorl	%r9d, %r9d
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L10
	movups	56(%rsp), %xmm0
	movups	72(%rsp), %xmm1
	leaq	136(%rsp), %rdx
	leaq	120(%rsp), %rsi
	movups	88(%rsp), %xmm2
	leaq	104(%rsp), %rdi
	movups	%xmm0, 104(%rsp)
	movups	%xmm1, 120(%rsp)
	movups	%xmm2, 136(%rsp)
	call	_Z41__device_stub__Z12MatMulKernel6MatrixS_S_R6MatrixS0_S0_
.L10:
	movq	96(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$2, %ecx
	movq	%rbp, %rdx
	call	cudaMemcpy@PLT
	movq	64(%rsp), %rdi
	call	cudaFree@PLT
	movq	80(%rsp), %rdi
	call	cudaFree@PLT
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$168, %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
.LFE2028:
	.size	_Z6MatMul6MatrixS_S_, .-_Z6MatMul6MatrixS_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"%f "
.LC3:
	.string	"\n"
	.section	.rodata
	.align 32
.LC0:
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1084227584
	.long	1086324736
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1084227584
	.long	1086324736
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.align 32
.LC1:
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1084227584
	.long	1086324736
	.long	1088421888
	.long	1090519040
	.long	1091567616
	.long	1092616192
	.long	1093664768
	.long	1094713344
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1084227584
	.long	1086324736
	.long	1088421888
	.long	1090519040
	.long	1091567616
	.long	1092616192
	.long	1093664768
	.long	1094713344
	.long	1065353216
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1084227584
	.long	1086324736
	.long	1088421888
	.long	1090519040
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	leaq	.LC0(%rip), %rsi
	movl	$16, %ecx
	leaq	.LC2(%rip), %r14
	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
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	subq	$208, %rsp
	.cfi_def_cfa_offset 256
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
	rep movsl
	leaq	72(%rsp), %rdi
	leaq	.LC1(%rip), %rsi
	movl	$32, %ecx
	rep movsl
	movl	$4, %esi
	movl	$32, %edi
	call	calloc@PLT
	movl	$2147483649, %edx
	leaq	8(%rsp), %rsi
	leaq	72(%rsp), %rcx
	salq	$3, %rdx
	movq	%rax, %r9
	movq	%rax, %rbp
	movabsq	$17179869188, %rdi
	movq	%rdx, %r8
	call	_Z6MatMul6MatrixS_S_
.L14:
	movq	%r12, %r13
	xorl	%ebx, %ebx
	salq	$4, %r13
	addq	%rbp, %r13
.L15:
	cvtss2sd	0(%r13,%rbx,4), %xmm0
	movq	%r14, %rsi
	movl	$2, %edi
	movb	$1, %al
	call	__printf_chk@PLT
	incq	%rbx
	cmpq	$8, %rbx
	jne	.L15
	leaq	.LC3(%rip), %rbx
	movl	$2, %edi
	xorl	%eax, %eax
	incq	%r12
	movq	%rbx, %rsi
	call	__printf_chk@PLT
	cmpq	$4, %r12
	jne	.L14
	xorl	%eax, %eax
	movq	%rbx, %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$208, %rsp
	.cfi_def_cfa_offset 48
	xorl	%eax, %eax
	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
.LFE2027:
	.size	main, .-main
	.section	.rodata.str1.1
.LC4:
	.string	"_Z12MatMulKernel6MatrixS_S_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2056:
	.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	.LC4(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z12MatMulKernel6MatrixS_S_(%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
.LFE2056:
	.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 <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <hip/hip_runtime.h>

#ifndef __HIPCC__ 
#define __HIPCC__
#endif


#include <hip/device_functions.h>
#include <hip/hip_runtime_api.h>

// Matrices are stored in row-major order: 
// M(row, col) = *(M.elements + row * M.width + col) 
typedef struct { 
	int width; 
	int height; 
	float* elements; 
} Matrix; 

// Thread block size 
#define BLOCK_SIZE 16 
// Forward declaration of the matrix multiplication kernel 
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix); 
void MatMul(const Matrix A, const Matrix B, Matrix C);

int main() {
	float A[16] = { 1, 2, 3, 4, 5, 6, 1, 2, 3, 4, 5, 6, 1, 2, 3, 4 };
	float B[32] = { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 1, 2, 3, 4, 5, 6, 7, 8 };
	int rowsA, colsA, rowsB, colsB;
	rowsA = 4;
	colsA = 4;
	rowsB = 4;
	colsB = 8;
	float *C = (float*)calloc(rowsA * colsB, sizeof(float));

	Matrix Am, Bm, Cm;
	Am.height = rowsA;
	Am.width = colsA;
	Am.elements = A;
	Bm.height = rowsB;
	Bm.width = colsB;
	Bm.elements = B;
	Cm.height = rowsA;
	Cm.width = colsB;
	Cm.elements = C;

	MatMul(Am, Bm, Cm);

	for (int i = 0; i < Cm.height; ++i){
		for (int j = 0; j < Cm.width; ++j){
			printf("%f ", Cm.elements[i * Cm.height + j]);
		}
		printf("\n");
	}
	printf("\n");
	//print_matrix(Cm, Cm.height, Cm.width);
	//C = matmul(A, rowsA, colsA, B, rowsB, colsB);

}

// Matrix multiplication - Host code 
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE 
void MatMul(const Matrix A, const Matrix B, Matrix C) { 
	
	// Load A and B to device memory 
	Matrix d_A; d_A.width = A.width; d_A.height = A.height; 
	size_t size = A.width * A.height * sizeof(float); 
	hipMalloc(&d_A.elements, size); 
	hipMemcpy(d_A.elements, A.elements, size, hipMemcpyHostToDevice); 
	
	Matrix d_B; d_B.width = B.width; d_B.height = B.height; 
	size = B.width * B.height * sizeof(float); 
	hipMalloc(&d_B.elements, size); 
	hipMemcpy(d_B.elements, B.elements, size, hipMemcpyHostToDevice); 
	
	// Allocate C in device memory 
	Matrix d_C; d_C.width = C.width; d_C.height = C.height; 
	size = C.width * C.height * sizeof(float); 
	hipMalloc(&d_C.elements, size); 
	
	// Invoke kernel 
	dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); 
	dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y); 
	MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C); 
	
	// Read C from device memory 
	hipMemcpy(C.elements, d_C.elements, size, hipMemcpyDeviceToHost); 
	
	// Free device memory 
	hipFree(d_A.elements); 
	hipFree(d_B.elements); 
	hipFree(d_C.elements); } 

// Matrix multiplication kernel called by MatMul() 

__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) { 
	// Each thread computes one element of C 
	// by accumulating results into Cvalue 
	float Cvalue = 0; 
	int row = blockIdx.y * blockDim.y + threadIdx.y; 
	int col = blockIdx.x * blockDim.x + threadIdx.x; 
	
	for (int e = 0; e < A.width; ++e) 
		Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col]; 
		C.elements[row * C.width + col] = Cvalue; 
}

	.text
	.file	"kernel.hip"
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI0_0:
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40400000                      # float 3
	.long	0x40800000                      # float 4
.LCPI0_1:
	.long	0x40a00000                      # float 5
	.long	0x40c00000                      # float 6
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
.LCPI0_2:
	.long	0x40400000                      # float 3
	.long	0x40800000                      # float 4
	.long	0x40a00000                      # float 5
	.long	0x40c00000                      # float 6
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	subq	$192, %rsp
	.cfi_def_cfa_offset 224
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movaps	.LCPI0_0(%rip), %xmm0           # xmm0 = [1.0E+0,2.0E+0,3.0E+0,4.0E+0]
	movq	%rsp, %r14
	movaps	%xmm0, (%r14)
	movaps	.LCPI0_1(%rip), %xmm1           # xmm1 = [5.0E+0,6.0E+0,1.0E+0,2.0E+0]
	movaps	%xmm1, 16(%r14)
	movaps	.LCPI0_2(%rip), %xmm1           # xmm1 = [3.0E+0,4.0E+0,5.0E+0,6.0E+0]
	movaps	%xmm1, 32(%r14)
	movaps	%xmm0, 48(%r14)
	leaq	64(%rsp), %r15
	movl	$16, %ecx
	movl	$.L__const.main.B, %esi
	movq	%r15, %rdi
	rep;movsq (%rsi), %es:(%rdi)
	movl	$32, %edi
	movl	$4, %esi
	callq	calloc
	movq	%rax, %rbx
	movabsq	$17179869192, %rdx              # imm = 0x400000008
	leaq	-4(%rdx), %rdi
	movq	%r14, %rsi
	movq	%r15, %rcx
	movq	%rdx, %r8
	movq	%rax, %r9
	callq	_Z6MatMul6MatrixS_S_
	xorl	%r14d, %r14d
.LBB0_1:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_2 Depth 2
	xorl	%r15d, %r15d
.LBB0_2:                                #   Parent Loop BB0_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rbx,%r15,4), %xmm0
	movl	$.L.str, %edi
	movb	$1, %al
	callq	printf
	incq	%r15
	cmpq	$8, %r15
	jne	.LBB0_2
# %bb.3:                                #   in Loop: Header=BB0_1 Depth=1
	movl	$10, %edi
	callq	putchar@PLT
	incq	%r14
	addq	$16, %rbx
	cmpq	$4, %r14
	jne	.LBB0_1
# %bb.4:
	movl	$10, %edi
	callq	putchar@PLT
	xorl	%eax, %eax
	addq	$192, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.globl	_Z6MatMul6MatrixS_S_            # -- Begin function _Z6MatMul6MatrixS_S_
	.type	_Z6MatMul6MatrixS_S_,@function
_Z6MatMul6MatrixS_S_:                   # @_Z6MatMul6MatrixS_S_
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, 48(%rsp)                   # 8-byte Spill
	movq	%r8, %r12
	movq	%r8, 32(%rsp)                   # 8-byte Spill
	movq	%rcx, 40(%rsp)                  # 8-byte Spill
	movq	%rdx, %rbp
	movq	%rsi, 24(%rsp)                  # 8-byte Spill
	movq	%rdi, %r15
	movq	%rdi, %rax
	shrq	$32, %rax
	movq	%rdx, %r14
	shrq	$32, %r14
	shrq	$32, %r12
	leaq	80(%rsp), %rbx
	movq	%rdi, -8(%rbx)
	imull	%r15d, %eax
	movslq	%eax, %r13
	shlq	$2, %r13
	movq	%rbx, %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	24(%rsp), %rsi                  # 8-byte Reload
	movq	%r13, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	64(%rsp), %rbx
	movq	%rbp, -8(%rbx)
	imull	%ebp, %r14d
	movslq	%r14d, %r13
	shlq	$2, %r13
	movq	%rbx, %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	movq	(%rbx), %rdi
	movq	40(%rsp), %rsi                  # 8-byte Reload
	movq	%r13, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	16(%rsp), %rdi
	movq	32(%rsp), %rax                  # 8-byte Reload
	movq	%rax, -8(%rdi)
	imull	%eax, %r12d
	movslq	%r12d, %r12
	shlq	$2, %r12
	movq	%r12, %rsi
	callq	hipMalloc
	shrl	$4, %ebp
	shrq	$4, %r15
	movabsq	$1152921500311879680, %rdi      # imm = 0xFFFFFFF00000000
	andq	%r15, %rdi
	orq	%rbp, %rdi
	movabsq	$68719476752, %rdx              # imm = 0x1000000010
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	72(%rsp), %rdi
	movq	80(%rsp), %rsi
	movq	56(%rsp), %rdx
	movq	64(%rsp), %rcx
	movq	8(%rsp), %r8
	movq	16(%rsp), %r9
	callq	_Z27__device_stub__MatMulKernel6MatrixS_S_
.LBB1_2:
	movq	16(%rsp), %rsi
	movq	48(%rsp), %rdi                  # 8-byte Reload
	movq	%r12, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movq	80(%rsp), %rdi
	callq	hipFree
	movq	64(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z6MatMul6MatrixS_S_, .Lfunc_end1-_Z6MatMul6MatrixS_S_
	.cfi_endproc
                                        # -- End function
	.globl	_Z27__device_stub__MatMulKernel6MatrixS_S_ # -- Begin function _Z27__device_stub__MatMulKernel6MatrixS_S_
	.type	_Z27__device_stub__MatMulKernel6MatrixS_S_,@function
_Z27__device_stub__MatMulKernel6MatrixS_S_: # @_Z27__device_stub__MatMulKernel6MatrixS_S_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$128, %rsp
	.cfi_def_cfa_offset 176
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	80(%rsp), %rax
	movq	%rdi, (%rax)
	movq	%rsi, 8(%rax)
	leaq	64(%rsp), %rsi
	movq	%rdx, (%rsi)
	movq	%rcx, 8(%rsi)
	leaq	48(%rsp), %rcx
	movq	%r8, (%rcx)
	movq	%r9, 8(%rcx)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rsi, 8(%rbx)
	movq	%rcx, 16(%rbx)
	leaq	32(%rsp), %r14
	leaq	16(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%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	$_Z12MatMulKernel6MatrixS_S_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$144, %rsp
	.cfi_adjust_cfa_offset -144
	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	_Z27__device_stub__MatMulKernel6MatrixS_S_, .Lfunc_end2-_Z27__device_stub__MatMulKernel6MatrixS_S_
	.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	.LBB3_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB3_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z12MatMulKernel6MatrixS_S_, %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_end3:
	.size	__hip_module_ctor, .Lfunc_end3-__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	.LBB4_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
.LBB4_2:
	retq
.Lfunc_end4:
	.size	__hip_module_dtor, .Lfunc_end4-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	.L__const.main.B,@object        # @__const.main.B
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
.L__const.main.B:
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40400000                      # float 3
	.long	0x40800000                      # float 4
	.long	0x40a00000                      # float 5
	.long	0x40c00000                      # float 6
	.long	0x40e00000                      # float 7
	.long	0x41000000                      # float 8
	.long	0x41100000                      # float 9
	.long	0x41200000                      # float 10
	.long	0x41300000                      # float 11
	.long	0x41400000                      # float 12
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40400000                      # float 3
	.long	0x40800000                      # float 4
	.long	0x40a00000                      # float 5
	.long	0x40c00000                      # float 6
	.long	0x40e00000                      # float 7
	.long	0x41000000                      # float 8
	.long	0x41100000                      # float 9
	.long	0x41200000                      # float 10
	.long	0x41300000                      # float 11
	.long	0x41400000                      # float 12
	.long	0x3f800000                      # float 1
	.long	0x40000000                      # float 2
	.long	0x40400000                      # float 3
	.long	0x40800000                      # float 4
	.long	0x40a00000                      # float 5
	.long	0x40c00000                      # float 6
	.long	0x40e00000                      # float 7
	.long	0x41000000                      # float 8
	.size	.L__const.main.B, 128

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

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

	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"_Z12MatMulKernel6MatrixS_S_"
	.size	.L__unnamed_1, 28

	.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 _Z27__device_stub__MatMulKernel6MatrixS_S_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z12MatMulKernel6MatrixS_S_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z12MatMulKernel6MatrixS_S_ ; -- Begin function _Z12MatMulKernel6MatrixS_S_
	.globl	_Z12MatMulKernel6MatrixS_S_
	.p2align	8
	.type	_Z12MatMulKernel6MatrixS_S_,@function
_Z12MatMulKernel6MatrixS_S_:            ; @_Z12MatMulKernel6MatrixS_S_
; %bb.0:
	s_clause 0x2
	s_load_b32 s4, s[0:1], 0x3c
	s_load_b32 s6, s[0:1], 0x0
	s_load_b64 s[2:3], s[0:1], 0x28
	v_bfe_u32 v2, v0, 10, 10
	v_and_b32_e32 v3, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s5, s4, 16
	s_and_b32 s4, s4, 0xffff
	s_delay_alu instid0(VALU_DEP_1)
	v_mad_u64_u32 v[0:1], null, s15, s5, v[2:3]
	v_mad_u64_u32 v[1:2], null, s14, s4, v[3:4]
	s_cmp_lt_i32 s6, 1
	s_cbranch_scc1 .LBB0_3
; %bb.1:                                ; %.lr.ph
	s_load_b64 s[8:9], s[0:1], 0x8
	s_delay_alu instid0(VALU_DEP_2)
	v_mul_lo_u32 v2, v0, s6
	s_clause 0x1
	s_load_b32 s7, s[0:1], 0x10
	s_load_b64 s[4:5], s[0:1], 0x18
	v_mov_b32_e32 v6, 0
	v_mov_b32_e32 v4, v1
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	v_lshlrev_b64 v[2:3], 2, v[2:3]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s8, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s9, v3, vcc_lo
.LBB0_2:                                ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v5, 31, v4
	s_add_i32 s6, s6, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_eq_u32 s6, 0
	v_lshlrev_b64 v[7:8], 2, v[4:5]
	v_add_nc_u32_e32 v4, s7, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, s4, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s5, v8, vcc_lo
	global_load_b32 v5, v[2:3], off
	global_load_b32 v7, v[7:8], off
	v_add_co_u32 v2, vcc_lo, v2, 4
	v_add_co_ci_u32_e32 v3, vcc_lo, 0, v3, vcc_lo
	s_waitcnt vmcnt(0)
	v_fmac_f32_e32 v6, v5, v7
	s_cbranch_scc0 .LBB0_2
	s_branch .LBB0_4
.LBB0_3:
	v_mov_b32_e32 v6, 0
.LBB0_4:                                ; %Flow29
	s_load_b32 s0, s[0:1], 0x20
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[2:3], null, v0, s0, v[1:2]
	v_ashrrev_i32_e32 v3, 31, v2
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[2:3]
	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_b32 v[0:1], v6, off
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z12MatMulKernel6MatrixS_S_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 304
		.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 9
		.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	_Z12MatMulKernel6MatrixS_S_, .Lfunc_end0-_Z12MatMulKernel6MatrixS_S_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 340
; NumSgprs: 18
; NumVgprs: 9
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 9
; 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:
      - .offset:         0
        .size:           16
        .value_kind:     by_value
      - .offset:         16
        .size:           16
        .value_kind:     by_value
      - .offset:         32
        .size:           16
        .value_kind:     by_value
      - .offset:         48
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         52
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         56
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         60
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         62
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         64
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         66
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         68
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         70
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         96
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         112
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 304
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z12MatMulKernel6MatrixS_S_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z12MatMulKernel6MatrixS_S_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     9
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
