#include "NeuralNetwork.cuh"

/**
  *	Creates a neural network with the specified number of layers
  *	and neurons
  *	Parameter layers: the number of layers in the neural network
  *	Parameter neurons: an array with the number of neurons for each layer
  *	Returns: a NeuralNet with the specified layers/neurons
  */

NeuralNet* createNeuralNet(int layers, int* neurons, activation** activations){
	// Seeds the randomizer
	srand(time(NULL));

	// Creates the neural net
	NeuralNet* nn;
	cudaMallocManaged(&nn, sizeof(NeuralNet));
	
	// Sets the attributes of the neural net
	nn->layers = layers;
	cudaMallocManaged(&nn->neurons, layers*sizeof(int));
	for(int layer=0; layer<layers; layer++){
		nn->neurons[layer] = neurons[layer];
	}

	// Allocates memory for the activation function enum
	cudaMallocManaged(&nn->activations, 
			(layers - 1) * sizeof(activation*));

	// Allocates memory for the weights/biases and assigns random values
	cudaMallocManaged(&nn->biases, (layers - 1) * sizeof(double*));
	cudaMallocManaged(&nn->weights, (layers - 1) * sizeof(double**));
	
	for(int layer = 0; layer < layers-1; layer++){

		cudaMallocManaged(&nn->activations[layer], 
				neurons[layer + 1] *sizeof(activation));

		cudaMallocManaged(&nn->biases[layer], 
				neurons[layer+1] * sizeof(double));
		cudaMallocManaged(&nn->weights[layer], 
				neurons[layer] * sizeof(double*));

		for(int neuron1 = 0; neuron1 < neurons[layer+1]; neuron1++){
			nn->biases[layer][neuron1] = double(rand())/RAND_MAX;
			nn->activations[layer][neuron1] = 
				activations[layer][neuron1];
		}

		for(int neuron1 = 0; neuron1 < neurons[layer]; neuron1++){
			cudaMallocManaged(&nn->weights[layer][neuron1], 
					neurons[layer+1] * sizeof(double));
			
			for(int neuron2 = 0; neuron2 < neurons[layer+1]; 
					neuron2++){
				nn->weights[layer][neuron1][neuron2] = 
					double(rand())/RAND_MAX;
			}
		}
	}

	return nn;
}

void freeNeuralNet(NeuralNet* nn){
	for(int layer = 0; layer < nn->layers-1; layer++){
		for(int neuron1=0; neuron1 < nn->neurons[layer]; neuron1++){
			cudaFree(nn->weights[layer][neuron1]);
		}
		cudaFree(nn->activations[layer]);
		cudaFree(nn->biases[layer]);
		cudaFree(nn->weights[layer]);
	}
	cudaFree(nn->neurons);
	cudaFree(nn->activations);
	cudaFree(nn->biases);
	cudaFree(nn->weights);
	cudaFree(nn);
}

void printNeuralNet(NeuralNet* nn){
	printf("Layers = %d\n", nn->layers);
	for(int layer = 0; layer < nn->layers; layer++){

		printf("\nLayer %d Neurons %d\n", layer, nn->neurons[layer]);

		// Prints the biases
		if(layer != 0){
			printf("Activations for this layer:\n");
			for(int neuron1 = 0; neuron1 < nn->neurons[layer]; 
					neuron1++){
				printf("%d\t", nn->activations[layer-1][neuron1]);
			}
			printf("\n");
			printf("Biases for this layer:\n");
			for(int neuron1 = 0; neuron1 < nn->neurons[layer]; 
					neuron1++){
				printf("%f\t", nn->biases[layer-1][neuron1]);
			}
			printf("\n");
		}
		else{
			printf("No biases in this layer\n");
		}

		// Prints the weights
		if(layer != nn->layers - 1){
			printf("Weights for this layer:\n");
			for(int neuron1 = 0; neuron1 < nn->neurons[layer];
					neuron1++){
				for(int neuron2 = 0; 
						neuron2 < nn->neurons[layer+1];
						neuron2++){
					printf("%f\t", nn->weights[layer]\
							[neuron1][neuron2]);
				}
				printf("\n");
			}
		}
	}
}


	.file	"tmpxft_002405ab_00000000-6_NeuralNetwork.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2032:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2032:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0, @function
_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0:
.LFB2178:
	.cfi_startproc
	movl	$1, %edx
	jmp	cudaMallocManaged@PLT
	.cfi_endproc
.LFE2178:
	.size	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0, .-_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0
	.globl	_Z15createNeuralNetiPiPP10activation
	.type	_Z15createNeuralNetiPiPP10activation, @function
_Z15createNeuralNetiPiPP10activation:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%edi, %ebx
	xorl	%edi, %edi
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	leaq	16(%rsp), %rdi
	movl	$40, %esi
	call	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	16(%rsp), %rdi
	movslq	%ebx, %rsi
	salq	$2, %rsi
	movl	%ebx, (%rdi)
	addq	$8, %rdi
	call	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	16(%rsp), %rdi
	xorl	%eax, %eax
