#include "MemoryCell.cuh"

long long MemoryCell::n = 0;

MemoryCell::MemoryCell(int c) {
	nConnections = c;
	activationIn = 0; activationInPrime = 0;
	activationOut = 0; activationOutPrime = 0;
	state = 0; previousState = 0;
	feedback = 0; previousFeedback = 0;
	bias = 0;

	default_random_engine g(time(0) + (n++));
	normal_distribution<double> d(0, 1);

	cellFeedbackWeight = d(g);
	cellFeedbackPartial = 0;
	inputFeedbackPartial = 0;
	forgetFeedbackPartial = 0;

	cellDataWeight = (double *)malloc(sizeof(double) * c);
	cellDataPartial = (double *)malloc(sizeof(double) * c);
	forgetDataPartial = (double *)malloc(sizeof(double) * c);
	inputDataPartial = (double *)malloc(sizeof(double) * c);

	for (int i = 0; i < c; i++) {
		cellDataWeight[i] = (d(g));
		cellDataPartial[i] = (0);
		forgetDataPartial[i] = (0);
		inputDataPartial[i] = (0);
	}
}

MemoryCell::~MemoryCell() {}

__device__ double MemoryCell::activateIn(double data) {
	activationIn = activationFunction(data);
	activationInPrime = activationFunctionPrime(data);
	return activationIn;
}

__device__ double MemoryCell::activateOut(double data) {
	activationOut = activationFunction(data);
	activationOutPrime = activationFunctionPrime(data);
	return activationOut;
}

MemoryCell *MemoryCell::copyToGPU(MemoryCell *memory) {
	MemoryCell *memoryCell;
	cudaMalloc((void **)&memoryCell, (sizeof(MemoryCell)));
	cudaDeviceSynchronize();
	cudaMemcpy(memoryCell, memory, sizeof(MemoryCell), cudaMemcpyHostToDevice);
	cudaDeviceSynchronize();

	double *cdw, *idp, *fdp, *cdp;
	cudaMalloc((void **)&cdw, (sizeof(double) * memory->nConnections));
	cudaMalloc((void **)&idp, (sizeof(double) * memory->nConnections));
	cudaMalloc((void **)&fdp, (sizeof(double) * memory->nConnections));
	cudaMalloc((void **)&cdp, (sizeof(double) * memory->nConnections));
	cudaDeviceSynchronize();

	cudaMemcpy(cdw, memory->cellDataWeight, (sizeof(double) * memory->nConnections), cudaMemcpyHostToDevice);
	cudaMemcpy(idp, memory->inputDataPartial, (sizeof(double) * memory->nConnections), cudaMemcpyHostToDevice);
	cudaMemcpy(fdp, memory->forgetDataPartial, (sizeof(double) * memory->nConnections), cudaMemcpyHostToDevice);
	cudaMemcpy(cdp, memory->cellDataPartial, (sizeof(double) * memory->nConnections), cudaMemcpyHostToDevice);
	cudaDeviceSynchronize();

	cudaMemcpy(&(memoryCell->cellDataWeight), &cdw, sizeof(double *), cudaMemcpyHostToDevice);
	cudaMemcpy(&(memoryCell->inputDataPartial), &idp, sizeof(double *), cudaMemcpyHostToDevice);
	cudaMemcpy(&(memoryCell->forgetDataPartial), &fdp, sizeof(double *), cudaMemcpyHostToDevice);
	cudaMemcpy(&(memoryCell->cellDataPartial), &cdp, sizeof(double *), cudaMemcpyHostToDevice);
	cudaDeviceSynchronize();

	return memoryCell;
}

