//#include "cuda_runtime.h"
//#include "device_launch_parameters.h"
//#include <cuda.h>
//#include <fstream>
//#include <iostream>
//#include <stdio.h>
//#include <ctime>
//
//using namespace std;
//
//__global__ void matrixMultiply(float * A, float * B, float * C,
//	int numARows, int numAColumns,
//	int numBRows, int numBColumns,
//	int numCRows, int numCColumns) {
//		int tx = threadIdx.x + blockIdx.x*blockDim.x;
//		int ty = threadIdx.y + blockIdx.y*blockDim.y;
//		
//		float p = 0;
//
//		if (tx < numBColumns && ty < numARows){
//			for (int k = 0; k < numAColumns; k++){
//				p += A[ty*numAColumns + k]*B[k*numBColumns + tx];
//			}
//			C[ty*numBColumns + tx] = p;
//		}
//}
//
//void calc(char *file)
//{
//	float * hostA;
//	float * hostB;
//	float * hostC;
//	float * deviceA;
//	float * deviceB;
//	float * deviceC;
//	int numARows;
//	int numAColumns;
//	int numBRows;
//	int numBColumns;
//	int numCRows;
//	int numCColumns;
//
//	float * result;
//
//	 Charger le fichier d'entree
//	char * in0 = new char();
//	strcpy(in0, file);
//	strcat(in0, "/input0.raw");
//	ifstream fin0(in0);
//	fin0 >> numARows >> numAColumns;
//	hostA = (float*)malloc(numARows*numAColumns*sizeof(float));
//	for (int i = 0; i < numARows*numAColumns; i++){
//		fin0 >> hostA[i];
//	}
//	fin0.close();
//	
//	char * in1 = new char();
//	strcpy(in1, file);
//	strcat(in1, "/input1.raw");
//	ifstream fin1(in1);
//	fin1 >> numBRows >> numBColumns;
//	hostB = (float*)malloc(numBRows*numBColumns*sizeof(float));
//	for (int i = 0; i < numBRows*numBColumns; i++)
//		fin1 >> hostB[i];
//	fin1.close();
//
//	 Initialiser numCRows et numCColumns
//	numCRows = numARows;
//	numCColumns = numBColumns;
//	 Allouer hostC
//	hostC = (float*)malloc(numCRows*numCColumns*sizeof(float));
//
//	 Afficher les informations sur la matrice
//	 Allouer la memoire sur GPU
//	cudaMalloc((float**)&deviceA, numARows*numAColumns*sizeof(float));
//	cudaMalloc((float**)&deviceB, numBRows*numBColumns*sizeof(float));
//	cudaMalloc((float**)&deviceC, numCRows*numCColumns*sizeof(float));
//
//	 Copier la memoire sur le GPU
//	cudaMemcpy(deviceA, hostA, numARows*numAColumns*sizeof(float), cudaMemcpyHostToDevice);
//	cudaMemcpy(deviceB, hostB, numBRows*numBColumns*sizeof(float), cudaMemcpyHostToDevice);
//
//	 Initialise la grille et les dimensions de chaque bloc
//	int gridX = ceil((double)numBColumns/16.);
//	int gridY = ceil((double)numARows/16.);
//	dim3 dimGrid(gridX, gridY, 1);
//	dim3 dimBlock(16 , 16, 1);
//
//	 Execute le kernel
//	matrixMultiply<<<dimGrid , dimBlock>>>(deviceA, deviceB, deviceC, numARows, numAColumns, numBRows, numBColumns, numCRows, numCColumns);
//
//	cudaThreadSynchronize();
//
//	 Charge le resultat en memoire CPU
//	cudaMemcpy(hostC, deviceC, numCRows*numCColumns*sizeof(float), cudaMemcpyDeviceToHost);
//
//	TEST
//	char * out = new char();
//	strcpy(out, file);
//	strcat(out, "/output.raw");
//	ifstream fout(out) ;
//	fout >> numCRows >> numCColumns;
//	result = (float*)malloc(numCRows*numCColumns*sizeof(float));
//	for (int i = 0; i < numCRows*numCColumns; i++)
//		fout >> result[i];
//	fout.close();
//
//	for (int i = 0; i < numCRows*numCColumns; i++){
//		printf("%d %f \n", i, fabs(result[i]-hostC[i]));
//	}
//	 Libere la memoire
//	free(hostA);
//	free(hostB);
//	free(hostC);
//	cudaFree(deviceA);
//	cudaFree(deviceB);
//	cudaFree(deviceC);
//	free(result);
//
//	printf("\n%d %d\n%d %d\n%d %d\n", numARows, numAColumns, numBRows, numBColumns, numCRows, numCColumns);
//	printf("%d %d\n", gridY, gridX);
//}
//
//int main()
//{
//	clock_t tbegin, tend;
//	tbegin = clock();
//	calc("mp2_data/5");
//	system("pause");
//
//	calc("ex2_data2/0");
//	printf("0\n");
//
//	calc("ex2_data2/1");
//	printf("1\n");
//
//	calc("ex2_data2/2");
//	printf("2\n");
//
//	calc("ex2_data2/3");
//	printf("3\n");
//
//	calc("ex2_data2/4");
//	printf("4\n");
//
//	calc("ex2_data2/5");
//	printf("5\n");
//	
//	calc("ex2_data2/6");
//	printf("6\n");
//	
//	calc("ex2_data2/7");
//	printf("7\n");
//	
//	calc("ex2_data2/8");
//	printf("8\n");
//	
//	calc("ex2_data2/9");
//	printf("9\n");
//	tend = clock();
//
//	printf("%f\n", (float)(tend-tbegin)/CLOCKS_PER_SEC);
//	system("pause");
//	return 0;
//}

	.file	"tmpxft_002d6036_00000000-6_matrix-simple.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>
