//#include <cudaDefs.h>
//
//cudaError_t error = cudaSuccess;
//cudaDeviceProp deviceProp = cudaDeviceProp();
//
///*
// ============================================================================
// Name        : VSB-PAII.cu
// Author      : Dave
// Version     :
// Copyright   : Your copyright notice
// Description : CUDA compute reciprocals
// ============================================================================
// */
//
// /*
//  ============================================================================
//  Name        : cuda1.cu
//  Author      : david
//  Version     :
//  Copyright   : Your copyright notice
//  Description : CUDA compute reciprocals
//  ============================================================================
//  */
//
//
//#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)
//
//constexpr unsigned int THREAD_PER_BLOCK = 256;
//constexpr unsigned int MEMBLOCK = 2;
//
//__global__ void VectorAdd(int* __restrict__ A, int* __restrict__ B, int* __restrict__ C, int M)
//{
//	unsigned int offset = threadIdx.x + blockIdx.x *  THREAD_PER_BLOCK;
//	unsigned int skip = gridDim.x * THREAD_PER_BLOCK;
//
//	while (offset < M)
//	{
//		C[offset] = A[offset] + B[offset];
//		offset += skip;
//	}
//}
//__global__ void VectorAddMN(int* __restrict__ A, int* __restrict__ B, int* __restrict__ C, int M, int N)
//{
//	const int i = blockIdx.x * blockDim.x + threadIdx.x;
//	const int j = blockIdx.y * blockDim.y + threadIdx.y;
//	if (i < M && j < N)
//	{
//		C[i * M + j] = A[i * M + j] + B[i * M + j];
//	}
//}
//
//__global__ void VectorAddMN_2(int *A, int *B, int *C, int M, int N)
//{
//	const int i = blockIdx.x * blockDim.x + threadIdx.x;
//	const int j = blockIdx.y * blockDim.y + threadIdx.y;
//	if (i < M && j < N)
//	{
//
//		for (int i = 0; i < M * N; ++i) {
//			int row = i / N;
//			int col = i % N;
//
//			C[row + col] = A[row + col] + B[row + col];
//		}
//	}
//}
//void CV1_1()
//{
//
//	//Time
//	cudaEvent_t startEvent, stopEvent;
//	float elapsedTime;
//	cudaEventCreate(&startEvent);
//	cudaEventCreate(&stopEvent);
//	cudaEventRecord(startEvent, 0);
//
//	
//	const unsigned int M = 10;
//	const unsigned int bytes = M * sizeof(int);
//
//	//Host allocate
//	int *A_vectorHost = (int*)malloc(bytes);
//	int *B_vectorHost = (int*)malloc(bytes);
//	int *C_vectorHost = (int*)malloc(bytes);
//
//
//	//Allocate the DEVICE memory to be able to copy data from HOST.
//	int *A_vectorDevice;
//	cudaMalloc((void**)&A_vectorDevice, bytes);
//
//	int *B_vectorDevice;
//	cudaMalloc((void**)&B_vectorDevice, bytes);
//
//	int *C_vectorDevice;
//	cudaMalloc((void**)&C_vectorDevice, bytes);
//
//
//	for (int i = 0; i < M; i++)
//	{
//		A_vectorHost[i] = i;
//		B_vectorHost[i] = i;
//		C_vectorHost[i] = 0;
//	}
//
//	cudaMemcpy(A_vectorDevice, A_vectorHost, bytes, cudaMemcpyHostToDevice);
//	cudaMemcpy(B_vectorDevice, B_vectorHost, bytes, cudaMemcpyHostToDevice);
//	cudaMemcpy(C_vectorDevice, C_vectorHost, bytes, cudaMemcpyHostToDevice);
//
//	dim3 dimBlock(THREAD_PER_BLOCK, THREAD_PER_BLOCK);
//	dim3 dimGrid(1, MEMBLOCK);
//
//
//	VectorAdd <<< dimGrid, dimBlock >>> (A_vectorDevice, B_vectorDevice, C_vectorDevice, M);
//
//	
//
//
//	cudaMemcpy(C_vectorHost, C_vectorDevice, bytes, cudaMemcpyDeviceToHost);
//
//	for (int i = 0; i < 10; i++)
//	{
//		printf("\n %d", C_vectorHost[i]);
//	}
//
//	free(A_vectorHost);
//	free(B_vectorHost);
//	free(C_vectorHost);
//
//	A_vectorHost = NULL;
//	B_vectorHost = NULL;
//	C_vectorHost = NULL;
//
//	cudaFree(A_vectorDevice);
//	cudaFree(B_vectorDevice);
//	cudaFree(C_vectorDevice);
//
//
//	//Time
//	cudaEventRecord(stopEvent, 0);
//	cudaEventSynchronize(stopEvent);
//	cudaEventElapsedTime(&elapsedTime, startEvent, stopEvent);
//
//	printf("Time to get device properties: %f ms", elapsedTime);
//	cudaEventDestroy(startEvent);
//	cudaEventDestroy(stopEvent);
//}
//
//void CV1_2()
//{
//	const unsigned int M = 10;
//	const unsigned int N = 10;
//
//	//Host allocate
//	int *A_Host = (int*)malloc(M*N * sizeof(int));
//	int *B_Host = (int*)malloc(M*N * sizeof(int));
//	int *C_Host = (int*)malloc(M*N * sizeof(int));
//
//	for (int i = 0; i < M; i++)
//	{
//		for (int j = 0; j < N; j++)
//		{
//			A_Host[i * M + j] = i * j;
//			B_Host[i * M + j] = i * j;
//
//		}
//	}
//
//	int *A_Device;
//	cudaMalloc(&A_Device, M*N * sizeof(int));
//
//	int *B_Device;
//	cudaMalloc(&B_Device, M*N * sizeof(int));
//	int *C_Device;
//	cudaMalloc(&C_Device, M*N * sizeof(int));
//
//	cudaMemcpy(A_Device, A_Host, M*N * sizeof(int), cudaMemcpyHostToDevice);
//	cudaMemcpy(B_Device, B_Host, M*N * sizeof(int), cudaMemcpyHostToDevice);
//	cudaMemcpy(C_Device, C_Host, M * sizeof(int), cudaMemcpyHostToDevice);
//
//	dim3 dimBlock(M, N);
//	dim3 dimGrid(1, 1);
//
//	VectorAddMN << <dimGrid, dimBlock >> > (A_Device, B_Device, C_Device, M, N);
//
//	cudaMemcpy(C_Host, C_Device, M*N * sizeof(int), cudaMemcpyDeviceToHost);
//
//
//	for (int i = 0; i < M; i++)
//	{
//		std::cout << i << "| ";
//		for (int j = 0; j < N; j++)
//		{
//			std::cout << C_Host[i * M + j] << "   ";
//		}
//		std::cout << std::endl;
//	}
//
//	free(A_Host);
//	free(B_Host);
//	free(C_Host);
//	cudaFree(A_Device);
//	cudaFree(B_Device);
//	cudaFree(C_Device);
//};
//
//
//void CV1_2_2()
//{
//	const unsigned int M = 10;
//	const unsigned int N = 10;
//
//	// M rows N cols
//	int** A_Host = new int*[M];
//	int** B_Host = new int*[M];
//	int** C_Host = new int*[M];
//
//
//	A_Host[0] = new int[M * N];
//	B_Host[0] = new int[M * N];
//	C_Host[0] = new int[M * N];
//
//	for (int i = 1; i < M; ++i) {
//		A_Host[i] = A_Host[i - 1] + N;
//		B_Host[i] = B_Host[i - 1] + N;
//		C_Host[i] = C_Host[i - 1] + N;
//	}
//
//	//Fill
//	for (int i = 0; i < M; ++i) {
//		for (int j = 0; j < N; ++j) {
//			A_Host[i][j] = i * j;
//			B_Host[i][j] = i * j;
//		}
//	}
//
//
//	int *A_Device;
//	cudaMalloc((void **)&A_Device, sizeof(int) * M * N);
//
//
//	int *B_Device;
//	cudaMalloc((void **)&B_Device, sizeof(int) * M * N);
//
//
//	int *C_Device;
//	cudaMalloc((void **)&C_Device, sizeof(int) * M * N);
//
//	cudaMemcpy(A_Device, A_Host[0], M*N * sizeof(int), cudaMemcpyHostToDevice);
//	cudaMemcpy(B_Device, B_Host[0], M*N * sizeof(int), cudaMemcpyHostToDevice);
//	cudaMemcpy(C_Device, C_Host[0], M*N * sizeof(int), cudaMemcpyHostToDevice);
//
//	dim3 dimBlock(M, N);
//	dim3 dimGrid(1, 1);
//
//	VectorAddMN_2 <<< dimGrid, dimBlock >>> (A_Device, B_Device, C_Device, M, N);
//
//	cudaMemcpy(C_Host, C_Device, M*N * sizeof(int), cudaMemcpyDeviceToHost);
//
//
//	for (int i = 0; i < M; i++)
//	{
//		std::cout << i << "| ";
//		for (int j = 0; j < N; j++)
//		{
//			std::cout << C_Host[i][j] << "  ";
//		}
//		std::cout << std::endl;
//	}
//
//	free(A_Host);
//	free(B_Host);
//	free(C_Host);
//	cudaFree(A_Device);
//	cudaFree(B_Device);
//	cudaFree(C_Device);
//};
////
////int main(int argc, char *argv[])
////{
////
////	CV1_1();
////
////}
////
//

	.file	"tmpxft_0021aba3_00000000-6_Runner_1.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 <cudaDefs.h>