.L4:
	cmpl	%eax, %ebx
	jle	.L17
	movl	0(%rbp,%rax,4), %ecx
	movq	8(%rdi), %rdx
	movl	%ecx, (%rdx,%rax,4)
	incq	%rax
	jmp	.L4
.L17:
	leal	-1(%rbx), %eax
	addq	$32, %rdi
	movslq	%eax, %rbx
	movl	%eax, 4(%rsp)
	salq	$3, %rbx
	movq	%rbx, %rsi
	call	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	16(%rsp), %rax
	movq	%rbx, %rsi
	leaq	16(%rax), %rdi
	call	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	16(%rsp), %rax
	movq	%rbx, %rsi
	xorl	%ebx, %ebx
	leaq	24(%rax), %rdi
	call	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0
.L6:
	movq	16(%rsp), %rax
	cmpl	%ebx, 4(%rsp)
	jle	.L18
	movq	32(%rax), %rdi
	movslq	4(%rbp,%rbx,4), %rsi
	leaq	0(,%rbx,8), %r12
	xorl	%r13d, %r13d
	salq	$2, %rsi
	addq	%r12, %rdi
	call	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	16(%rsp), %rax
	movslq	4(%rbp,%rbx,4), %rsi
	movq	16(%rax), %rdi
	salq	$3, %rsi
	addq	%r12, %rdi
	call	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0
	movq	16(%rsp), %rax
	movslq	0(%rbp,%rbx,4), %rsi
	movq	24(%rax), %rdi
	salq	$3, %rsi
	addq	%r12, %rdi
	call	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0
.L7:
	cmpl	%r13d, 4(%rbp,%rbx,4)
	jle	.L19
	call	rand@PLT
	movl	%eax, %ecx
	movq	16(%rsp), %rax
	cvtsi2sdl	%ecx, %xmm0
	divsd	.LC0(%rip), %xmm0
	movq	16(%rax), %rdx
	movq	32(%rax), %rax
	movq	(%rdx,%r12), %rdx
	movq	(%rax,%r12), %rax
	movsd	%xmm0, (%rdx,%r13,8)
	movq	(%r14,%rbx,8), %rdx
	movl	(%rdx,%r13,4), %edx
	movl	%edx, (%rax,%r13,4)
	incq	%r13
	jmp	.L7
.L19:
	xorl	%r15d, %r15d
.L9:
	movq	16(%rsp), %rax
	cmpl	%r15d, 0(%rbp,%rbx,4)
	jle	.L20
	movq	24(%rax), %rax
	movslq	4(%rbp,%rbx,4), %rsi
	leaq	0(,%r15,8), %r13
	movq	(%rax,%r12), %rdi
	salq	$3, %rsi
	addq	%r13, %rdi
	call	_Z17cudaMallocManagedI2nnE9cudaErrorPPT_mj.constprop.0.isra.0
	xorl	%ecx, %ecx
.L10:
	cmpl	%ecx, 4(%rbp,%rbx,4)
	jle	.L21
	movq	%rcx, 8(%rsp)
	call	rand@PLT
	movq	8(%rsp), %rcx
	movl	%eax, %esi
	movq	16(%rsp), %rax
	cvtsi2sdl	%esi, %xmm0
	divsd	.LC0(%rip), %xmm0
	movq	24(%rax), %rax
	movq	(%rax,%r12), %rax
	movq	(%rax,%r13), %rax
	movsd	%xmm0, (%rax,%rcx,8)
	incq	%rcx
	jmp	.L10
.L21:
	incq	%r15
	jmp	.L9
.L20:
	incq	%rbx
	jmp	.L6
