/*
   Copyright 2019 Cleuton Sampaio

   Licensed under the Apache License, Version 2.0 (the "License");
   you may not use this file except in compliance with the License.
   You may obtain a copy of the License at

     http://www.apache.org/licenses/LICENSE-2.0

   Unless required by applicable law or agreed to in writing, software
   distributed under the License is distributed on an "AS IS" BASIS,
   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
   See the License for the specific language governing permissions and
   limitations under the License.
*/
/* 
    Matrix multiplication sample using CUDA 
    this sample works for squared matrices!
*/

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include <stdio.h>

// CUDA Kernel function: 

__global__ void matmul(float *A, float* B, float *C, int size)
{

    // Row and Column indexes: 
    int row = blockIdx.y*blockDim.y+threadIdx.y;
    int col = blockIdx.x*blockDim.x+threadIdx.x;

    // Are they bellow the maximum?
    if (col < size && row < size) {
       float result = 0;
       for(int ix=0;ix<size;ix++) {
          result += A[row*size+ix]*B[ix*size+col];
       }
       C[row*size+col] = result;
    }

}

int main()
{
    // Matrices and constants
    int size = 3;
    int total = size*size;
    float cpu_A[] = {-1,2,4,0,5,3,6,2,1};
    float cpu_B[] = {3,0,2,3,4,5,4,7,2};
    float cpu_C[total];

    // Allocate device memory:
    float* gpu_A;
    float* gpu_B;
    float* gpu_C;
    int msize = total * sizeof(float);
    cudaMalloc((void**)&gpu_A, msize);
    cudaMemcpy(gpu_A,cpu_A,msize,cudaMemcpyHostToDevice);
    cudaMalloc((void**)&gpu_B, msize);
    cudaMemcpy(gpu_B,cpu_B,msize,cudaMemcpyHostToDevice);
    cudaMalloc((void**)&gpu_C,msize);

    // Blocks & grids:
    dim3 blocks(size,size);
    dim3 grid(1,1);

    // Call the kernel:
    matmul<<<grid,blocks>>>(gpu_A,gpu_B,gpu_C,size);

    // Get the result Matrix:
    cudaMemcpy(cpu_C,gpu_C,msize,cudaMemcpyDeviceToHost);
    std::cout << cpu_C[0] << '\t' << cpu_C[1] << '\t' << cpu_C[2] << std::endl
              << cpu_C[3] << '\t' << cpu_C[4] << '\t' << cpu_C[5] << std::endl
              << cpu_C[6] << '\t' << cpu_C[7] << '\t' << cpu_C[8] << std::endl;

    //Free device matrices
    cudaFree(gpu_A);
    cudaFree(gpu_B);
    cudaFree(gpu_C);
    
}

	.file	"tmpxft_0034a0ef_00000000-6_matmul.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3638:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3638:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z30__device_stub__Z6matmulPfS_S_iPfS_S_i
	.type	_Z30__device_stub__Z6matmulPfS_S_iPfS_S_i, @function