//
//cudaError_t error = cudaSuccess;
//cudaDeviceProp deviceProp = cudaDeviceProp();
//
///*
// ============================================================================
// Name        : VSB-PAII.cu
// Author      : Dave
// Version     :
// Copyright   : Your copyright notice
// Description : CUDA compute reciprocals
// ============================================================================
// */
//
// /*
//  ============================================================================
//  Name        : cuda1.cu
//  Author      : david
//  Version     :
//  Copyright   : Your copyright notice
//  Description : CUDA compute reciprocals
//  ============================================================================
//  */
//
//
//#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)
//
//constexpr unsigned int THREAD_PER_BLOCK = 256;
//constexpr unsigned int MEMBLOCK = 2;
//
//__global__ void VectorAdd(int* __restrict__ A, int* __restrict__ B, int* __restrict__ C, int M)
//{
//	unsigned int offset = threadIdx.x + blockIdx.x *  THREAD_PER_BLOCK;
//	unsigned int skip = gridDim.x * THREAD_PER_BLOCK;
//
//	while (offset < M)
//	{
//		C[offset] = A[offset] + B[offset];
//		offset += skip;
//	}
//}
//__global__ void VectorAddMN(int* __restrict__ A, int* __restrict__ B, int* __restrict__ C, int M, int N)
//{
//	const int i = blockIdx.x * blockDim.x + threadIdx.x;
//	const int j = blockIdx.y * blockDim.y + threadIdx.y;
//	if (i < M && j < N)
//	{
//		C[i * M + j] = A[i * M + j] + B[i * M + j];
//	}
//}
//
//__global__ void VectorAddMN_2(int *A, int *B, int *C, int M, int N)
//{
//	const int i = blockIdx.x * blockDim.x + threadIdx.x;
//	const int j = blockIdx.y * blockDim.y + threadIdx.y;
//	if (i < M && j < N)
//	{
//
//		for (int i = 0; i < M * N; ++i) {
//			int row = i / N;
//			int col = i % N;
//
//			C[row + col] = A[row + col] + B[row + col];
//		}
//	}
//}
//void CV1_1()
//{
//
//	//Time
//	cudaEvent_t startEvent, stopEvent;
//	float elapsedTime;
//	cudaEventCreate(&startEvent);
//	cudaEventCreate(&stopEvent);
//	cudaEventRecord(startEvent, 0);
//
//	
//	const unsigned int M = 10;
//	const unsigned int bytes = M * sizeof(int);
//
//	//Host allocate
//	int *A_vectorHost = (int*)malloc(bytes);
//	int *B_vectorHost = (int*)malloc(bytes);
//	int *C_vectorHost = (int*)malloc(bytes);
//
//
//	//Allocate the DEVICE memory to be able to copy data from HOST.
//	int *A_vectorDevice;
//	cudaMalloc((void**)&A_vectorDevice, bytes);
//
//	int *B_vectorDevice;
//	cudaMalloc((void**)&B_vectorDevice, bytes);
//
//	int *C_vectorDevice;
//	cudaMalloc((void**)&C_vectorDevice, bytes);
//
//
//	for (int i = 0; i < M; i++)
//	{
//		A_vectorHost[i] = i;
//		B_vectorHost[i] = i;
//		C_vectorHost[i] = 0;
//	}
//
//	cudaMemcpy(A_vectorDevice, A_vectorHost, bytes, cudaMemcpyHostToDevice);
//	cudaMemcpy(B_vectorDevice, B_vectorHost, bytes, cudaMemcpyHostToDevice);
//	cudaMemcpy(C_vectorDevice, C_vectorHost, bytes, cudaMemcpyHostToDevice);
//
//	dim3 dimBlock(THREAD_PER_BLOCK, THREAD_PER_BLOCK);
//	dim3 dimGrid(1, MEMBLOCK);
//
//
//	VectorAdd <<< dimGrid, dimBlock >>> (A_vectorDevice, B_vectorDevice, C_vectorDevice, M);
//
//	
//
//
//	cudaMemcpy(C_vectorHost, C_vectorDevice, bytes, cudaMemcpyDeviceToHost);
//
//	for (int i = 0; i < 10; i++)
//	{
//		printf("\n %d", C_vectorHost[i]);
//	}
//
//	free(A_vectorHost);
//	free(B_vectorHost);
//	free(C_vectorHost);
//
//	A_vectorHost = NULL;
//	B_vectorHost = NULL;
//	C_vectorHost = NULL;
//
//	cudaFree(A_vectorDevice);
//	cudaFree(B_vectorDevice);
//	cudaFree(C_vectorDevice);
//
//
//	//Time
//	cudaEventRecord(stopEvent, 0);
//	cudaEventSynchronize(stopEvent);
//	cudaEventElapsedTime(&elapsedTime, startEvent, stopEvent);
//
//	printf("Time to get device properties: %f ms", elapsedTime);
//	cudaEventDestroy(startEvent);
//	cudaEventDestroy(stopEvent);
//}
//
//void CV1_2()
//{
//	const unsigned int M = 10;
//	const unsigned int N = 10;
//
//	//Host allocate
//	int *A_Host = (int*)malloc(M*N * sizeof(int));
//	int *B_Host = (int*)malloc(M*N * sizeof(int));
//	int *C_Host = (int*)malloc(M*N * sizeof(int));
//
//	for (int i = 0; i < M; i++)
//	{
//		for (int j = 0; j < N; j++)
//		{
//			A_Host[i * M + j] = i * j;
//			B_Host[i * M + j] = i * j;
//
//		}
//	}
//
//	int *A_Device;
//	cudaMalloc(&A_Device, M*N * sizeof(int));
//
//	int *B_Device;
//	cudaMalloc(&B_Device, M*N * sizeof(int));
//	int *C_Device;
//	cudaMalloc(&C_Device, M*N * sizeof(int));
//
//	cudaMemcpy(A_Device, A_Host, M*N * sizeof(int), cudaMemcpyHostToDevice);
//	cudaMemcpy(B_Device, B_Host, M*N * sizeof(int), cudaMemcpyHostToDevice);
//	cudaMemcpy(C_Device, C_Host, M * sizeof(int), cudaMemcpyHostToDevice);
//
//	dim3 dimBlock(M, N);
//	dim3 dimGrid(1, 1);
//
//	VectorAddMN << <dimGrid, dimBlock >> > (A_Device, B_Device, C_Device, M, N);
//
//	cudaMemcpy(C_Host, C_Device, M*N * sizeof(int), cudaMemcpyDeviceToHost);
//
//
//	for (int i = 0; i < M; i++)
//	{
//		std::cout << i << "| ";
//		for (int j = 0; j < N; j++)
//		{
//			std::cout << C_Host[i * M + j] << "   ";
//		}
//		std::cout << std::endl;
//	}
//
//	free(A_Host);
//	free(B_Host);
//	free(C_Host);
//	cudaFree(A_Device);
//	cudaFree(B_Device);
//	cudaFree(C_Device);
//};
//
//
//void CV1_2_2()
//{
//	const unsigned int M = 10;
//	const unsigned int N = 10;
//
//	// M rows N cols
//	int** A_Host = new int*[M];
//	int** B_Host = new int*[M];
//	int** C_Host = new int*[M];
//
//
//	A_Host[0] = new int[M * N];
//	B_Host[0] = new int[M * N];
//	C_Host[0] = new int[M * N];
//
//	for (int i = 1; i < M; ++i) {
//		A_Host[i] = A_Host[i - 1] + N;
//		B_Host[i] = B_Host[i - 1] + N;
//		C_Host[i] = C_Host[i - 1] + N;
//	}
//
//	//Fill
//	for (int i = 0; i < M; ++i) {
//		for (int j = 0; j < N; ++j) {
//			A_Host[i][j] = i * j;
//			B_Host[i][j] = i * j;
//		}
//	}
//
//
//	int *A_Device;
//	cudaMalloc((void **)&A_Device, sizeof(int) * M * N);
//
//
//	int *B_Device;
//	cudaMalloc((void **)&B_Device, sizeof(int) * M * N);
//
//
//	int *C_Device;
//	cudaMalloc((void **)&C_Device, sizeof(int) * M * N);
//
//	cudaMemcpy(A_Device, A_Host[0], M*N * sizeof(int), cudaMemcpyHostToDevice);
//	cudaMemcpy(B_Device, B_Host[0], M*N * sizeof(int), cudaMemcpyHostToDevice);
//	cudaMemcpy(C_Device, C_Host[0], M*N * sizeof(int), cudaMemcpyHostToDevice);
//
//	dim3 dimBlock(M, N);
//	dim3 dimGrid(1, 1);
//
//	VectorAddMN_2 <<< dimGrid, dimBlock >>> (A_Device, B_Device, C_Device, M, N);
//
//	cudaMemcpy(C_Host, C_Device, M*N * sizeof(int), cudaMemcpyDeviceToHost);
//
//
//	for (int i = 0; i < M; i++)
//	{
//		std::cout << i << "| ";
//		for (int j = 0; j < N; j++)
//		{
//			std::cout << C_Host[i][j] << "  ";
//		}
//		std::cout << std::endl;
//	}
//
//	free(A_Host);
//	free(B_Host);
//	free(C_Host);
//	cudaFree(A_Device);
//	cudaFree(B_Device);
//	cudaFree(C_Device);
//};
////
////int main(int argc, char *argv[])
////{
////
////	CV1_1();
////
////}
////
//

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