MemoryCell *MemoryCell::copyFromGPU(MemoryCell *memory) {

	MemoryCell *memoryCell;
	memoryCell = (MemoryCell *)malloc((sizeof(MemoryCell)));
	cudaDeviceSynchronize();
	cudaMemcpy(memoryCell, memory, sizeof(MemoryCell), cudaMemcpyDeviceToHost);
	cudaDeviceSynchronize();

	double *cdw, *idp, *fdp, *cdp;
	cdw = (double *)malloc(sizeof(double) * memoryCell->nConnections);
	idp = (double *)malloc(sizeof(double) * memoryCell->nConnections);
	fdp = (double *)malloc(sizeof(double) * memoryCell->nConnections);
	cdp = (double *)malloc(sizeof(double) * memoryCell->nConnections);

	cudaMemcpy(cdw, memoryCell->cellDataWeight, (sizeof(double) * memoryCell->nConnections), cudaMemcpyDeviceToHost);
	cudaMemcpy(idp, memoryCell->inputDataPartial, (sizeof(double) * memoryCell->nConnections), cudaMemcpyDeviceToHost);
	cudaMemcpy(fdp, memoryCell->forgetDataPartial, (sizeof(double) * memoryCell->nConnections), cudaMemcpyDeviceToHost);
	cudaMemcpy(cdp, memoryCell->cellDataPartial, (sizeof(double) * memoryCell->nConnections), cudaMemcpyDeviceToHost);
	cudaDeviceSynchronize();

	memcpy(&(memoryCell->cellDataWeight), &cdw, (sizeof(double *)));
	memcpy(&(memoryCell->inputDataPartial), &idp, (sizeof(double *)));
	memcpy(&(memoryCell->forgetDataPartial), &fdp, (sizeof(double *)));
	memcpy(&(memoryCell->cellDataPartial), &cdp, (sizeof(double *)));

	return memoryCell;
}


	.file	"tmpxft_00317760_00000000-6_MemoryCell.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.align 2
	.globl	_ZN10MemoryCellD2Ev
	.type	_ZN10MemoryCellD2Ev, @function
_ZN10MemoryCellD2Ev:
.LFB4723:
	.cfi_startproc
	endbr64
	leaq	16+_ZTV10MemoryCell(%rip), %rax
	movq	%rax, (%rdi)
	jmp	_ZN8BaseNodeD2Ev@PLT
	.cfi_endproc
.LFE4723:
	.size	_ZN10MemoryCellD2Ev, .-_ZN10MemoryCellD2Ev
	.globl	_ZN10MemoryCellD1Ev
	.set	_ZN10MemoryCellD1Ev,_ZN10MemoryCellD2Ev
	.align 2
	.globl	_ZN10MemoryCellD0Ev
	.type	_ZN10MemoryCellD0Ev, @function
_ZN10MemoryCellD0Ev:
.LFB4725:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movq	%rdi, 8(%rsp)
	call	_ZN10MemoryCellD1Ev
	movq	8(%rsp), %rdi
	movl	$152, %esi
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	jmp	_ZdlPvm@PLT
	.cfi_endproc
.LFE4725:
	.size	_ZN10MemoryCellD0Ev, .-_ZN10MemoryCellD0Ev
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4732:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4732:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.globl	_ZN10MemoryCell10activateInEd
	.type	_ZN10MemoryCell10activateInEd, @function
_ZN10MemoryCell10activateInEd:
.LFB4726:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE4726:
	.size	_ZN10MemoryCell10activateInEd, .-_ZN10MemoryCell10activateInEd
	.align 2
	.globl	_ZN10MemoryCell11activateOutEd
	.type	_ZN10MemoryCell11activateOutEd, @function
_ZN10MemoryCell11activateOutEd:
.LFB5438:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN10MemoryCell10activateInEd
	.cfi_endproc
.LFE5438:
	.size	_ZN10MemoryCell11activateOutEd, .-_ZN10MemoryCell11activateOutEd
	.align 2
	.globl	_ZN10MemoryCell9copyToGPUEPS_
	.type	_ZN10MemoryCell9copyToGPUEPS_, @function
