//#pragma comment (lib, "cublas.lib")
//#include "stdio.h"
//#include <cuda.h>
//using namespace std;
//#include <ctime>
//#include "cuda_runtime.h"
//#include "curand_kernel.h"
//#include "device_launch_parameters.h"
//#include <stdio.h>
//#include <stdlib.h>
//
//#include <string>
//#include <iomanip>
//#include <time.h>
//#include <iostream>
//#include <cmath>
//#include <math.h>
//
//#define TRAIN_NUM 60000
//#define TEST_NUM 10000
//#define ROW 4
//#define COL 4
//#define CONV_SIZE 2
//#define POOL_SIZE 1
//#define FC1_SIZE 45
//#define FC2_SIZE 10
//#define CONV_W_SIZE 3
//#define CONV_W_NUM 1
//
//__device__ float _alpha = 20;
//__device__ int _train_label[TRAIN_NUM];
//__device__ float _train_image[ROW][COL];
//__device__ float _conv_w[CONV_W_NUM][CONV_W_SIZE][CONV_W_SIZE];
//__device__ float _conv_z[CONV_W_NUM][CONV_SIZE][CONV_SIZE];
//
////__global__ void _input_conv(float _train_image[4][4], float _conv_w[1][3][3], float _conv_z[1][2][2])
//__global__ void _input_conv()
//{
//	int ix = threadIdx.x + blockDim.x * blockIdx.x;
//	int iy = threadIdx.y + blockDim.y * blockIdx.y;
//	int iz = threadIdx.z + blockDim.z * blockIdx.z;
//	if (ix < CONV_W_NUM && iy < CONV_SIZE && iz < CONV_SIZE)
//	{
//		_conv_z[ix][iy][iz] = 0;
//		// #pragma unroll
//		for (int l = 0;l < CONV_W_SIZE;l++)
//			for (int m = 0;m < CONV_W_SIZE;m++)
//				_conv_z[ix][iy][iz] += _train_image[iy + l][iz + m] * _conv_w[ix][l][m];
//		_conv_z[ix][iy][iz] += _alpha;
//		//_conv_a[ix][iy][iz] = _sigmoid(_conv_z[ix][iy][iz]);
//	}
//}
//int main() {
//	float train_image[ROW][COL] = { { 3, 1, 2, 4 },
//	{ 2, 4, 3, 1 },
//	{ 1, 5, 2, 3 },
//	{ 2, 3, 4, 1 } };
//	float conv_w[CONV_W_NUM][CONV_W_SIZE][CONV_W_SIZE] = { {
//	{1, 2, 3},
//	{4, 3, 1},
//	{1, 2, 4}}};
//
//	float** h_weight = new float* [3];
//			for (int i = 0;i < 3;i++) {
//				h_weight[i] = new float[3];
//			}
//	float conv_z[CONV_W_NUM][CONV_SIZE][CONV_SIZE];
//
//	float train_label[2] = { 3,2 };
//	//float* _train_image;
//	//float* _train_label;
//	//float* _conv_w;
//	//float* _conv_z;
//
//	//cudaMalloc(&_conv_w, CONV_W_NUM * CONV_W_SIZE * CONV_W_SIZE * sizeof(float));
//	//cudaMalloc(&_conv_z, CONV_W_NUM * CONV_SIZE * CONV_SIZE * sizeof(float));
//	//cudaMalloc(&_train_image, ROW*COL*sizeof(float));
//	//cudaMalloc(&_train_label, 2*sizeof(float));
//
//	float alpha = 10.0;
//	cudaMemcpyToSymbol(_alpha, &alpha, sizeof(float));
//	cudaMemcpyToSymbol(_train_image, train_image, ROW * COL * sizeof(float));
//	cudaMemcpyToSymbol(_conv_w, conv_w, CONV_W_NUM * CONV_W_SIZE * CONV_W_SIZE * sizeof(float));
//	//cudaMemcpy(_train_label, train_label, 2 * sizeof(float), cudaMemcpyHostToDevice);
//	//cudaMemcpy(_train_image, train_image, ROW * COL * sizeof(float), cudaMemcpyHostToDevice);
//	//cudaMemcpy(_conv_w, conv_w, CONV_W_NUM*CONV_W_SIZE*CONV_W_SIZE*sizeof(float), cudaMemcpyHostToDevice);
//	dim3 grid2(1, 2, 2);
//
//	//_input_conv << <1, grid2>> > ((float (*)[4])_train_image, (float (*)[3][3])_conv_w, (float (*)[2][2])_conv_z);
//	_input_conv << <1, grid2>> > ();
//	cudaMemcpyFromSymbol(&conv_z, _conv_z, CONV_W_NUM * CONV_SIZE * CONV_SIZE*sizeof(float));
//	for (int i = 0;i < CONV_SIZE;i++) {
//		for (int j = 0;j < CONV_SIZE;j++) {
//			cout << conv_z[0][i][j]<<" ";
//		}
//		cout << endl;
//	}
//	return 0;
//}
	.file	"tmpxft_00225805_00000000-6_testCNN.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>