//#include "cuda_runtime.h"
//#include "device_launch_parameters.h"
//#include <cuda.h>
//#include <fstream>
//#include <iostream>
//#include <stdio.h>
//#include <ctime>
//
//using namespace std;
//
//__global__ void matrixMultiply(float * A, float * B, float * C,
//	int numARows, int numAColumns,
//	int numBRows, int numBColumns,
//	int numCRows, int numCColumns) {
//		int tx = threadIdx.x + blockIdx.x*blockDim.x;
//		int ty = threadIdx.y + blockIdx.y*blockDim.y;
//		
//		float p = 0;
//
//		if (tx < numBColumns && ty < numARows){
//			for (int k = 0; k < numAColumns; k++){
//				p += A[ty*numAColumns + k]*B[k*numBColumns + tx];
//			}
//			C[ty*numBColumns + tx] = p;
//		}
//}
//
//void calc(char *file)
//{
//	float * hostA;
//	float * hostB;
//	float * hostC;
//	float * deviceA;
//	float * deviceB;
//	float * deviceC;
//	int numARows;
//	int numAColumns;
//	int numBRows;
//	int numBColumns;
//	int numCRows;
//	int numCColumns;
//
//	float * result;
//
//	 Charger le fichier d'entree
//	char * in0 = new char();
//	strcpy(in0, file);
//	strcat(in0, "/input0.raw");
//	ifstream fin0(in0);
//	fin0 >> numARows >> numAColumns;
//	hostA = (float*)malloc(numARows*numAColumns*sizeof(float));
//	for (int i = 0; i < numARows*numAColumns; i++){
//		fin0 >> hostA[i];
//	}
//	fin0.close();
//	
//	char * in1 = new char();
//	strcpy(in1, file);
//	strcat(in1, "/input1.raw");
//	ifstream fin1(in1);
//	fin1 >> numBRows >> numBColumns;
//	hostB = (float*)malloc(numBRows*numBColumns*sizeof(float));
//	for (int i = 0; i < numBRows*numBColumns; i++)
//		fin1 >> hostB[i];
//	fin1.close();
//
//	 Initialiser numCRows et numCColumns
//	numCRows = numARows;
//	numCColumns = numBColumns;
//	 Allouer hostC
//	hostC = (float*)malloc(numCRows*numCColumns*sizeof(float));
//
//	 Afficher les informations sur la matrice
//	 Allouer la memoire sur GPU
//	cudaMalloc((float**)&deviceA, numARows*numAColumns*sizeof(float));
//	cudaMalloc((float**)&deviceB, numBRows*numBColumns*sizeof(float));
//	cudaMalloc((float**)&deviceC, numCRows*numCColumns*sizeof(float));
//
//	 Copier la memoire sur le GPU
//	cudaMemcpy(deviceA, hostA, numARows*numAColumns*sizeof(float), cudaMemcpyHostToDevice);
//	cudaMemcpy(deviceB, hostB, numBRows*numBColumns*sizeof(float), cudaMemcpyHostToDevice);
//
//	 Initialise la grille et les dimensions de chaque bloc
//	int gridX = ceil((double)numBColumns/16.);
//	int gridY = ceil((double)numARows/16.);
//	dim3 dimGrid(gridX, gridY, 1);
//	dim3 dimBlock(16 , 16, 1);
//
//	 Execute le kernel
//	matrixMultiply<<<dimGrid , dimBlock>>>(deviceA, deviceB, deviceC, numARows, numAColumns, numBRows, numBColumns, numCRows, numCColumns);
//
//	cudaThreadSynchronize();
//
//	 Charge le resultat en memoire CPU
//	cudaMemcpy(hostC, deviceC, numCRows*numCColumns*sizeof(float), cudaMemcpyDeviceToHost);
//
//	TEST
//	char * out = new char();
//	strcpy(out, file);
//	strcat(out, "/output.raw");
//	ifstream fout(out) ;
//	fout >> numCRows >> numCColumns;
//	result = (float*)malloc(numCRows*numCColumns*sizeof(float));
//	for (int i = 0; i < numCRows*numCColumns; i++)
//		fout >> result[i];
//	fout.close();
//
//	for (int i = 0; i < numCRows*numCColumns; i++){
//		printf("%d %f \n", i, fabs(result[i]-hostC[i]));
//	}
//	 Libere la memoire
//	free(hostA);
//	free(hostB);
//	free(hostC);
//	cudaFree(deviceA);
//	cudaFree(deviceB);
//	cudaFree(deviceC);
//	free(result);
//
//	printf("\n%d %d\n%d %d\n%d %d\n", numARows, numAColumns, numBRows, numBColumns, numCRows, numCColumns);
//	printf("%d %d\n", gridY, gridX);
//}
//
//int main()
//{
//	clock_t tbegin, tend;
//	tbegin = clock();
//	calc("mp2_data/5");
//	system("pause");
//
//	calc("ex2_data2/0");
//	printf("0\n");
//
//	calc("ex2_data2/1");
//	printf("1\n");
//
//	calc("ex2_data2/2");
//	printf("2\n");
//
//	calc("ex2_data2/3");
//	printf("3\n");
//
//	calc("ex2_data2/4");
//	printf("4\n");
//
//	calc("ex2_data2/5");
//	printf("5\n");
//	
//	calc("ex2_data2/6");
//	printf("6\n");
//	
//	calc("ex2_data2/7");
//	printf("7\n");
//	
//	calc("ex2_data2/8");
//	printf("8\n");
//	
//	calc("ex2_data2/9");
//	printf("9\n");
//	tend = clock();
//
//	printf("%f\n", (float)(tend-tbegin)/CLOCKS_PER_SEC);
//	system("pause");
//	return 0;
//}

	.text
	.file	"matrix-simple.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
