#include<iostream>
#include<cstdlib>
#include<cmath>
#include<time.h>
using namespace std;


__global__ void matrixVectorMultiplication(int *a, int *b, int *c, int n)
{
    int row=threadIdx.x+blockDim.x*blockIdx.x;
    int sum=0;
   
    if(row<n){
        for(int j=0;j<n;j++)
        {
            sum=sum+a[(j*n)+row]*b[j];
        }
    c[row]=sum;
    }
}
int main()
{
    int *a,*b,*c;
    int *a_dev,*b_dev,*c_dev;
    int n=10;
    
    a=new int[n*n];
    b=new int[n];
    c=new int[n];
    int *d=new int[n];
    int size=n*sizeof(int);
    cudaMalloc(&a_dev,size*size);
    cudaMalloc(&b_dev,size);
    cudaMalloc(&c_dev,size);
    

    cout<<"\n\nMatrix is :\n\n";
    for(int i=0;i<n;i++)
    {
        for(int j=0;j<n;j++)
        {
            a[i*n+j]= i*n+j+1; //rand()%n;
            cout<<a[i*n+j]<<" ";
        }
        
        b[i]= i+1; //rand()%n;
        cout<<"\n";
       // d[i]=a[i]+b[i];
    }
    
    cout<<"\n\nVector is: \n\n";
    for(int i=0;i<n;i++)
        cout<<b[i]<<" ";
    cout<<"\n\n";
    cudaMemcpy(a_dev,a,size*size,cudaMemcpyHostToDevice);
    cudaMemcpy(b_dev,b,size,cudaMemcpyHostToDevice);
    
    
    dim3 threadsPerBlock(n, n);
    dim3 blocksPerGrid(1, 1);
    
    if(n*n>512){
        threadsPerBlock.x=512;
        threadsPerBlock.y=512;
        blocksPerGrid.x=ceil((double)n/(double)threadsPerBlock.x);
        blocksPerGrid.y=ceil((double)n/(double)threadsPerBlock.y);
    }
    
    matrixVectorMultiplication<<<n/256 +1,256>>>(a_dev,b_dev,c_dev,n);
    
    cudaMemcpy(c,c_dev,size,cudaMemcpyDeviceToHost);
    
    //CPU matrixVector multiplication
    clock_t t=clock();
    int sum=0;
    for(int row=0;row<n;row++)
    {
        sum=0;
        for(int col=0;col<n;col++)
        {
              sum=sum+a[col*n+row]*b[col];  
            
        }
      d[row]=sum;
    }
    t=clock()-t;
        cout<<"\nCPU Time Elapsed:  "<<((double)t);      //((double)t)/CLOCKS_PER_SEC;

    
    int error=0;
    cout<<"\n\n";
    for(int i=0;i<n;i++){
        error+=d[i]-c[i];
       cout<<" gpu "<<c[i]<<" CPU "<<d[i]<<endl;
    }
    
    cout<<"\nError : "<<error<<"\n\n";
    
    
    return 0;
}