_ZN10MemoryCell9copyToGPUEPS_:
.LFB4728:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	$152, %esi
	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
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
	subq	$48, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rdi
	leaq	8(%rsp), %r14
	leaq	16(%rsp), %r13
	call	cudaMalloc@PLT
	leaq	24(%rsp), %r12
	leaq	32(%rsp), %rbp
	call	cudaDeviceSynchronize@PLT
	movq	(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbx, %rsi
	movl	$152, %edx
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movslq	8(%rbx), %rsi
	movq	%r14, %rdi
	salq	$3, %rsi
	call	cudaMalloc@PLT
	movslq	8(%rbx), %rsi
	movq	%r13, %rdi
	salq	$3, %rsi
	call	cudaMalloc@PLT
	movslq	8(%rbx), %rsi
	movq	%r12, %rdi
	salq	$3, %rsi
	call	cudaMalloc@PLT
	movslq	8(%rbx), %rsi
	movq	%rbp, %rdi
	salq	$3, %rsi
	call	cudaMalloc@PLT
	call	cudaDeviceSynchronize@PLT
	movslq	8(%rbx), %rdx
	movq	16(%rbx), %rsi
	movl	$1, %ecx
	movq	8(%rsp), %rdi
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movslq	8(%rbx), %rdx
	movq	32(%rbx), %rsi
	movl	$1, %ecx
	movq	16(%rsp), %rdi
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movslq	8(%rbx), %rdx
	movq	40(%rbx), %rsi
	movl	$1, %ecx
	movq	24(%rsp), %rdi
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movslq	8(%rbx), %rdx
	movq	24(%rbx), %rsi
	movl	$1, %ecx
	movq	32(%rsp), %rdi
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movq	(%rsp), %rax
	movl	$1, %ecx
	movq	%r14, %rsi
	movl	$8, %edx
	leaq	16(%rax), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rax
	movl	$1, %ecx
	movq	%r13, %rsi
	movl	$8, %edx
	leaq	32(%rax), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rax
	movl	$1, %ecx
	movq	%r12, %rsi
	movl	$8, %edx
	leaq	40(%rax), %rdi
	call	cudaMemcpy@PLT
	movq	(%rsp), %rax
	movl	$1, %ecx
	movq	%rbp, %rsi
	movl	$8, %edx
	leaq	24(%rax), %rdi
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movq	(%rsp), %rax
	movq	40(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L10
	call	__stack_chk_fail@PLT
.L10:
	addq	$48, %rsp
	.cfi_def_cfa_offset 48
	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
	ret
	.cfi_endproc
.LFE4728:
	.size	_ZN10MemoryCell9copyToGPUEPS_, .-_ZN10MemoryCell9copyToGPUEPS_
	.align 2
	.globl	_ZN10MemoryCell11copyFromGPUEPS_
	.type	_ZN10MemoryCell11copyFromGPUEPS_, @function
_ZN10MemoryCell11copyFromGPUEPS_:
.LFB4729:
	.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
	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	%rdi, %rbp
	movl	$152, %edi
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	pushq	%rax
	.cfi_def_cfa_offset 64
	call	malloc@PLT
	movq	%rax, %rbx
	call	cudaDeviceSynchronize@PLT
	movq	%rbp, %rsi
	movl	$2, %ecx
	movq	%rbx, %rdi
	movl	$152, %edx
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movslq	8(%rbx), %rbp
	salq	$3, %rbp
	movq	%rbp, %rdi
	call	malloc@PLT
	movq	%rbp, %rdi
	movq	%rax, %r15
	call	malloc@PLT
	movq	%rbp, %rdi
	movq	%rax, %r14
	call	malloc@PLT
	movq	%rbp, %rdi
	movq	%rax, %r13
	call	malloc@PLT
	movq	16(%rbx), %rsi
	movq	%rbp, %rdx
	movq	%r15, %rdi
	movl	$2, %ecx
	movq	%rax, %r12
	call	cudaMemcpy@PLT
	movslq	8(%rbx), %rdx
	movq	32(%rbx), %rsi
	movq	%r14, %rdi
	movl	$2, %ecx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movslq	8(%rbx), %rdx
	movq	40(%rbx), %rsi
	movq	%r13, %rdi
	movl	$2, %ecx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	movslq	8(%rbx), %rdx
	movq	24(%rbx), %rsi
	movq	%r12, %rdi
	movl	$2, %ecx
	salq	$3, %rdx
	call	cudaMemcpy@PLT
	call	cudaDeviceSynchronize@PLT
	movq	%r15, 16(%rbx)
	movq	%rbx, %rax
	movq	%r14, 32(%rbx)
	movq	%r13, 40(%rbx)
	movq	%r12, 24(%rbx)
	popq	%rdx
	.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
.LFE4729:
	.size	_ZN10MemoryCell11copyFromGPUEPS_, .-_ZN10MemoryCell11copyFromGPUEPS_
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4755:
	.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
.LFE4755:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text._ZSt18generate_canonicalIdLm53ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_,"axG",@progbits,_ZSt18generate_canonicalIdLm53ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_,comdat
	.weak	_ZSt18generate_canonicalIdLm53ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	.type	_ZSt18generate_canonicalIdLm53ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_, @function
_ZSt18generate_canonicalIdLm53ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_:
.LFB5396:
	.cfi_startproc
	endbr64
	imulq	$16807, (%rdi), %rax
	movl	$2147483647, %ecx
	xorl	%edx, %edx
	divq	%rcx
	movq	%rdx, %rax
	decq	%rax
	js	.L18
	cvtsi2sdq	%rax, %xmm0
	jmp	.L19
.L18:
	movq	%rax, %rcx
	andl	$1, %eax
	shrq	%rcx
	orq	%rax, %rcx
	cvtsi2sdq	%rcx, %xmm0
	addsd	%xmm0, %xmm0
.L19:
	imulq	$16807, %rdx, %rax
	movl	$2147483647, %ecx
	xorps	%xmm1, %xmm1
	xorl	%edx, %edx
	addsd	%xmm1, %xmm0
	divq	%rcx
	movq	%rdx, (%rdi)
	decq	%rdx
	js	.L20
	cvtsi2sdq	%rdx, %xmm1
	jmp	.L21
.L20:
	movq	%rdx, %rax
	andl	$1, %edx
	shrq	%rax
	orq	%rdx, %rax
	cvtsi2sdq	%rax, %xmm1
	addsd	%xmm1, %xmm1
.L21:
	mulsd	.LC2(%rip), %xmm1
	addsd	%xmm1, %xmm0
	divsd	.LC3(%rip), %xmm0
	comisd	.LC4(%rip), %xmm0
	jb	.L17
	movsd	.LC0(%rip), %xmm0
.L17:
	ret
	.cfi_endproc
.LFE5396:
	.size	_ZSt18generate_canonicalIdLm53ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_, .-_ZSt18generate_canonicalIdLm53ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	.section	.text._ZNSt19normal_distributionIdEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEdRT_RKNS0_10param_typeE,"axG",@progbits,_ZNSt19normal_distributionIdEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEdRT_RKNS0_10param_typeE,comdat
	.align 2
	.weak	_ZNSt19normal_distributionIdEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEdRT_RKNS0_10param_typeE
	.type	_ZNSt19normal_distributionIdEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEdRT_RKNS0_10param_typeE, @function
_ZNSt19normal_distributionIdEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEdRT_RKNS0_10param_typeE:
.LFB5233:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	movq	%rdx, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rdi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 64
	cmpb	$0, 24(%rdi)
	je	.L29
	movb	$0, 24(%rdi)
	movsd	16(%rdi), %xmm1
	jmp	.L26
.L33:
	xorps	%xmm0, %xmm0
	ucomisd	%xmm0, %xmm2
	jp	.L28
	jne	.L28
.L29:
	movq	%r12, %rdi
	call	_ZSt18generate_canonicalIdLm53ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	movq	%r12, %rdi
	addsd	%xmm0, %xmm0
	subsd	.LC4(%rip), %xmm0
	movsd	%xmm0, 8(%rsp)
	call	_ZSt18generate_canonicalIdLm53ESt26linear_congruential_engineImLm16807ELm0ELm2147483647EEET_RT1_
	movsd	8(%rsp), %xmm2
	addsd	%xmm0, %xmm0
	mulsd	%xmm2, %xmm2
	movaps	%xmm0, %xmm1
	subsd	.LC4(%rip), %xmm1
	movaps	%xmm1, %xmm0
	mulsd	%xmm1, %xmm0
	addsd	%xmm0, %xmm2
	comisd	.LC4(%rip), %xmm2
	ja	.L29
	jmp	.L33
.L28:
	movaps	%xmm2, %xmm0
	movsd	%xmm1, 24(%rsp)
	movsd	%xmm2, 16(%rsp)
	call	log@PLT
	mulsd	.LC5(%rip), %xmm0
	movsd	16(%rsp), %xmm2
	divsd	%xmm2, %xmm0
	call	sqrt@PLT
	movsd	8(%rsp), %xmm2
	movsd	24(%rsp), %xmm1
	movb	$1, 24(%rbx)
	mulsd	%xmm0, %xmm2
	mulsd	%xmm0, %xmm1
	movsd	%xmm2, 16(%rbx)
.L26:
	mulsd	8(%rbp), %xmm1
	addsd	0(%rbp), %xmm1
	addq	$32, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	movaps	%xmm1, %xmm0
	ret
	.cfi_endproc
.LFE5233:
	.size	_ZNSt19normal_distributionIdEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEdRT_RKNS0_10param_typeE, .-_ZNSt19normal_distributionIdEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEdRT_RKNS0_10param_typeE
	.text
	.align 2
	.globl	_ZN10MemoryCellC2Ei
	.type	_ZN10MemoryCellC2Ei, @function
_ZN10MemoryCellC2Ei:
.LFB4720:
	.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	%rdi, %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
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%esi, %ebx
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	leaq	16(%rsp), %r15
	call	_ZN8BaseNodeC2Ev@PLT
	leaq	16+_ZTV10MemoryCell(%rip), %rax
	xorl	%edi, %edi
	movl	%ebx, 8(%r14)
	movq	%rax, (%r14)
	movq	$0x000000000, 64(%r14)
	movq	$0x000000000, 72(%r14)
	movq	$0x000000000, 80(%r14)
	movq	$0x000000000, 88(%r14)
	movq	$0x000000000, 96(%r14)
	movq	$0x000000000, 104(%r14)
	movq	$0x000000000, 112(%r14)
	movq	$0x000000000, 120(%r14)
	movq	$0x000000000, 56(%r14)
	call	time@PLT
	movq	_ZN10MemoryCell1nE(%rip), %rdx
	movb	$0, 48(%rsp)
	movq	$0x000000000, 24(%rsp)
	movq	$0x000000000, 40(%rsp)
	leaq	1(%rdx), %rsi
	addq	%rdx, %rax
	xorl	%edx, %edx
	movq	%rsi, _ZN10MemoryCell1nE(%rip)
	movl	$2147483647, %esi
	divq	%rsi
	movl	$1, %eax
	movq	%r15, %rsi
	testq	%rdx, %rdx
	cmove	%rax, %rdx
	movq	.LC4(%rip), %rax
	movq	%rax, 32(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, %rdi
	movq	%rdx, 16(%rsp)
	movq	%rax, %rdx
	movq	%rax, 8(%rsp)
	call	_ZNSt19normal_distributionIdEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEdRT_RKNS0_10param_typeE
	movslq	%ebx, %rax
	movq	$0x000000000, 128(%r14)
	movq	$0x000000000, 136(%r14)
	salq	$3, %rax
	movq	$0x000000000, 144(%r14)
	movq	%rax, %rdi
	movsd	%xmm0, 48(%r14)
	movq	%rax, (%rsp)
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 16(%r14)
	movq	%rax, %rbp
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 24(%r14)
	movq	%rax, %r12
	call	malloc@PLT
	movq	(%rsp), %rdi
	movq	%rax, 40(%r14)
	movq	%rax, %r13
	call	malloc@PLT
	movq	%rax, 32(%r14)
	xorl	%r14d, %r14d
.L35:
	cmpl	%r14d, %ebx
	jle	.L39
	movq	8(%rsp), %rdx
	movq	%r15, %rsi
	movq	%rax, (%rsp)
	movq	%rdx, %rdi
	call	_ZNSt19normal_distributionIdEclISt26linear_congruential_engineImLm16807ELm0ELm2147483647EEEEdRT_RKNS0_10param_typeE
	movq	(%rsp), %rax
	movq	$0x000000000, (%r12,%r14,8)
	movq	$0x000000000, 0(%r13,%r14,8)
	movq	$0x000000000, (%rax,%r14,8)
	movsd	%xmm0, 0(%rbp,%r14,8)
	incq	%r14
	jmp	.L35
.L39:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L37
	call	__stack_chk_fail@PLT
.L37:
	addq	$72, %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
.LFE4720:
	.size	_ZN10MemoryCellC2Ei, .-_ZN10MemoryCellC2Ei
	.globl	_ZN10MemoryCellC1Ei
	.set	_ZN10MemoryCellC1Ei,_ZN10MemoryCellC2Ei
	.weak	_ZTS10MemoryCell
	.section	.rodata._ZTS10MemoryCell,"aG",@progbits,_ZTS10MemoryCell,comdat
	.align 8
	.type	_ZTS10MemoryCell, @object
	.size	_ZTS10MemoryCell, 13
_ZTS10MemoryCell:
	.string	"10MemoryCell"
	.weak	_ZTI10MemoryCell
	.section	.data.rel.ro._ZTI10MemoryCell,"awG",@progbits,_ZTI10MemoryCell,comdat
	.align 8
	.type	_ZTI10MemoryCell, @object
	.size	_ZTI10MemoryCell, 40
_ZTI10MemoryCell:
	.quad	_ZTVN10__cxxabiv121__vmi_class_type_infoE+16
	.quad	_ZTS10MemoryCell
	.long	0
	.long	1
	.quad	_ZTI8BaseNode
	.quad	0
	.weak	_ZTV10MemoryCell
	.section	.data.rel.ro.local._ZTV10MemoryCell,"awG",@progbits,_ZTV10MemoryCell,comdat
	.align 8
	.type	_ZTV10MemoryCell, @object
	.size	_ZTV10MemoryCell, 32
_ZTV10MemoryCell:
	.quad	0
	.quad	_ZTI10MemoryCell
	.quad	_ZN10MemoryCellD1Ev
	.quad	_ZN10MemoryCellD0Ev
	.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
	.globl	_ZN10MemoryCell1nE
	.bss
	.align 8
	.type	_ZN10MemoryCell1nE, @object
	.size	_ZN10MemoryCell1nE, 8
_ZN10MemoryCell1nE:
	.zero	8
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC0:
	.long	-1
	.long	1072693247
	.align 8
.LC2:
	.long	-8388608
	.long	1105199103
	.align 8
.LC3:
	.long	-16777216
	.long	1137704959
	.align 8
.LC4:
	.long	0
	.long	1072693248
	.align 8
.LC5:
	.long	0
	.long	-1073741824
	.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 MEMORYCELL_H_
#define MEMORYCELL_H_

#include "BaseNode.cuh"
#include <vector>
#include <math.h>
#include <time.h>
#include <iostream>
#include <stdlib.h>
#include <string.h>
#include <hip/hip_runtime.h>
#include <random>
using namespace std;

class MemoryCell : BaseNode {
private:
	static long long n;
public:
	int nConnections;
	double *cellDataWeight, *cellDataPartial,
		*inputDataPartial, *forgetDataPartial;
	double cellFeedbackWeight, bias;
	double activationIn, activationInPrime,
		activationOut, activationOutPrime,
		state, previousState,
		feedback, previousFeedback,
		cellFeedbackPartial;
	double inputFeedbackPartial,
		forgetFeedbackPartial;
	__device__ double activateIn(double data);
	__device__ double activateOut(double data);
	MemoryCell(int c);
	virtual ~MemoryCell();
	static MemoryCell *copyToGPU(MemoryCell *memory);
	static MemoryCell *copyFromGPU(MemoryCell *memory);
};

#endif /* MEMORYCELL_H_ */

	.text
	.file	"MemoryCell.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.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
