diff --git a/Assignment_3_120040077/main.o b/Assignment_3_120040077/main.o
deleted file mode 100644
index b3c85c0..0000000
Binary files a/Assignment_3_120040077/main.o and /dev/null differ
diff --git a/Assignment_3_120040077/main1.cu b/Assignment_3_120040077/main1.cu
new file mode 100644
index 0000000..f69df91
--- /dev/null
+++ b/Assignment_3_120040077/main1.cu
@@ -0,0 +1,350 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <inttypes.h>
+#include <cuda.h>
+#include <sys/types.h>
+#include <sys/stat.h>
+#include <fcntl.h>
+#include <unistd.h>
+#include <time.h>
+
+//page size is 32 bytes
+#define PAGESIZE (32)
+//32 KB in the shared memory
+#define PHYSICAL_MEM_SIZE (32768)
+//128 KB of secondary storage
+#define STORAGE_SIZE (131072)
+//memory segment size for each thread
+#define MEMORY_SEGMENT (32768)
+
+//number of pages in shared memory
+#define PHYSICAL_PAGE_NUM (PHYSICAL_MEM_SIZE/PAGESIZE)
+//number of pages in global memory
+#define STORAGE_PAGE_NUM (STORAGE_MEM_SIZE/PAGESIZE)
+
+#define DATAFILE "./data.bin"
+#define OUTFILE "./snapshot.bin"
+typedef unsigned char uchar;
+typedef uint32_t u32;
+
+/*----------------------------------The macro below are for page table--------------------------------*/
+
+#define INVALID_VALUE (0xffffffff)
+
+#define THREAD_PID_0					(0)
+#define THREAD_PID_1					(1)
+#define THREAD_PID_2					(2)
+#define THREAD_PID_3					(3)
+
+#define PID_BIT_START					(0)
+#define PID_BIT_LEN 					(2)
+#define VIRPAGE_BIT_START				(PID_BIT_START+PID_BIT_LEN)
+#define VIRPAGE_BIT_LEN					(12)
+#define COUNTER_BIT_START				(VIRPAGE_BIT_START+VIRPAGE_BIT_LEN)
+#define COUNTER_BIT_LEN					(18)
+
+#define FULL_MASK						(0xFFFFFFFF)
+#define PID_MASK						(0x3)
+#define VIRPAGE_MASK					(0x3FFC)
+#define COUNTER_MASK					(0xFFFFC000)
+
+#define GET_PID(x)						((x&PID_MASK)>>PID_BIT_START)
+#define GET_VIRPAGE(x)					((x&VIRPAGE_MASK)>>VIRPAGE_BIT_START)
+#define GET_COUNTER(x)					((x&COUNTER_MASK)>>COUNTER_BIT_START)
+
+#define CLEAR_PID(x)					(x&(~PID_MASK))
+#define CLEAR_VIRPAGE(x)				(x&(~VIRPAGE_MASK))
+#define CLEAR_COUNTER(x)				(x&(~COUNTER_MASK))
+
+#define SET_PID(src,value)				(CLEAR_PID(src)|(value<<PID_BIT_START))
+#define SET_VIRPAGE(src,value)			(CLEAR_VIRPAGE(src)|(value<<VIRPAGE_BIT_START))
+#define SET_COUNTER(src,value)			(CLEAR_COUNTER(src)|(value<<COUNTER_BIT_START))
+
+/*----------------------------------The macro below are for storage table-------------------------------------*/
+
+#define FIRST_PID_START					(0)
+#define FIRST_PID_LEN					(2)
+#define FIRST_VIRPAGE_START				(FIRST_PID_START+FIRST_PID_LEN)
+#define FIRST_VIRPAGE_LEN				(12)
+#define SECOND_PID_START				(FIRST_VIRPAGE_START+FIRST_VIRPAGE_LEN)
+#define SECOND_PID_LEN					(2)
+#define SECOND_VIRPAGE_START			(SECOND_PID_START+SECOND_PID_LEN)
+#define SECOND_VIRPAGE_LEN				(12)
+
+#define FIRST_PID_MASK					(0x3)
+#define FIRST_VIRPAGE_MASK				(0x3FFC)
+#define SECOND_PID_MASK					(0xC000)
+#define SECOND_VIRPAGE_MASK				(0xFFF0000)
+
+#define GET_FIRST_PID(x)				((x&FIRST_PID_MASK)>>FIRST_PID_START)
+#define GET_FIRST_VIRPAGE(x)			((x&FIRST_VIRPAGE_MASK)>>FIRST_VIRPAGE_START)
+#define GET_SECOND_PID(x)				((x&SECOND_PID_MASK)>>SECOND_PID_START)
+#define GET_SECOND_VIRPAGE(x)			((x&SECOND_VIRPAGE_MASK)>>SECOND_VIRPAGE_START)
+
+#define CLEAR_FIRST_PID(x)				(x&(~FIRST_PID_MASK))
+#define CLEAR_FIRST_VIRPAGE(x)			(x&(~FIRST_VIRPAGE_MASK))
+#define CLEAR_SECOND_PID(x)				(x&(~SECOND_PID_MASK))
+#define CLEAR_SECOND_VIRPAGE(x)			(x&(~SECOND_VIRPAGE_MASK))
+
+#define SET_FIRST_PID(src,value)		(CLEAR_FIRST_PID(src)|(value<<FIRST_PID_START))
+#define SET_FIRST_VIRPAGE(src,value)	(CLEAR_FIRST_VIRPAGE(src)|(value<<FIRST_VIRPAGE_START))
+#define SET_SECOND_PID(src,value)		(CLEAR_SECOND_PID(src)|(value<<SECOND_PID_START))
+#define SET_SECOND_VIRPAGE(src,value)	(CLEAR_SECOND_VIRPAGE(src)|(value<<SECOND_VIRPAGE_START))
+
+#define GET_STORAGE_PID(num)	 		((num%2==0)?GET_FIRST_PID(STORAGE_TABLE[num/2]):GET_SECOND_PID(STORAGE_TABLE[num/2]))
+#define GET_STORAGE_VIRPAGE(num)	 	((num%2==0)?GET_FIRST_VIRPAGE(STORAGE_TABLE[num/2]):GET_SECOND_VIRPAGE(STORAGE_TABLE[num/2]))
+#define SET_STORAGE_PID(num,value)		STORAGE_TABLE[num/2]=((num%2==0)?SET_FIRST_PID(STORAGE_TABLE[num/2],value):SET_SECOND_PID(STORAGE_TABLE[num/2],value))
+#define SET_STORAGE_VIRPAGE(num,value)	STORAGE_TABLE[num/2]=((num%2==0)?SET_FIRST_VIRPAGE(STORAGE_TABLE[num/2],value):SET_SECOND_VIRPAGE(STORAGE_TABLE[num/2],value))
+
+#define __LOCK(); for(int p=0;p<4;p++){if(threadIdx.x==p){
+#define __UNLOCK(); }__syncthreads();}
+
+//#define __GET_BASE() ((threadIdx.x)*MEMORY_SEGMENT)
+#define __GET_BASE() (p*MEMORY_SEGMENT)
+
+//Storage table
+
+__device__ __managed__ u32 *STORAGE_TABLE;
+__device__ __managed__ u32 STORAGE_COUNT;
+
+//Page-fault times
+__device__ __managed__ u32 PAGEFAULT = 0;
+
+//secondary memory
+__device__ __managed__ uchar storage[STORAGE_SIZE];
+
+//data input and output
+__device__ __managed__ uchar results[STORAGE_SIZE];
+__device__ __managed__ uchar input[STORAGE_SIZE];
+
+//shared memory for page table and memory-occupied table
+extern __shared__ u32 pt[];
+
+
+__device__ u32 paging(uchar *buffer, u32 page_num, u32 offset)
+{
+	u32 free_page_num=INVALID_VALUE;
+	u32 lru_time=1;
+	u32 lru_page_num=INVALID_VALUE;
+	u32 valid_page_num=INVALID_VALUE;
+	for(int i=0;i<PHYSICAL_PAGE_NUM;i++){
+		if((GET_PID(pt[i])==threadIdx.x) && (GET_VIRPAGE(pt[i])==page_num) && (GET_COUNTER(pt[i])>0)){
+			pt[i]=SET_COUNTER(pt[i],1);
+			valid_page_num=i;
+		}
+		else if(GET_COUNTER(pt[i])>0)
+			pt[i]=SET_COUNTER(pt[i],GET_COUNTER(pt[i])+1);
+
+		//get free page
+		if(GET_COUNTER(pt[i])==0)
+			free_page_num=i;
+		//get LRU page
+		if(GET_COUNTER(pt[i])>lru_time){
+			lru_time=GET_COUNTER(pt[i]);
+			lru_page_num=i;
+		}
+	}
+	if(valid_page_num!=INVALID_VALUE){
+		return valid_page_num*PAGESIZE+offset;
+	}
+	else if(free_page_num!=INVALID_VALUE){
+		pt[free_page_num]=SET_PID(pt[free_page_num],threadIdx.x);
+		pt[free_page_num]=SET_VIRPAGE(pt[free_page_num],page_num);
+		pt[free_page_num]=SET_COUNTER(pt[free_page_num],1);
+/*		for(int i=0;i<PHYSICAL_PAGE_NUM;i++){
+			if((GET_COUNTER(pt[i])>0)||(i==free_page_num))
+				pt[i]=SET_COUNTER(pt[i],GET_COUNTER(pt[i])+1);
+		}*/
+		PAGEFAULT++;
+		return free_page_num*PAGESIZE+offset;
+	}
+	else{
+		u32 swap_out_start=INVALID_VALUE;
+		u32 swap_in_start=INVALID_VALUE;
+		u32 phy_start=lru_page_num*PAGESIZE;
+		for(int i=0;i<STORAGE_COUNT;i++){
+			if((GET_STORAGE_PID(i)==GET_PID(pt[lru_page_num])) && (GET_STORAGE_VIRPAGE(i)==GET_VIRPAGE(pt[lru_page_num])))
+				swap_out_start=i*PAGESIZE;
+
+		
+			if((GET_STORAGE_PID(i)==threadIdx.x) && (GET_STORAGE_VIRPAGE(i)==page_num))
+				swap_in_start=i*PAGESIZE;
+		}
+		//if the storage does not contain the page to be swapped out
+		if(swap_out_start==INVALID_VALUE){
+			swap_out_start=STORAGE_COUNT*PAGESIZE;
+			SET_STORAGE_PID(STORAGE_COUNT,GET_PID(pt[lru_page_num]));
+			SET_STORAGE_VIRPAGE(STORAGE_COUNT,GET_VIRPAGE(pt[lru_page_num]));
+			STORAGE_COUNT++;
+		}
+		//if the page to be swapped in is not in the storage
+		if(swap_in_start==INVALID_VALUE){
+			for(int i=0;i<PAGESIZE;i++){
+				storage[swap_out_start+i]=buffer[phy_start+i];
+				buffer[phy_start+i]=0;
+			}
+		}
+		//if the page to be swapped in us in the storage
+		else{
+			for(int i=0;i<PAGESIZE;i++){
+				storage[swap_out_start+i]=buffer[phy_start+i];
+				buffer[phy_start+i]=storage[swap_in_start+i];
+			}
+		}
+		pt[lru_page_num]=SET_PID(pt[lru_page_num],threadIdx.x);
+		pt[lru_page_num]=SET_VIRPAGE(pt[lru_page_num],page_num);
+		pt[lru_page_num]=SET_COUNTER(pt[lru_page_num],1);
+/*		for(int i=0;i<PHYSICAL_PAGE_NUM;i++){
+			if((GET_COUNTER(pt[i])>0)||(i==lru_page_num))
+				pt[i]=SET_COUNTER(pt[i],GET_COUNTER(pt[i])+1);
+		}*/
+		PAGEFAULT++;
+		return phy_start+offset; 
+	}
+
+}
+
+__device__ uchar Gread(uchar *buffer,u32 addr)
+{
+	u32 page_num  = addr/PAGESIZE;
+	u32 offset 	  = addr%PAGESIZE;
+
+	//addr means the addr in shared memory
+	addr = paging(buffer, page_num, offset);
+	return buffer[addr];
+}
+__device__ void Gwrite(uchar *buffer, u32 addr, uchar value)
+{
+	u32 page_num  = addr/PAGESIZE;
+	u32 offset    = addr%PAGESIZE;
+
+	//addr means the addr in shared memory
+	addr = paging(buffer, page_num, offset);
+	buffer[addr] = value;
+}
+
+__device__ void snapshot(uchar *results, uchar *buffer, int offset, int input_size)
+{
+	for(int i=0;i<input_size;i++)
+		results[i] = Gread(buffer, i+offset);
+}
+
+__device__ void init_pageTable(int pt_entries)
+{
+	PAGEFAULT=0;
+	STORAGE_TABLE = pt+PHYSICAL_PAGE_NUM;
+	STORAGE_COUNT=0;
+	for(int i=0;i<PHYSICAL_PAGE_NUM+2048;i++){
+		pt[i]=0;
+	}
+	for(int i=0;i<STORAGE_SIZE;i++){
+		results[i]=0;
+		storage[i]=0;
+	}
+}
+
+int load_binaryFile(const char *filename, uchar *input, int size)
+{
+	int fd=0;
+	int sizeread=0;
+	int sizehasread=0;
+	fd=open(filename,O_RDONLY);
+	if(fd==-1){
+		perror("open data.bin error");
+		return -1;
+	}
+	while((sizeread=read(fd,input,size))!=-1){
+		sizehasread+=sizeread;
+		size-=sizeread;
+		input+=sizeread;
+		if(sizeread==0){
+			close(fd);
+			return sizehasread;
+		}
+	}
+	close(fd);
+	perror("read data.bin error");
+	return -1;
+}
+
+int write_binaryFile(const char *filename, uchar *output, int size)
+{
+	int fd=0;
+	int sizewritten=0;
+	int sizehaswritten=0;
+	fd=open(filename,O_WRONLY|O_CREAT,S_IRUSR|S_IWUSR|S_IRGRP|S_IWGRP|S_IROTH);
+	if(fd==-1){
+		perror("open snapshot.bin error");
+		return -1;
+	}
+	while((sizewritten=write(fd,output,size))!=-1){
+		sizehaswritten+=sizewritten;
+		size-=sizewritten;
+		output+=sizewritten;
+		if(sizewritten==0){
+			close(fd);
+			return sizehaswritten;
+		}
+	}
+	close(fd);
+	perror("write snapshot.bin error");
+	return -1;
+}
+
+__global__ void mykernel(int input_size)
+{
+	//Regard shared memory as physical memory
+	__shared__ uchar data[PHYSICAL_MEM_SIZE];
+
+	//get page table entries
+	int pt_entries=PHYSICAL_MEM_SIZE/PAGESIZE;
+	//We should initialize the page table
+	if(threadIdx.x==0)
+		init_pageTable(pt_entries);
+
+	
+	//####GWrite/Gread code section start####
+	__LOCK();
+	for(int i=0;i<input_size;i++)
+		Gwrite(data,i+__GET_BASE(),input[i+__GET_BASE()]);
+	__UNLOCK();
+
+
+	for(int i=input_size-1;i>=input_size-10;i--){
+	__LOCK();
+		int value = Gread(data,i+__GET_BASE());
+	__UNLOCK();
+	}
+
+	__LOCK();
+	snapshot(results+__GET_BASE(),data,__GET_BASE(),input_size);
+	__UNLOCK();
+	//####GWrite/Gread code section end####
+	printf("this thread pid = %d, total pagefault times=%u\n",threadIdx.x,PAGEFAULT);
+	return;
+}
+
+int main()
+{
+	clock_t t;
+	t=clock();
+
+	//Load data.bin into input buffer
+	int input_size = load_binaryFile(DATAFILE, input, STORAGE_SIZE);
+
+	printf("The read size is %d\n", input_size);
+
+	//main procedure
+	cudaSetDevice(4);
+	mykernel<<<1,4,16384>>>(input_size/4);
+	cudaDeviceSynchronize();
+	cudaDeviceReset();
+
+	//write binary file from results buffer
+	write_binaryFile(OUTFILE, results, input_size);
+
+	t=clock()-t;
+	printf("total elapsed time = %f\n",((float)t)/CLOCKS_PER_SEC);
+
+	return 0;
+}
\ No newline at end of file
diff --git a/Assignment_3_120040077/main1.o b/Assignment_3_120040077/main1.o
new file mode 100644
index 0000000..a67eb35
Binary files /dev/null and b/Assignment_3_120040077/main1.o differ
diff --git a/Assignment_3_120040077/report/report.tex b/Assignment_3_120040077/report/report.tex
new file mode 100644
index 0000000..81280fe
--- /dev/null
+++ b/Assignment_3_120040077/report/report.tex
@@ -0,0 +1,18 @@
+\documentclass{article}
+\usepackage[utf8]{inputenc}
+\usepackage{amsmath}
+\usepackage{color}
+\title{CSC3150 hw5 report}
+\author{Zhouliang Yu}
+\date{November 2021}
+
+\begin{document}
+
+\maketitle
+
+\section{How to Compile}
+
+\section{How to Design My Program}
+\subsection{}
+
+\end{document}
diff --git a/Assignment_3_120040077/snapshot.bin b/Assignment_3_120040077/snapshot.bin
index 8d757bb..ef88f4d 100644
--- a/Assignment_3_120040077/snapshot.bin
+++ b/Assignment_3_120040077/snapshot.bin
@@ -1358,4 +1358,4 @@ B'	c;U4
?5[$VaS	X"3BM
aU!3
 X4#38Z]NNQ
$/GAc[;/AONJ<E$

CN[.7!R5%K-*WdH
Z)F4(:55R_,15JR6b7T?N`NP&G3[@+'<%;U?BQ#1N-8
!=%b;?8 :/
 SGSc$6OF6	B	I*Q=GV9;EB&2I0NMO&+(CQ+O%I18:2B=c'5X]W!FA>d6'@G4OdX
,5=^#4/'3N^EH/>-!E`\c/&0
Z</L?-S6D!K,CYT@I4@
 '	JP8S@Z$J9N)d9.b"1`")AOUb=E+*dH36+:!D(=MJWRCdOaZ`bca6	-H%V!7>*+I.(	3b6U]S)</X8J;\<[cI0(=
!PS[EY[=6FH49G%:E=9<LM)X;&>X)Ed.&Gb^aK<?EW.36[*>4:O#>G=E(SUO\a2$*R02)F!\O7AD-DP
;>`]I\*NM
-#1<)Q[ b6KAc*-?
]`bA''+bDd3V3X,[# 0Nb\)22Uac1#Z+U9*#UP&K`	\RR(;OS3QT#MS$@N0 H"D,M3(VA%b0b&UUHd_1/ .7Y39X,1)$`EZ-O30?-a=D]
\ No newline at end of file
+#1<)Q[ b6KAc*-?
]`bA''+bDd3V3X,[# 0Nb\)22Uac1#Z+U9*#UP&K`	\RR(;OS3QT#MS$@N0 H"D,M3(VA%b0b&UUHd_1/
\ No newline at end of file
diff --git a/Assignment_3_120040077/user_program.o b/Assignment_3_120040077/user_program.o
deleted file mode 100644
index 399b6eb..0000000
Binary files a/Assignment_3_120040077/user_program.o and /dev/null differ
diff --git a/Assignment_3_120040077/virtual_memory.cu b/Assignment_3_120040077/virtual_memory.cu
index 7cf3c55..ecbfa50 100644
--- a/Assignment_3_120040077/virtual_memory.cu
+++ b/Assignment_3_120040077/virtual_memory.cu
@@ -42,14 +42,7 @@ __device__ bool check_page_fault(VirtualMemory *vm, u32 page_num) {
   return false;
 }
 
-__device__ void move_to_memory(VirtualMemory *vm, u32 frame_num, u32 page_num) {
-  u32 original_page_num = vm->invert_page_table[frame_num];
-  for (int i = 0; i < 32; i ++) {
-    vm->storage[original_page_num * 32 + i] = vm -> buffer[frame_num * 32 + i];
-    vm->buffer[frame_num * 32 + i] = vm -> storage[page_num * 32 + i];
-  }
-  vm -> invert_page_table[frame_num] = page_num;
-}
+
 
 __device__ int find_frame_number(VirtualMemory *vm, u32 page_num) {
   for (int i = 0; i < vm->PAGE_ENTRIES; i++) {
@@ -70,13 +63,7 @@ __device__ int find_frame_num_in_frame_table(VirtualMemory *vm, u32 frame_num)
   return -1;
 }
 
-__device__ void change_frame_table_valid_to_invalid(VirtualMemory *vm, u32 frame_num) {
-  int tempt = vm->invert_page_table[vm->PAGE_ENTRIES + find_frame_num_in_frame_table(vm, frame_num)];
-  for (int i = find_frame_num_in_frame_table(vm, frame_num); i < vm -> PAGE_ENTRIES - 1; i ++) {
-    vm->invert_page_table[i + vm->PAGE_ENTRIES] = vm->invert_page_table[i + vm->PAGE_ENTRIES + 1];
-  }
-  vm -> invert_page_table[2 * vm->PAGE_ENTRIES - 1] = tempt;
-}
+
 
 __device__ uchar vm_read(VirtualMemory *vm, u32 addr)
 {
@@ -84,30 +71,55 @@ __device__ uchar vm_read(VirtualMemory *vm, u32 addr)
   u32 page_num = addr / 32;
   u32 page_offset = addr % 32;
   u32 frame_num; 
+  bool isFault = false;
 
-  if (!check_page_fault(vm, page_num)) {
+/** check wethere there is a fault in page table*/
+  for (int i = 0; i < vm->PAGE_ENTRIES; i++) {
+    if (vm -> invert_page_table[i] == page_num) {
+      isFault = false;
+    }else {
+      isFault = true;
+    }
+  }
+  
+
+  if (isFault == false) {
     frame_num = vm -> invert_page_table[vm->PAGE_ENTRIES];
-    move_to_memory(vm, frame_num, page_num);
+    /** move to memory*/
+    u32 original_page_num = vm->invert_page_table[frame_num];
+    for (int i = 0; i < 32; i ++) {
+    vm->storage[original_page_num * 32 + i] = vm -> buffer[frame_num * 32 + i];
+    vm->buffer[frame_num * 32 + i] = vm -> storage[page_num * 32 + i];
+  }
+  vm -> invert_page_table[frame_num] = page_num;
   }else {
-    frame_num = find_frame_number(vm, page_num);
+    // frame_num = find_frame_number(vm, page_num);
+    /** find if there exists frame number*/
+    for (int i = 0; i < vm->PAGE_ENTRIES; i++)
+    {
+      if (vm->invert_page_table[i] == page_num)
+      {
+        frame_num = i;
+      }
+    }
+    frame_num = -1; //out of index or not found
   }
-  change_frame_table_valid_to_invalid(vm, frame_num);
-}
-
-__device__ void move_to_storage(VirtualMemory *vm, u32 frame_num){
-  u32 page_num = vm -> invert_page_table[frame_num];
-  for (int i = 0; i < 32; i ++) {
-    vm->storage[page_num * 32 + i] = vm->buffer[frame_num * 32 + i];
+  
+int tempt = vm->invert_page_table[vm->PAGE_ENTRIES + find_frame_num_in_frame_table(vm, frame_num)];
+  for (int i = find_frame_num_in_frame_table(vm, frame_num); i < vm -> PAGE_ENTRIES - 1; i ++) {
+    vm->invert_page_table[i + vm->PAGE_ENTRIES] = vm->invert_page_table[i + vm->PAGE_ENTRIES + 1];
   }
+  vm -> invert_page_table[2 * vm->PAGE_ENTRIES - 1] = tempt;
 }
 
-
-__device__ void check_frame_full(VirtualMemory *vm, u32 page_num, u32 frame_num){
-  if (vm->invert_page_table[frame_num] != 0x80000000)
-  {
-    move_to_storage(vm, frame_num);
-  }
-}
+// /** swap the page from buffer to the storage*/
+// __device__ void swap(VirtualMemory *vm, u32 frame_num){
+//   // u32 page_num = vm -> invert_page_table[frame_num];
+//   for (int i = 0; i < 32; i ++) {
+//     // vm->storage[vm -> invert_page_table[frame_num] * 32 + i] = vm->buffer[frame_num * 32 + i];
+//     vm->storage[(vm->invert_page_table[frame_num]) * 32 + i] = vm->buffer[frame_num * 32 + i]
+//   }
+// }
 
 
 __device__ void vm_write(VirtualMemory *vm, u32 addr, uchar value) {
@@ -115,23 +127,37 @@ __device__ void vm_write(VirtualMemory *vm, u32 addr, uchar value) {
   u32 page_num = addr / 32;
   u32 page_offset = addr % 32;
   u32 frame_num;
+
   if(!check_page_fault(vm, page_num)) {
-    frame_num = vm->invert_page_table[vm->PAGE_ENTRIES];
-    check_frame_full(vm,page_num,frame_num);
+    /** put the frame number as the top one*/
+  frame_num = vm->invert_page_table[vm->PAGE_ENTRIES];
+    
+/** check if frame is full*/
+  if (vm->invert_page_table[frame_num] != 0x80000000) 
+  {
+    /** if is full move to the storage*/
+	// swap(vm, frame_num);
+    for (int i = 0; i < 32; i++)
+    {
+      // vm->storage[vm -> invert_page_table[frame_num] * 32 + i] = vm->buffer[frame_num * 32 + i];
+      vm->storage[(vm->invert_page_table[frame_num]) * 32 + i] = vm->buffer[frame_num * 32 + i]
+    }
+  }
     vm->invert_page_table[frame_num] = page_num;
   }
   else{
-    frame_num = find_frame_number(vm, page_num);
+    /** LRU*/
+    // frame_num = find_frame_number(vm, page_num);
+
   }
   vm->buffer[frame_num * 32 + page_offset] = value;
-  change_frame_table_valid_to_invalid(vm, frame_num);
-}
 
-__device__ void move_to_result_buffer(VirtualMemory *vm, uchar* result, u32 page_num){
-  u32 frame_num = find_frame_number(vm, page_num);
-  for (int i = 0; i < 32; i++){
-    result[page_num * 32 + i] = vm -> buffer[frame_num * 32 + i]; // load element from vm buffer to result buffer in global memory
+/** change from invalid to valid*/
+  int tempt = vm->invert_page_table[vm->PAGE_ENTRIES + find_frame_num_in_frame_table(vm, frame_num)];
+  for (int i = find_frame_num_in_frame_table(vm, frame_num); i < vm -> PAGE_ENTRIES - 1; i ++) {
+    vm->invert_page_table[i + vm->PAGE_ENTRIES] = vm->invert_page_table[i + vm->PAGE_ENTRIES + 1];
   }
+  vm -> invert_page_table[2 * vm->PAGE_ENTRIES - 1] = tempt;
 }
 
 
@@ -145,12 +171,36 @@ __device__ void vm_snapshot(VirtualMemory *vm, uchar *results, int offset,
     u32 frame_num;
     if (!check_page_fault(vm, page_num)) {
       frame_num = vm -> invert_page_table[vm -> PAGE_ENTRIES];
-      move_to_memory(vm, frame_num, page_num);
+      u32 original_page_num = vm->invert_page_table[frame_num];
+    for (int i = 0; i < 32; i ++) {
+    vm->storage[original_page_num * 32 + i] = vm -> buffer[frame_num * 32 + i];
+    vm->buffer[frame_num * 32 + i] = vm -> storage[page_num * 32 + i];
+  }
+  vm -> invert_page_table[frame_num] = page_num;
     }else{
-      frame_num = find_frame_number(vm, page_num);
+      // frame_num = find_frame_number(vm, page_num);
+      /** find if there exists frame number*/
+      for (int i = 0; i < vm->PAGE_ENTRIES; i++)
+      {
+        if (vm->invert_page_table[i] == page_num)
+        {
+          frame_num = i;
+        }
+      }
+      frame_num = -1; //out of index or not found
     }
-    move_to_result_buffer(vm, results, page_num);
-    change_frame_table_valid_to_invalid(vm, frame_num);
+ 
+/** move to the results buffer*/
+ for (int i = 0; i < 32; i++){
+    results[page_num * 32 + i] = vm -> buffer[frame_num * 32 + i]; // load element from vm buffer to result buffer in global memory
+  }    
+
+/** change from invalid to valid*/
+    int tempt = vm->invert_page_table[vm->PAGE_ENTRIES + find_frame_num_in_frame_table(vm, frame_num)];
+  for (int i = find_frame_num_in_frame_table(vm, frame_num); i < vm -> PAGE_ENTRIES - 1; i ++) {
+    vm->invert_page_table[i + vm->PAGE_ENTRIES] = vm->invert_page_table[i + vm->PAGE_ENTRIES + 1];
+  }
+  vm -> invert_page_table[2 * vm->PAGE_ENTRIES - 1] = tempt;
   }
 }
 
diff --git a/Assignment_3_120040077/virtual_memory.o b/Assignment_3_120040077/virtual_memory.o
deleted file mode 100644
index 588a1ac..0000000
Binary files a/Assignment_3_120040077/virtual_memory.o and /dev/null differ
diff --git a/Assignment_3_120040077/vm b/Assignment_3_120040077/vm
deleted file mode 100644
index 56551d1..0000000
Binary files a/Assignment_3_120040077/vm and /dev/null differ