/*
Output
==11960== NVPROF is profiling process 11960, command: ./a.out


Matrix is :

1 2 3 4 5 6 7 8 9 10 
11 12 13 14 15 16 17 18 19 20 
21 22 23 24 25 26 27 28 29 30 
31 32 33 34 35 36 37 38 39 40 
41 42 43 44 45 46 47 48 49 50 
51 52 53 54 55 56 57 58 59 60 
61 62 63 64 65 66 67 68 69 70 
71 72 73 74 75 76 77 78 79 80 
81 82 83 84 85 86 87 88 89 90 
91 92 93 94 95 96 97 98 99 100 


Vector is: 

1 2 3 4 5 6 7 8 9 10 


CPU Time Elapsed:  3

 gpu 3355 CPU 3355
 gpu 3410 CPU 3410
 gpu 3465 CPU 3465
 gpu 3520 CPU 3520
 gpu 3575 CPU 3575
 gpu 3630 CPU 3630
 gpu 3685 CPU 3685
 gpu 3740 CPU 3740
 gpu 3795 CPU 3795
 gpu 3850 CPU 3850

Error : 0

==11960== Profiling application: ./a.out
==11960== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   48.69%  4.1910us         1  4.1910us  4.1910us  4.1910us  matrixVectorMultiplication(int*, int*, int*, int)
                   30.86%  2.6560us         2  1.3280us  1.0880us  1.5680us  [CUDA memcpy HtoD]
                   20.45%  1.7600us         1  1.7600us  1.7600us  1.7600us  [CUDA memcpy DtoH]
      API calls:   99.71%  200.37ms         3  66.791ms  5.1320us  200.36ms  cudaMalloc
                    0.16%  313.26us        97  3.2290us     124ns  206.98us  cuDeviceGetAttribute
                    0.05%  104.47us         1  104.47us  104.47us  104.47us  cuDeviceTotalMem
                    0.03%  64.413us         1  64.413us  64.413us  64.413us  cuDeviceGetName
                    0.03%  62.632us         3  20.877us  19.062us  22.537us  cudaMemcpy
                    0.02%  32.072us         1  32.072us  32.072us  32.072us  cudaLaunchKernel
                    0.00%  4.1760us         1  4.1760us  4.1760us  4.1760us  cuDeviceGetPCIBusId
                    0.00%  2.1090us         3     703ns     121ns  1.7920us  cuDeviceGetCount
                    0.00%     988ns         2     494ns     137ns     851ns  cuDeviceGet
                    0.00%     210ns         1     210ns     210ns     210ns  cuDeviceGetUuid

*/
	.file	"tmpxft_003a55bc_00000000-6_vectorMatMul.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	_Z51__device_stub__Z26matrixVectorMultiplicationPiS_S_iPiS_S_i
	.type	_Z51__device_stub__Z26matrixVectorMultiplicationPiS_S_iPiS_S_i, @function
_Z51__device_stub__Z26matrixVectorMultiplicationPiS_S_iPiS_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	_Z26matrixVectorMultiplicationPiS_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	_Z51__device_stub__Z26matrixVectorMultiplicationPiS_S_iPiS_S_i, .-_Z51__device_stub__Z26matrixVectorMultiplicationPiS_S_iPiS_S_i
	.globl	_Z26matrixVectorMultiplicationPiS_S_i
	.type	_Z26matrixVectorMultiplicationPiS_S_i, @function
_Z26matrixVectorMultiplicationPiS_S_i:
.LFB3661:
	.cfi_startproc
	endbr64
	jmp	_Z51__device_stub__Z26matrixVectorMultiplicationPiS_S_iPiS_S_i
	.cfi_endproc
.LFE3661:
	.size	_Z26matrixVectorMultiplicationPiS_S_i, .-_Z26matrixVectorMultiplicationPiS_S_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"\n\nMatrix is :\n\n"
.LC1:
	.string	" "
.LC2:
	.string	"\n"
.LC3:
	.string	"\n\nVector is: \n\n"
.LC4:
	.string	"\n\n"
.LC5:
	.string	"\nCPU Time Elapsed:  "
.LC6:
	.string	" gpu "
.LC7:
	.string	" CPU "
