/*

#ifndef __CUDACC__
#define __CUDACC__
#endif


#include "cuda_runtime.h"
#include "device_launch_parameters.h"


#include <stdio.h>
#include <conio.h>

const int TILE_WIDTH=2;
const int width=4;

__global__ void matrixmul(int *d_M,int *d_N,int *d_P)
{
	__shared__ int dS_M[TILE_WIDTH][TILE_WIDTH];
	__shared__ int dS_N[TILE_WIDTH][TILE_WIDTH];
	
	int by= blockIdx.y;
	int ty=threadIdx.y;
	int bx=blockIdx.x;
	int tx=threadIdx.x;

	int Row= by*TILE_WIDTH+ty;
	int Col= bx*TILE_WIDTH+tx;
	int pvalue=0;
	for(int m=0;m<(width/TILE_WIDTH);m++)
	{   dS_M[ty][tx]=d_M[Row*width + (m*TILE_WIDTH+tx)] ;
	    dS_N[ty][tx]=d_N[Col+(m*TILE_WIDTH+ty)*width] ;
		__syncthreads();
		for(int k=0;k<TILE_WIDTH;k++)
		{
			pvalue += dS_M[ty][k]*dS_N[k][tx] ;
			
		__syncthreads();
		}
		d_P [Row*width +Col]=pvalue;

	}

}
int main()
{
	int i,j,size;
	int h_M[width][width],h_N[width][width],h_P[width][width];
	int *d_M,*d_N,*d_P;
	
	printf("\n Enter according to width= %d \n ", width);
	
	for(i=0;i<width;i++)
	{
		for(j=0;j<width;j++)
			{ h_M[i][j]=2;
		      h_N[i][j]=2;
		     if(i==j)
		       {
				   h_N[i][j]= 1;
			     h_M[i][j]= 1;
			  }
   }
	}

	

	printf("\nh_M array : \n");
	for(i=0;i<width;i++)
	{   printf("\n");
		for(j=0;j<width;j++)
			printf("%d ",h_M[i][j]);
	}
	
	printf("\n\nh_N array : \n");
	for(i=0;i<width;i++)
	{   printf("\n");
		for(j=0;j<width;j++)
			printf("%d ",h_N[i][j]);
	}
	
	

	size=sizeof(int)*width*width;
	cudaMalloc((void**)&d_M,size);
	cudaMalloc((void**)&d_N,size);
	cudaMalloc((void**)&d_P,size);
	cudaMemcpy(d_M,h_M,size,cudaMemcpyHostToDevice);
	cudaMemcpy(d_N,h_N,size,cudaMemcpyHostToDevice);
	dim3 dimGrid((width/TILE_WIDTH),(width/TILE_WIDTH),1);
	dim3 dimBlock(TILE_WIDTH,TILE_WIDTH,1);

	matrixmul<<<dimGrid,dimBlock>>>(d_M,d_N,d_P);
	
	cudaMemcpy(h_P,d_P,size,cudaMemcpyDeviceToHost);
	cudaFree(d_M);
	cudaFree(d_N);
	cudaFree(d_P);

	printf("\n Resultant Array ( h_M * h_N ): \n");
	for(i=0;i<width;i++)
	{   printf("\n");
		for(j=0;j<width;j++)
			printf("%d ",h_P[i][j]);
	}
	
  getch();
  return 0;
}

*/
	.file	"tmpxft_0035170a_00000000-6_Matrix_Mul.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2033:
	.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>
/*

#ifndef __CUDACC__
#define __CUDACC__
#endif


#include "cuda_runtime.h"
#include "device_launch_parameters.h"


#include <stdio.h>
#include <conio.h>

const int TILE_WIDTH=2;
const int width=4;

__global__ void matrixmul(int *d_M,int *d_N,int *d_P)
{
	__shared__ int dS_M[TILE_WIDTH][TILE_WIDTH];
	__shared__ int dS_N[TILE_WIDTH][TILE_WIDTH];
	
	int by= blockIdx.y;
	int ty=threadIdx.y;
	int bx=blockIdx.x;
	int tx=threadIdx.x;

	int Row= by*TILE_WIDTH+ty;
	int Col= bx*TILE_WIDTH+tx;
	int pvalue=0;
	for(int m=0;m<(width/TILE_WIDTH);m++)
	{   dS_M[ty][tx]=d_M[Row*width + (m*TILE_WIDTH+tx)] ;
	    dS_N[ty][tx]=d_N[Col+(m*TILE_WIDTH+ty)*width] ;
		__syncthreads();
		for(int k=0;k<TILE_WIDTH;k++)
		{
			pvalue += dS_M[ty][k]*dS_N[k][tx] ;
			
		__syncthreads();
		}
		d_P [Row*width +Col]=pvalue;

	}

}
int main()
{
	int i,j,size;
	int h_M[width][width],h_N[width][width],h_P[width][width];
	int *d_M,*d_N,*d_P;
	
	printf("\n Enter according to width= %d \n ", width);
	
	for(i=0;i<width;i++)
	{
		for(j=0;j<width;j++)
			{ h_M[i][j]=2;
		      h_N[i][j]=2;
		     if(i==j)
		       {
				   h_N[i][j]= 1;
			     h_M[i][j]= 1;
			  }
   }
	}

	

	printf("\nh_M array : \n");
	for(i=0;i<width;i++)
	{   printf("\n");
		for(j=0;j<width;j++)
			printf("%d ",h_M[i][j]);
	}
	
	printf("\n\nh_N array : \n");
	for(i=0;i<width;i++)
	{   printf("\n");
		for(j=0;j<width;j++)
			printf("%d ",h_N[i][j]);
	}
	
	

	size=sizeof(int)*width*width;
	cudaMalloc((void**)&d_M,size);
	cudaMalloc((void**)&d_N,size);
	cudaMalloc((void**)&d_P,size);
	cudaMemcpy(d_M,h_M,size,cudaMemcpyHostToDevice);
	cudaMemcpy(d_N,h_N,size,cudaMemcpyHostToDevice);
	dim3 dimGrid((width/TILE_WIDTH),(width/TILE_WIDTH),1);
	dim3 dimBlock(TILE_WIDTH,TILE_WIDTH,1);

	matrixmul<<<dimGrid,dimBlock>>>(d_M,d_N,d_P);
	
	cudaMemcpy(h_P,d_P,size,cudaMemcpyDeviceToHost);
	cudaFree(d_M);
	cudaFree(d_N);
	cudaFree(d_P);

	printf("\n Resultant Array ( h_M * h_N ): \n");
	for(i=0;i<width;i++)
	{   printf("\n");
		for(j=0;j<width;j++)
			printf("%d ",h_P[i][j]);
	}
	
  getch();
  return 0;
}

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

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

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

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

	.end_amdgpu_metadata