.L18:
	movq	24(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	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
.LFE2027:
	.size	_Z15createNeuralNetiPiPP10activation, .-_Z15createNeuralNetiPiPP10activation
	.globl	_Z13freeNeuralNetP2nn
	.type	_Z13freeNeuralNetP2nn, @function
_Z13freeNeuralNetP2nn:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
.L23:
	movl	(%rbx), %eax
	decl	%eax
	cmpl	%ebp, %eax
	jle	.L25
	leaq	0(,%rbp,4), %r14
	leaq	0(,%rbp,8), %r13
	xorl	%r12d, %r12d
.L26:
	movq	8(%rbx), %rax
	cmpl	%r12d, (%rax,%r14)
	jle	.L28
	movq	24(%rbx), %rax
	movq	(%rax,%r13), %rax
	movq	(%rax,%r12,8), %rdi
	incq	%r12
	call	cudaFree@PLT
	jmp	.L26
.L28:
	movq	32(%rbx), %rax
	incq	%rbp
	movq	(%rax,%r13), %rdi
	call	cudaFree@PLT
	movq	16(%rbx), %rax
	movq	(%rax,%r13), %rdi
	call	cudaFree@PLT
	movq	24(%rbx), %rax
	movq	(%rax,%r13), %rdi
	call	cudaFree@PLT
	jmp	.L23
.L25:
	movq	8(%rbx), %rdi
	call	cudaFree@PLT
	movq	32(%rbx), %rdi
	call	cudaFree@PLT
	movq	16(%rbx), %rdi
	call	cudaFree@PLT
	movq	24(%rbx), %rdi
	call	cudaFree@PLT
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	cudaFree@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z13freeNeuralNetP2nn, .-_Z13freeNeuralNetP2nn
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Layers = %d\n"
.LC2:
	.string	"\nLayer %d Neurons %d\n"
.LC3:
	.string	"Activations for this layer:\n"
.LC4:
	.string	"%d\t"
.LC5:
	.string	"\n"
.LC6:
	.string	"Biases for this layer:\n"
.LC7:
	.string	"%f\t"
.LC8:
	.string	"No biases in this layer\n"
.LC9:
	.string	"Weights for this layer:\n"
	.text
	.globl	_Z14printNeuralNetP2nn
	.type	_Z14printNeuralNetP2nn, @function
_Z14printNeuralNetP2nn:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	leaq	.LC1(%rip), %rsi
	xorl	%eax, %eax
	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
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdi, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movl	(%rdi), %edx
	movl	$2, %edi
	call	__printf_chk@PLT
.L30:
	movl	%ebp, %r13d
	cmpl	%ebp, (%rbx)
	jle	.L44
	movq	8(%rbx), %rax
	leaq	.LC2(%rip), %rsi
	movl	%r13d, %edx
	movl	$2, %edi
	leaq	0(,%rbp,4), %r12
	movl	(%rax,%rbp,4), %ecx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %rsi
	testq	%rbp, %rbp
	je	.L43
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%r15d, %r15d
	call	__printf_chk@PLT
	leaq	-8(,%rbp,8), %r14
.L32:
	movq	8(%rbx), %rax
	cmpl	%r15d, (%rax,%r12)
	jle	.L45
	movq	32(%rbx), %rax
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	movq	(%rax,%r14), %rax
	movl	(%rax,%r15,4), %edx
	xorl	%eax, %eax
	incq	%r15
	call	__printf_chk@PLT
	jmp	.L32
.L45:
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%r15d, %r15d
	call	__printf_chk@PLT
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L34:
	movq	8(%rbx), %rax
	cmpl	%r15d, (%rax,%r12)
	jle	.L46
	movq	16(%rbx), %rax
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	movq	(%rax,%r14), %rax
	movsd	(%rax,%r15,8), %xmm0
	movb	$1, %al
	incq	%r15
	call	__printf_chk@PLT
	jmp	.L34
.L46:
	leaq	.LC5(%rip), %rsi
.L43:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	(%rbx), %eax
	decl	%eax
	cmpl	%r13d, %eax
	je	.L37
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%r13d, %r13d
	call	__printf_chk@PLT
	leaq	4(%r12), %rax
	leaq	0(,%rbp,8), %r14
	movq	%rax, (%rsp)
.L38:
	movq	8(%rbx), %rax
	cmpl	%r13d, (%rax,%r12)
	jle	.L37
	leaq	0(,%r13,8), %r8
	xorl	%r15d, %r15d
.L40:
	movq	8(%rbx), %rax
	movq	(%rsp), %rcx
	cmpl	%r15d, (%rax,%rcx)
	jle	.L47
	movq	24(%rbx), %rax
	movl	$2, %edi
	movq	%r8, 8(%rsp)
	leaq	.LC7(%rip), %rsi
	movq	(%rax,%r14), %rax
	movq	(%rax,%r8), %rax
	movsd	(%rax,%r15,8), %xmm0
	movb	$1, %al
	incq	%r15
	call	__printf_chk@PLT
	movq	8(%rsp), %r8
	jmp	.L40
.L47:
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incq	%r13
	call	__printf_chk@PLT
	jmp	.L38
.L37:
	incq	%rbp
	jmp	.L30
.L44:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	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
.LFE2029:
	.size	_Z14printNeuralNetP2nn, .-_Z14printNeuralNetP2nn
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2055:
	.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
.LFE2055:
	.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
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.long	-4194304
	.long	1105199103
	.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:

#ifndef NEURALNET_CUH
#define NEURALNET_CUH

#include <stdio.h>
#include <stdlib.h>
#include <time.h>

#define NUMBLOCKS 4
#define BLOCKSIZE 64

enum activation{LINEAR, BINARY_STEP, LOGISTIC, TANH, 
	RELU, LEAKYRELU};

typedef struct nn{
	int layers;
	int* neurons;
	double** biases;
	double*** weights;
	activation** activations;
} NeuralNet;

NeuralNet* createNeuralNet(int layers, int* neurons, activation** activations);

void printNeuralNet(NeuralNet* nn);

void freeNeuralNet(NeuralNet* nn);

#endif


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