.LC8:
	.string	"\nError : "
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3635:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$400, %edi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	xorl	%r12d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$104, %rsp
	.cfi_def_cfa_offset 160
	movq	%fs:40, %rax
	movq	%rax, 88(%rsp)
	xorl	%eax, %eax
	call	_Znam@PLT
	movl	$40, %edi
	movq	%rax, 8(%rsp)
	call	_Znam@PLT
	movl	$40, %edi
	movq	%rax, %rbx
	call	_Znam@PLT
	movl	$40, %edi
	movq	%rax, %rbp
	call	_Znam@PLT
	leaq	40(%rsp), %rdi
	movl	$1600, %esi
	movq	%rax, %r13
	call	cudaMalloc@PLT
	leaq	48(%rsp), %rdi
	movl	$40, %esi
	call	cudaMalloc@PLT
	leaq	56(%rsp), %rdi
	movl	$40, %esi
	call	cudaMalloc@PLT
	leaq	.LC0(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.L9:
	imulq	$40, %r12, %r14
	movq	8(%rsp), %rax
	movl	%r12d, 28(%rsp)
	imull	$10, %r12d, %r15d
	addq	%r14, %rax
	movl	$1, %r14d
	movq	%rax, 16(%rsp)
.L10:
	movq	16(%rsp), %rax
	leal	(%r15,%r14), %esi
	leaq	_ZSt4cout(%rip), %rdi
	movl	%esi, -4(%rax,%r14,4)
	incq	%r14
	call	_ZNSolsEi@PLT
	leaq	.LC1(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$11, %r14
	jne	.L10
	movl	28(%rsp), %r15d
	leaq	.LC2(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	incl	%r15d
	movl	%r15d, (%rbx,%r12,4)
	incq	%r12
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$10, %r12
	jne	.L9
	leaq	.LC3(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	xorl	%r14d, %r14d
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.L12:
	movl	(%rbx,%r14,4), %esi
	leaq	_ZSt4cout(%rip), %rdi
	incq	%r14
	call	_ZNSolsEi@PLT
	leaq	.LC1(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cmpq	$10, %r14
	jne	.L12
	leaq	.LC4(%rip), %r14
	leaq	_ZSt4cout(%rip), %rdi
	movq	%r14, %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	8(%rsp), %rsi
	movq	40(%rsp), %rdi
	movl	$1, %ecx
	movl	$1600, %edx
	call	cudaMemcpy@PLT
	movq	48(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$40, %edx
	call	cudaMemcpy@PLT
	movl	$16777217, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	salq	$8, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	movabsq	$4294967297, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	movq	56(%rsp), %rdx
	movq	48(%rsp), %rsi
	movl	$10, %ecx
	movq	40(%rsp), %rdi
	call	_Z51__device_stub__Z26matrixVectorMultiplicationPiS_S_iPiS_S_i
.L13:
	movq	56(%rsp), %rsi
	movl	$40, %edx
	movl	$2, %ecx
	movq	%rbp, %rdi
	call	cudaMemcpy@PLT
	call	clock@PLT
	xorl	%edx, %edx
	movq	%rax, %r15
.L14:
	movq	8(%rsp), %rax
	xorl	%ecx, %ecx
	leaq	(%rax,%rdx,4), %rsi
	xorl	%eax, %eax
.L15:
	movl	(%rsi), %edi
	imull	(%rbx,%rcx,4), %edi
	incq	%rcx
	addq	$40, %rsi
	addl	%edi, %eax
	cmpq	$10, %rcx
	jne	.L15
	movl	%eax, 0(%r13,%rdx,4)
	incq	%rdx
	cmpq	$10, %rdx
	jne	.L14
	call	clock@PLT
	leaq	.LC5(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	subq	%r15, %rax
	xorl	%r15d, %r15d
	movq	%rax, %rbx
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	cvtsi2sdq	%rbx, %xmm0
	xorl	%ebx, %ebx
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%r14, %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
.L17:
	movl	0(%r13,%r15), %r12d
	leaq	.LC6(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	movl	%r12d, %eax
	subl	0(%rbp,%r15), %eax
	addl	%eax, %ebx
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	0(%rbp,%r15), %esi
	addq	$4, %r15
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	leaq	.LC7(%rip), %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%r12d, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	cmpq	$40, %r15
	jne	.L17
	leaq	.LC8(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movl	%ebx, %esi
	movq	%rax, %rdi
	call	_ZNSolsEi@PLT
	movq	%r14, %rsi
	movq	%rax, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L18
	call	__stack_chk_fail@PLT
.L18:
	addq	$104, %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
.LFE3635:
	.size	main, .-main
	.section	.rodata.str1.1
.LC9:
	.string	"_Z26matrixVectorMultiplicationPiS_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	.LC9(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z26matrixVectorMultiplicationPiS_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:


#include <hip/hip_runtime.h>
#include<iostream>
#include<cstdlib>
#include<cmath>
#include<time.h>
using namespace std;


__global__ void matrixVectorMultiplication(int *a, int *b, int *c, int n)
{
    int row=threadIdx.x+blockDim.x*blockIdx.x;
    int sum=0;
   
    if(row<n){
        for(int j=0;j<n;j++)
        {
            sum=sum+a[(j*n)+row]*b[j];
        }
    c[row]=sum;
    }
}
int main()
{
    int *a,*b,*c;
    int *a_dev,*b_dev,*c_dev;
    int n=10;
    
    a=new int[n*n];
    b=new int[n];
    c=new int[n];
    int *d=new int[n];
    int size=n*sizeof(int);
    hipMalloc(&a_dev,size*size);
    hipMalloc(&b_dev,size);
    hipMalloc(&c_dev,size);
    

    cout<<"\n\nMatrix is :\n\n";
    for(int i=0;i<n;i++)
    {
        for(int j=0;j<n;j++)
        {
            a[i*n+j]= i*n+j+1; //rand()%n;
            cout<<a[i*n+j]<<" ";
        }
        
        b[i]= i+1; //rand()%n;
        cout<<"\n";
       // d[i]=a[i]+b[i];
    }
    
    cout<<"\n\nVector is: \n\n";
    for(int i=0;i<n;i++)
        cout<<b[i]<<" ";
    cout<<"\n\n";
    hipMemcpy(a_dev,a,size*size,hipMemcpyHostToDevice);
    hipMemcpy(b_dev,b,size,hipMemcpyHostToDevice);
    
    
    dim3 threadsPerBlock(n, n);
    dim3 blocksPerGrid(1, 1);
    
    if(n*n>512){
        threadsPerBlock.x=512;
        threadsPerBlock.y=512;
        blocksPerGrid.x=ceil((double)n/(double)threadsPerBlock.x);
        blocksPerGrid.y=ceil((double)n/(double)threadsPerBlock.y);
    }
    
    matrixVectorMultiplication<<<n/256 +1,256>>>(a_dev,b_dev,c_dev,n);
    
    hipMemcpy(c,c_dev,size,hipMemcpyDeviceToHost);
    
    //CPU matrixVector multiplication
    clock_t t=clock();
    int sum=0;
    for(int row=0;row<n;row++)
    {
        sum=0;
        for(int col=0;col<n;col++)
        {
              sum=sum+a[col*n+row]*b[col];  
            
        }
      d[row]=sum;
    }
    t=clock()-t;
        cout<<"\nCPU Time Elapsed:  "<<((double)t);      //((double)t)/CLOCKS_PER_SEC;

    
    int error=0;
    cout<<"\n\n";
    for(int i=0;i<n;i++){
        error+=d[i]-c[i];
       cout<<" gpu "<<c[i]<<" CPU "<<d[i]<<endl;
    }
    
    cout<<"\nError : "<<error<<"\n\n";
    
    
    return 0;
}

/*
Output
==11960== NVPROF is profiling process 11960, command: ./a.out


Matrix is :

1 2 3 4 5 6 7 8 9 10 
11 12 13 14 15 16 17 18 19 20 
21 22 23 24 25 26 27 28 29 30 
31 32 33 34 35 36 37 38 39 40 
41 42 43 44 45 46 47 48 49 50 
51 52 53 54 55 56 57 58 59 60 
61 62 63 64 65 66 67 68 69 70 
71 72 73 74 75 76 77 78 79 80 
81 82 83 84 85 86 87 88 89 90 
91 92 93 94 95 96 97 98 99 100 


Vector is: 

1 2 3 4 5 6 7 8 9 10 


CPU Time Elapsed:  3

 gpu 3355 CPU 3355
 gpu 3410 CPU 3410
 gpu 3465 CPU 3465
 gpu 3520 CPU 3520
 gpu 3575 CPU 3575
 gpu 3630 CPU 3630
 gpu 3685 CPU 3685
 gpu 3740 CPU 3740
 gpu 3795 CPU 3795
 gpu 3850 CPU 3850

Error : 0

==11960== Profiling application: ./a.out
==11960== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   48.69%  4.1910us         1  4.1910us  4.1910us  4.1910us  matrixVectorMultiplication(int*, int*, int*, int)
                   30.86%  2.6560us         2  1.3280us  1.0880us  1.5680us  [CUDA memcpy HtoD]
                   20.45%  1.7600us         1  1.7600us  1.7600us  1.7600us  [CUDA memcpy DtoH]
      API calls:   99.71%  200.37ms         3  66.791ms  5.1320us  200.36ms  cudaMalloc
                    0.16%  313.26us        97  3.2290us     124ns  206.98us  cuDeviceGetAttribute
                    0.05%  104.47us         1  104.47us  104.47us  104.47us  cuDeviceTotalMem
                    0.03%  64.413us         1  64.413us  64.413us  64.413us  cuDeviceGetName
                    0.03%  62.632us         3  20.877us  19.062us  22.537us  cudaMemcpy
                    0.02%  32.072us         1  32.072us  32.072us  32.072us  cudaLaunchKernel
                    0.00%  4.1760us         1  4.1760us  4.1760us  4.1760us  cuDeviceGetPCIBusId
                    0.00%  2.1090us         3     703ns     121ns  1.7920us  cuDeviceGetCount
                    0.00%     988ns         2     494ns     137ns     851ns  cuDeviceGet
                    0.00%     210ns         1     210ns     210ns     210ns  cuDeviceGetUuid

*/
	.text
	.file	"vectorMatMul.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z41__device_stub__matrixVectorMultiplicationPiS_S_i # -- Begin function _Z41__device_stub__matrixVectorMultiplicationPiS_S_i
	.type	_Z41__device_stub__matrixVectorMultiplicationPiS_S_i,@function
_Z41__device_stub__matrixVectorMultiplicationPiS_S_i: # @_Z41__device_stub__matrixVectorMultiplicationPiS_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	$_Z26matrixVectorMultiplicationPiS_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	_Z41__device_stub__matrixVectorMultiplicationPiS_S_i, .Lfunc_end0-_Z41__device_stub__matrixVectorMultiplicationPiS_S_i
	.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	$40, %rsp
	.cfi_def_cfa_offset 96
	.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	$400, %edi                      # imm = 0x190
	callq	_Znam
	movq	%rax, %r15
	movl	$40, %edi
	callq	_Znam
	movq	%rax, %r12
	movl	$40, %edi
	callq	_Znam
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movl	$40, %edi
	callq	_Znam
	movq	%rax, (%rsp)                    # 8-byte Spill
	leaq	24(%rsp), %rdi
	movl	$1600, %esi                     # imm = 0x640
	callq	hipMalloc
	leaq	16(%rsp), %rdi
	movl	$40, %esi
	callq	hipMalloc
	leaq	8(%rsp), %rdi
	movl	$40, %esi
	callq	hipMalloc
	movl	$_ZSt4cout, %edi
	movl	$.L.str, %esi
	movl	$15, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$1, %r13d
	xorl	%ebx, %ebx
	movq	%r15, %rbp
.LBB1_1:                                # %.preheader98
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_2 Depth 2
	xorl	%r14d, %r14d
.LBB1_2:                                #   Parent Loop BB1_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	leal	(%r14,%r13), %esi
	movl	%esi, (%rbp,%r14,4)
	movl	$_ZSt4cout, %edi
	callq	_ZNSolsEi
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%r14
	cmpq	$10, %r14
	jne	.LBB1_2
# %bb.3:                                #   in Loop: Header=BB1_1 Depth=1
	leaq	1(%rbx), %r14
	movl	%r14d, (%r12,%rbx,4)
	movl	$_ZSt4cout, %edi
	movl	$.L.str.2, %esi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	addq	$10, %r13
	addq	$40, %rbp
	movq	%r14, %rbx
	cmpq	$10, %r14
	jne	.LBB1_1
# %bb.4:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.3, %esi
	movl	$15, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorl	%ebx, %ebx
.LBB1_5:                                # =>This Inner Loop Header: Depth=1
	movl	(%r12,%rbx,4), %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSolsEi
	movl	$.L.str.1, %esi
	movl	$1, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	incq	%rbx
	cmpq	$10, %rbx
	jne	.LBB1_5
# %bb.6:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$2, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	24(%rsp), %rdi
	movl	$1600, %edx                     # imm = 0x640
	movq	%r15, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	16(%rsp), %rdi
	movl	$40, %edx
	movq	%r12, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967297, %rdi               # imm = 0x100000001
	leaq	255(%rdi), %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	movq	32(%rsp), %r14                  # 8-byte Reload
	jne	.LBB1_8
# %bb.7:
	movq	24(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
	movl	$10, %ecx
	callq	_Z41__device_stub__matrixVectorMultiplicationPiS_S_i
.LBB1_8:
	movq	8(%rsp), %rsi
	movl	$40, %edx
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	xorl	%ebx, %ebx
	callq	clock
	movq	%rax, %r13
.LBB1_9:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB1_10 Depth 2
	movq	%r15, %rax
	xorl	%edx, %edx
	xorl	%ecx, %ecx
.LBB1_10:                               #   Parent Loop BB1_9 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movl	(%r12,%rdx,4), %esi
	imull	(%rax), %esi
	addl	%esi, %ecx
	incq	%rdx
	addq	$40, %rax
	cmpq	$10, %rdx
	jne	.LBB1_10
# %bb.11:                               #   in Loop: Header=BB1_9 Depth=1
	movq	(%rsp), %rax                    # 8-byte Reload
	movl	%ecx, (%rax,%rbx,4)
	incq	%rbx
	addq	$4, %r15
	cmpq	$10, %rbx
	jne	.LBB1_9
# %bb.12:
	callq	clock
	movq	%rax, %r15
	subq	%r13, %r15
	movl	$_ZSt4cout, %edi
	movl	$.L.str.5, %esi
	movl	$20, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	cvtsi2sd	%r15, %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movl	$_ZSt4cout, %edi
	movl	$.L.str.4, %esi
	movl	$2, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorl	%r13d, %r13d
	xorl	%ebp, %ebp
	movq	(%rsp), %rbx                    # 8-byte Reload
.LBB1_13:                               # =>This Inner Loop Header: Depth=1
	movl	(%rbx,%r13,4), %r15d
	addl	%r15d, %ebp
	subl	(%r14,%r13,4), %ebp
	movl	$_ZSt4cout, %edi
	movl	$.L.str.6, %esi
	movl	$5, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	(%r14,%r13,4), %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSolsEi
	movq	%rax, %r12
	movl	$.L.str.7, %esi
	movl	$5, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%r12, %rdi
	movl	%r15d, %esi
	callq	_ZNSolsEi
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	incq	%r13
	cmpq	$10, %r13
	jne	.LBB1_13
# %bb.14:
	movl	$_ZSt4cout, %edi
	movl	$.L.str.8, %esi
	movl	$9, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movl	$_ZSt4cout, %edi
	movl	%ebp, %esi
	callq	_ZNSolsEi
	movl	$.L.str.4, %esi
	movl	$2, %edx
	movq	%rax, %rdi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	xorl	%eax, %eax
	addq	$40, %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	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	$_Z26matrixVectorMultiplicationPiS_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	_Z26matrixVectorMultiplicationPiS_S_i,@object # @_Z26matrixVectorMultiplicationPiS_S_i
	.section	.rodata,"a",@progbits
	.globl	_Z26matrixVectorMultiplicationPiS_S_i
	.p2align	3, 0x0
_Z26matrixVectorMultiplicationPiS_S_i:
	.quad	_Z41__device_stub__matrixVectorMultiplicationPiS_S_i
	.size	_Z26matrixVectorMultiplicationPiS_S_i, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"\n\nMatrix is :\n\n"
	.size	.L.str, 16

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

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"\n"
	.size	.L.str.2, 2

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"\n\nVector is: \n\n"
	.size	.L.str.3, 16

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"\n\n"
	.size	.L.str.4, 3

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"\nCPU Time Elapsed:  "
	.size	.L.str.5, 21

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	" gpu "
	.size	.L.str.6, 6

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

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"\nError : "
	.size	.L.str.8, 10

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z26matrixVectorMultiplicationPiS_S_i"
	.size	.L__unnamed_1, 38

	.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 _Z41__device_stub__matrixVectorMultiplicationPiS_S_i
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z26matrixVectorMultiplicationPiS_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	_Z26matrixVectorMultiplicationPiS_S_i ; -- Begin function _Z26matrixVectorMultiplicationPiS_S_i
	.globl	_Z26matrixVectorMultiplicationPiS_S_i
	.p2align	8
	.type	_Z26matrixVectorMultiplicationPiS_S_i,@function
_Z26matrixVectorMultiplicationPiS_S_i:  ; @_Z26matrixVectorMultiplicationPiS_S_i
; %bb.0:
	s_clause 0x1
	s_load_b32 s3, s[0:1], 0x2c
	s_load_b32 s2, s[0:1], 0x18
	s_waitcnt lgkmcnt(0)
	s_and_b32 s3, s3, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s3, v[0:1]
	s_mov_b32 s3, exec_lo
	v_cmpx_gt_i32_e64 s2, v1
	s_cbranch_execz .LBB0_5
; %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_mov_b32_e32 v0, 0
	s_cmp_lt_i32 s2, 1
	s_cbranch_scc1 .LBB0_4
; %bb.2:                                ; %.lr.ph.preheader
	v_mov_b32_e32 v2, v1
	s_mov_b32 s3, s2
.LBB0_3:                                ; %.lr.ph
                                        ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1)
	v_ashrrev_i32_e32 v3, 31, v2
	s_waitcnt lgkmcnt(0)
	s_load_b32 s8, s[6:7], 0x0
	s_add_i32 s3, s3, -1
	s_add_u32 s6, s6, 4
	s_addc_u32 s7, s7, 0
	v_lshlrev_b64 v[3:4], 2, v[2:3]
	v_add_nc_u32_e32 v2, s2, v2
	s_cmp_eq_u32 s3, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_u32 v3, vcc_lo, s4, v3
	v_add_co_ci_u32_e32 v4, vcc_lo, s5, v4, vcc_lo
	global_load_b32 v5, v[3:4], off
	s_waitcnt vmcnt(0) lgkmcnt(0)
	v_mad_u64_u32 v[3:4], null, s8, v5, v[0:1]
	s_delay_alu instid0(VALU_DEP_1)
	v_mov_b32_e32 v0, v3
	s_cbranch_scc0 .LBB0_3
.LBB0_4:                                ; %Flow35
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_lshlrev_b64 v[1:2], 2, v[1:2]
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v1, vcc_lo, s0, v1
	s_delay_alu instid0(VALU_DEP_2)
	v_add_co_ci_u32_e32 v2, vcc_lo, s1, v2, vcc_lo
	global_store_b32 v[1:2], v0, off
.LBB0_5:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z26matrixVectorMultiplicationPiS_S_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 288
		.amdhsa_user_sgpr_count 15
		.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 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 6
		.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	_Z26matrixVectorMultiplicationPiS_S_i, .Lfunc_end0-_Z26matrixVectorMultiplicationPiS_S_i
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 252
; NumSgprs: 18
; NumVgprs: 6
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 0
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 6
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.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:           _Z26matrixVectorMultiplicationPiS_S_i
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z26matrixVectorMultiplicationPiS_S_i.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     6
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