//#pragma comment (lib, "cublas.lib")
//#include "stdio.h"
//#include <cuda.h>
//using namespace std;
//#include <ctime>
//#include "cuda_runtime.h"
//#include "curand_kernel.h"
//#include "device_launch_parameters.h"
//#include <stdio.h>
//#include <stdlib.h>
//
//#include <string>
//#include <iomanip>
//#include <time.h>
//#include <iostream>
//#include <cmath>
//#include <math.h>
//
//#define TRAIN_NUM 60000
//#define TEST_NUM 10000
//#define ROW 4
//#define COL 4
//#define CONV_SIZE 2
//#define POOL_SIZE 1
//#define FC1_SIZE 45
//#define FC2_SIZE 10
//#define CONV_W_SIZE 3
//#define CONV_W_NUM 1
//
//__device__ float _alpha = 20;
//__device__ int _train_label[TRAIN_NUM];
//__device__ float _train_image[ROW][COL];
//__device__ float _conv_w[CONV_W_NUM][CONV_W_SIZE][CONV_W_SIZE];
//__device__ float _conv_z[CONV_W_NUM][CONV_SIZE][CONV_SIZE];
//
////__global__ void _input_conv(float _train_image[4][4], float _conv_w[1][3][3], float _conv_z[1][2][2])
//__global__ void _input_conv()
//{
//	int ix = threadIdx.x + blockDim.x * blockIdx.x;
//	int iy = threadIdx.y + blockDim.y * blockIdx.y;
//	int iz = threadIdx.z + blockDim.z * blockIdx.z;
//	if (ix < CONV_W_NUM && iy < CONV_SIZE && iz < CONV_SIZE)
//	{
//		_conv_z[ix][iy][iz] = 0;
//		// #pragma unroll
//		for (int l = 0;l < CONV_W_SIZE;l++)
//			for (int m = 0;m < CONV_W_SIZE;m++)
//				_conv_z[ix][iy][iz] += _train_image[iy + l][iz + m] * _conv_w[ix][l][m];
//		_conv_z[ix][iy][iz] += _alpha;
//		//_conv_a[ix][iy][iz] = _sigmoid(_conv_z[ix][iy][iz]);
//	}
//}
//int main() {
//	float train_image[ROW][COL] = { { 3, 1, 2, 4 },
//	{ 2, 4, 3, 1 },
//	{ 1, 5, 2, 3 },
//	{ 2, 3, 4, 1 } };
//	float conv_w[CONV_W_NUM][CONV_W_SIZE][CONV_W_SIZE] = { {
//	{1, 2, 3},
//	{4, 3, 1},
//	{1, 2, 4}}};
//
//	float** h_weight = new float* [3];
//			for (int i = 0;i < 3;i++) {
//				h_weight[i] = new float[3];
//			}
//	float conv_z[CONV_W_NUM][CONV_SIZE][CONV_SIZE];
//
//	float train_label[2] = { 3,2 };
//	//float* _train_image;
//	//float* _train_label;
//	//float* _conv_w;
//	//float* _conv_z;
//
//	//cudaMalloc(&_conv_w, CONV_W_NUM * CONV_W_SIZE * CONV_W_SIZE * sizeof(float));
//	//cudaMalloc(&_conv_z, CONV_W_NUM * CONV_SIZE * CONV_SIZE * sizeof(float));
//	//cudaMalloc(&_train_image, ROW*COL*sizeof(float));
//	//cudaMalloc(&_train_label, 2*sizeof(float));
//
//	float alpha = 10.0;
//	cudaMemcpyToSymbol(_alpha, &alpha, sizeof(float));
//	cudaMemcpyToSymbol(_train_image, train_image, ROW * COL * sizeof(float));
//	cudaMemcpyToSymbol(_conv_w, conv_w, CONV_W_NUM * CONV_W_SIZE * CONV_W_SIZE * sizeof(float));
//	//cudaMemcpy(_train_label, train_label, 2 * sizeof(float), cudaMemcpyHostToDevice);
//	//cudaMemcpy(_train_image, train_image, ROW * COL * sizeof(float), cudaMemcpyHostToDevice);
//	//cudaMemcpy(_conv_w, conv_w, CONV_W_NUM*CONV_W_SIZE*CONV_W_SIZE*sizeof(float), cudaMemcpyHostToDevice);
//	dim3 grid2(1, 2, 2);
//
//	//_input_conv << <1, grid2>> > ((float (*)[4])_train_image, (float (*)[3][3])_conv_w, (float (*)[2][2])_conv_z);
//	_input_conv << <1, grid2>> > ();
//	cudaMemcpyFromSymbol(&conv_z, _conv_z, CONV_W_NUM * CONV_SIZE * CONV_SIZE*sizeof(float));
//	for (int i = 0;i < CONV_SIZE;i++) {
//		for (int j = 0;j < CONV_SIZE;j++) {
//			cout << conv_z[0][i][j]<<" ";
//		}
//		cout << endl;
//	}
//	return 0;
//}
	.text
	.file	"testCNN.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