_Z30__device_stub__Z6matmulPfS_S_iPfS_S_i:
.LFB3660:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L2
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z6matmulPfS_S_i(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L2:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L4
	call	__stack_chk_fail@PLT
.L4:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3660:
	.size	_Z30__device_stub__Z6matmulPfS_S_iPfS_S_i, .-_Z30__device_stub__Z6matmulPfS_S_iPfS_S_i
	.globl	_Z6matmulPfS_S_i
	.type	_Z6matmulPfS_S_i, @function
_Z6matmulPfS_S_i:
.LFB3661:
	.cfi_startproc
	endbr64
	jmp	_Z30__device_stub__Z6matmulPfS_S_iPfS_S_i
	.cfi_endproc
.LFE3661:
	.size	_Z6matmulPfS_S_i, .-_Z6matmulPfS_S_i
	.section	.rodata
	.align 32
.LC0:
	.long	-1082130432
	.long	1073741824
	.long	1082130432
	.long	0
	.long	1084227584
	.long	1077936128
	.long	1086324736
	.long	1073741824
	.long	1065353216
	.align 32
.LC1:
	.long	1077936128
	.long	0
	.long	1073741824
	.long	1077936128
	.long	1082130432
	.long	1084227584
	.long	1082130432
	.long	1088421888
	.long	1073741824
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3635:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	leaq	.LC0(%rip), %rsi
	movl	$9, %ecx
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	60(%rsp), %rdi
	rep movsl
	leaq	96(%rsp), %rdi
	leaq	.LC1(%rip), %rsi
	movl	$9, %ecx
	rep movsl
	leaq	8(%rsp), %rdi
	movl	$36, %esi
	call	cudaMalloc@PLT
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movl	$36, %edx
	leaq	60(%rsp), %rsi
	call	cudaMemcpy@PLT
	leaq	16(%rsp), %rdi
	movl	$36, %esi
	call	cudaMalloc@PLT
	movq	16(%rsp), %rdi
	movl	$1, %ecx
	movl	$36, %edx
	leaq	96(%rsp), %rsi
	call	cudaMemcpy@PLT
	leaq	24(%rsp), %rdi
	movl	$36, %esi
	call	cudaMalloc@PLT
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movabsq	$12884901891, %rdx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L9
	movq	24(%rsp), %rdx
	movq	16(%rsp), %rsi
	movl	$3, %ecx
	movq	8(%rsp), %rdi
	call	_Z30__device_stub__Z6matmulPfS_S_iPfS_S_i
.L9:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movl	$36, %edx
	leaq	132(%rsp), %rdi
	call	cudaMemcpy@PLT
	leaq	_ZSt4cout(%rip), %rdi
	cvtss2sd	132(%rsp), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movl	$9, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	cvtss2sd	136(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movl	$9, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	cvtss2sd	140(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	cvtss2sd	144(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movl	$9, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	cvtss2sd	148(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movl	$9, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	cvtss2sd	152(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	cvtss2sd	156(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movl	$9, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	cvtss2sd	160(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movl	$9, %esi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c@PLT
	cvtss2sd	164(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	xorl	%eax, %eax
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3635:
	.size	main, .-main
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"_Z6matmulPfS_S_i"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3663:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC2(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z6matmulPfS_S_i(%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
.LFE3663:
	.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:

/*
   Copyright 2019 Cleuton Sampaio

   Licensed under the Apache License, Version 2.0 (the "License");
   you may not use this file except in compliance with the License.
   You may obtain a copy of the License at

     http://www.apache.org/licenses/LICENSE-2.0

   Unless required by applicable law or agreed to in writing, software
   distributed under the License is distributed on an "AS IS" BASIS,
   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
   See the License for the specific language governing permissions and
   limitations under the License.
*/
/* 
    Matrix multiplication sample using CUDA 
    this sample works for squared matrices!
*/

#include "hip/hip_runtime.h"

#include <iostream>
#include <stdio.h>

// CUDA Kernel function: 

__global__ void matmul(float *A, float* B, float *C, int size)
{

    // Row and Column indexes: 
    int row = blockIdx.y*blockDim.y+threadIdx.y;
    int col = blockIdx.x*blockDim.x+threadIdx.x;

    // Are they bellow the maximum?
    if (col < size && row < size) {
       float result = 0;
       for(int ix=0;ix<size;ix++) {
          result += A[row*size+ix]*B[ix*size+col];
       }
       C[row*size+col] = result;
    }

}

int main()
{
    // Matrices and constants
    int size = 3;
    int total = size*size;
    float cpu_A[] = {-1,2,4,0,5,3,6,2,1};
    float cpu_B[] = {3,0,2,3,4,5,4,7,2};
    float cpu_C[total];

    // Allocate device memory:
    float* gpu_A;
    float* gpu_B;
    float* gpu_C;
    int msize = total * sizeof(float);
    hipMalloc((void**)&gpu_A, msize);
    hipMemcpy(gpu_A,cpu_A,msize,hipMemcpyHostToDevice);
    hipMalloc((void**)&gpu_B, msize);
    hipMemcpy(gpu_B,cpu_B,msize,hipMemcpyHostToDevice);
    hipMalloc((void**)&gpu_C,msize);

    // Blocks & grids:
    dim3 blocks(size,size);
    dim3 grid(1,1);

    // Call the kernel:
    matmul<<<grid,blocks>>>(gpu_A,gpu_B,gpu_C,size);

    // Get the result Matrix:
    hipMemcpy(cpu_C,gpu_C,msize,hipMemcpyDeviceToHost);
    std::cout << cpu_C[0] << '\t' << cpu_C[1] << '\t' << cpu_C[2] << std::endl
              << cpu_C[3] << '\t' << cpu_C[4] << '\t' << cpu_C[5] << std::endl
              << cpu_C[6] << '\t' << cpu_C[7] << '\t' << cpu_C[8] << std::endl;

    //Free device matrices
    hipFree(gpu_A);
    hipFree(gpu_B);
    hipFree(gpu_C);
    
}

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

                                        # End of file scope inline assembly
	.globl	_Z21__device_stub__matmulPfS_S_i # -- Begin function _Z21__device_stub__matmulPfS_S_i
	.type	_Z21__device_stub__matmulPfS_S_i,@function
_Z21__device_stub__matmulPfS_S_i:       # @_Z21__device_stub__matmulPfS_S_i
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	4(%rsp), %rdx
	movl	%ecx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z6matmulPfS_S_i, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z21__device_stub__matmulPfS_S_i, .Lfunc_end0-_Z21__device_stub__matmulPfS_S_i
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI1_0:
	.long	0xbf800000                      # float -1
	.long	0x40000000                      # float 2
	.long	0x40800000                      # float 4
	.long	0x00000000                      # float 0
.LCPI1_1:
	.long	0x40a00000                      # float 5
	.long	0x40400000                      # float 3
	.long	0x40c00000                      # float 6
	.long	0x40000000                      # float 2
.LCPI1_2:
	.long	0x40400000                      # float 3
	.long	0x00000000                      # float 0
	.long	0x40000000                      # float 2
	.long	0x40400000                      # float 3
.LCPI1_3:
	.long	0x40800000                      # float 4
	.long	0x40a00000                      # float 5
	.long	0x40800000                      # float 4
	.long	0x40e00000                      # float 7
	.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	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$168, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsp, %r12
	movaps	.LCPI1_0(%rip), %xmm0           # xmm0 = [-1.0E+0,2.0E+0,4.0E+0,0.0E+0]
	leaq	128(%rsp), %r14
	movaps	%xmm0, (%r14)
	movaps	.LCPI1_1(%rip), %xmm0           # xmm0 = [5.0E+0,3.0E+0,6.0E+0,2.0E+0]
	movaps	%xmm0, 16(%r14)
	movl	$1065353216, 32(%r14)           # imm = 0x3F800000
	movaps	.LCPI1_2(%rip), %xmm0           # xmm0 = [3.0E+0,0.0E+0,2.0E+0,3.0E+0]
	leaq	80(%rsp), %rbx
	movaps	%xmm0, (%rbx)
	movaps	.LCPI1_3(%rip), %xmm0           # xmm0 = [4.0E+0,5.0E+0,4.0E+0,7.0E+0]
	movaps	%xmm0, 16(%rbx)
	movl	$1073741824, 32(%rbx)           # imm = 0x40000000
	leaq	24(%rsp), %r15
	movl	$36, %esi
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r15), %rdi
	movl	$36, %edx
	movq	%r14, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	16(%rsp), %r14
	movl	$36, %esi
	movq	%r14, %rdi
	callq	hipMalloc
	movq	(%r14), %rdi
	movl	$36, %edx
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	leaq	8(%rsp), %rdi
	movl	$36, %esi
	callq	hipMalloc
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	movabsq	$12884901891, %rdx              # imm = 0x300000003
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB1_2
# %bb.1:
	movq	24(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
	movl	$3, %ecx
	callq	_Z21__device_stub__matmulPfS_S_i
.LBB1_2:
	movq	8(%rsp), %rsi
	leaq	32(%rsp), %rbx
	movl	$36, %edx
	movq	%rbx, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorps	%xmm0, %xmm0
	cvtss2sd	(%rbx), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movb	$9, 7(%rsp)
	movq	(%rax), %rax
	movq	-24(%rax), %rax
	cmpq	$0, 16(%rbx,%rax)
	je	.LBB1_4
# %bb.3:
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rbx, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rax, %rbx
	jmp	.LBB1_5
.LBB1_4:
	movq	%rbx, %rdi
	movl	$9, %esi
	callq	_ZNSo3putEc
.LBB1_5:                                # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit
	xorps	%xmm0, %xmm0
	cvtss2sd	36(%rsp), %xmm0
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movb	$9, 7(%rsp)
	movq	(%rax), %rax
	movq	-24(%rax), %rax
	cmpq	$0, 16(%rbx,%rax)
	je	.LBB1_7
# %bb.6:
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rbx, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rax, %rbx
	jmp	.LBB1_8
.LBB1_7:
	movq	%rbx, %rdi
	movl	$9, %esi
	callq	_ZNSo3putEc
.LBB1_8:                                # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit22
	xorps	%xmm0, %xmm0
	cvtss2sd	40(%rsp), %xmm0
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	xorps	%xmm0, %xmm0
	cvtss2sd	44(%rsp), %xmm0
	movq	%rax, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movb	$9, 7(%rsp)
	movq	(%rax), %rax
	movq	-24(%rax), %rax
	cmpq	$0, 16(%rbx,%rax)
	je	.LBB1_10
# %bb.9:
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rbx, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rax, %rbx
	jmp	.LBB1_11
.LBB1_10:
	movq	%rbx, %rdi
	movl	$9, %esi
	callq	_ZNSo3putEc
.LBB1_11:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit25
	xorps	%xmm0, %xmm0
	cvtss2sd	48(%rsp), %xmm0
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movb	$9, 7(%rsp)
	movq	(%rax), %rax
	movq	-24(%rax), %rax
	cmpq	$0, 16(%rbx,%rax)
	je	.LBB1_13
# %bb.12:
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rbx, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rax, %rbx
	jmp	.LBB1_14
.LBB1_13:
	movq	%rbx, %rdi
	movl	$9, %esi
	callq	_ZNSo3putEc
.LBB1_14:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit28
	xorps	%xmm0, %xmm0
	cvtss2sd	52(%rsp), %xmm0
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	xorps	%xmm0, %xmm0
	cvtss2sd	56(%rsp), %xmm0
	movq	%rax, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movb	$9, 7(%rsp)
	movq	(%rax), %rax
	movq	-24(%rax), %rax
	cmpq	$0, 16(%rbx,%rax)
	je	.LBB1_16
# %bb.15:
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rbx, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rax, %rbx
	jmp	.LBB1_17
.LBB1_16:
	movq	%rbx, %rdi
	movl	$9, %esi
	callq	_ZNSo3putEc
.LBB1_17:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit31
	xorps	%xmm0, %xmm0
	cvtss2sd	60(%rsp), %xmm0
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movb	$9, 7(%rsp)
	movq	(%rax), %rax
	movq	-24(%rax), %rax
	cmpq	$0, 16(%rbx,%rax)
	je	.LBB1_19
# %bb.18:
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rbx, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rax, %rbx
	jmp	.LBB1_20
.LBB1_19:
	movq	%rbx, %rdi
	movl	$9, %esi
	callq	_ZNSo3putEc
.LBB1_20:                               # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit34
	xorps	%xmm0, %xmm0
	cvtss2sd	64(%rsp), %xmm0
	movq	%rbx, %rdi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %rbx
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%rbx, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%rbx, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	%r12, %rsp
	xorl	%eax, %eax
	addq	$168, %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_end1:
	.size	main, .Lfunc_end1-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	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z6matmulPfS_S_i, %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_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__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	.LBB3_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
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z6matmulPfS_S_i,@object        # @_Z6matmulPfS_S_i
	.section	.rodata,"a",@progbits
	.globl	_Z6matmulPfS_S_i
	.p2align	3, 0x0
_Z6matmulPfS_S_i:
	.quad	_Z21__device_stub__matmulPfS_S_i
	.size	_Z6matmulPfS_S_i, 8

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

	.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__matmulPfS_S_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z6matmulPfS_S_i
	.addrsig_sym _ZSt4cout
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z6matmulPfS_S_i        ; -- Begin function _Z6matmulPfS_S_i
	.globl	_Z6matmulPfS_S_i
	.p2align	8
	.type	_Z6matmulPfS_S_i,@function
_Z6matmulPfS_S_i:                       ; @_Z6matmulPfS_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x2c
	s_load_b32 s2, s[0:1], 0x18
	v_bfe_u32 v1, v0, 10, 10
	v_and_b32_e32 v4, 0x3ff, v0
	s_waitcnt lgkmcnt(0)
	s_lshr_b32 s4, s3, 16
	s_and_b32 s3, s3, 0xffff
	v_mad_u64_u32 v[2:3], null, s15, s4, v[1:2]
	v_mad_u64_u32 v[0:1], null, s14, s3, v[4:5]
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_max_i32_e32 v1, v2, v0
	v_cmpx_gt_i32_e64 s2, v1
	s_cbranch_execz .LBB0_6
; %bb.1:                                ; %.preheader
	s_clause 0x1
	s_load_b128 s[4:7], s[0:1], 0x0
	s_load_b64 s[0:1], s[0:1], 0x10
	v_mul_lo_u32 v1, v2, s2
	s_cmp_lt_i32 s2, 1
	s_cbranch_scc1 .LBB0_4
; %bb.2:                                ; %.lr.ph.preheader
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_3) | instid1(VALU_DEP_3)
	v_ashrrev_i32_e32 v2, 31, v1
	v_mov_b32_e32 v6, 0
	v_mov_b32_e32 v4, v0
	s_mov_b32 s3, s2
	v_lshlrev_b64 v[2:3], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v2, vcc_lo, s4, v2
	v_add_co_ci_u32_e32 v3, vcc_lo, s5, v3, vcc_lo
.LBB0_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v5, 31, v4
	s_add_i32 s3, s3, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_eq_u32 s3, 0
	v_lshlrev_b64 v[7:8], 2, v[4:5]
	v_add_nc_u32_e32 v4, s2, v4
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v7, vcc_lo, s6, v7
	v_add_co_ci_u32_e32 v8, vcc_lo, s7, v8, vcc_lo
	global_load_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_3
	s_branch .LBB0_5
.LBB0_4:
	v_mov_b32_e32 v6, 0
.LBB0_5:                                ; %Flow46
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v0, v1, v0
	v_ashrrev_i32_e32 v1, 31, v0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v0, vcc_lo, s0, v0
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v1, vcc_lo, s1, v1, vcc_lo
	global_store_b32 v[0:1], v6, off
.LBB0_6:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z6matmulPfS_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 1
		.amdhsa_next_free_vgpr 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	_Z6matmulPfS_S_i, .Lfunc_end0-_Z6matmulPfS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 332
; 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:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     by_value
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         36
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         40
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         44
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         46
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         48
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         50
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         52
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         54
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         88
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         96
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 288
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z6matmulPfS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z6matmulPfS_S_i.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
