 const int RM_RF     = 6;
 
 // Declare variables
+__device__ __managed__ u32 currentTime = 0;
+__device__ __managed__ u32 fileCt = 0;
 __device__ __managed__ uchar *volume;
+// ==================================================================
+// How I use volume:
+// V--0          V--3072           V--36864
+// ------------------------------------------------------------------
+// | Temp | LS_S | LS_D | Metadata |              Data              |
+// ------------------------------------------------------------------
+//        ^--1024       ^--5120
+// Temp => For pending
+// LS_S => Array for sorting by LS_S
+// LS_D => Array for sorting by LS_D
+// Metadata =>
+//      Entry structure:
+//      V--0   V--21  V--23     V--27
+//      -----------------------------------
+//      | Name | Size | CreateT | ModifyT |
+//      -----------------------------------
+//      Name    := 21 bytes, for 20 char and '\0'
+//      Size    :=  2 bytes
+//      CreateT :=  4 bytes
+//      ModifyT :=  4 bytes
+//      Total size = 31 bytes
+// ==================================================================
 
 // Function
 // ******************************************************************

#define DATAFILE            "./data.bin"
#define OUTPUTFILE          "./snapshot.bin"
#define FILE_ENTRIES		1024
#define DNE					0xffffffff
typedef unsigned char uchar;
typedef uint32_t u32;


// ******************************************************************
// FS Helper function
__device__ inline int dstrlen(const char *str) {
	int len = 0;
	for (; str[len] != '\0'; len++);
	return len;
}
__device__ bool isNameMatched(int fid, const char *name) {
	bool isMatched = true;
	int len = dstrlen(name);
	for (int j = 0; j < len; j++) {
		if (volume[META_OFFSET+fid*META_SIZE+j] != name[j]) {
			isMatched = false;
			break;
		}
	}
	if (volume[META_OFFSET+fid*META_SIZE+len] != '\0')
		isMatched = false;
	return isMatched;
}
__device__ u32 createNewFile(const char *name) {
	u32 fid = fileCt++;
	int len = 0;
	for (; name[len] != '\0'; len++) {
		volume[META_OFFSET+fid*META_SIZE+len] = name[len];
	}
	volume[META_OFFSET+fid*META_SIZE+len] = '\0';
	volume[META_OFFSET+fid*META_SIZE+SIZE_OFFSET]   = 0;
	volume[META_OFFSET+fid*META_SIZE+SIZE_OFFSET+1] = 0;
	volume[META_OFFSET+fid*META_SIZE+CREATET_OFFSET]   = (currentTime & 0xFF000000) >> 24;
	volume[META_OFFSET+fid*META_SIZE+CREATET_OFFSET+1] = (currentTime & 0x00FF0000) >> 16;
	volume[META_OFFSET+fid*META_SIZE+CREATET_OFFSET+2] = (currentTime & 0x0000FF00) >> 8;
	volume[META_OFFSET+fid*META_SIZE+CREATET_OFFSET+3] = (currentTime & 0x000000FF);
	volume[META_OFFSET+fid*META_SIZE+MODIFYT_OFFSET]   = (currentTime & 0xFF000000) >> 24;
	volume[META_OFFSET+fid*META_SIZE+MODIFYT_OFFSET+1] = (currentTime & 0x00FF0000) >> 16;
	volume[META_OFFSET+fid*META_SIZE+MODIFYT_OFFSET+2] = (currentTime & 0x0000FF00) >> 8;
	volume[META_OFFSET+fid*META_SIZE+MODIFYT_OFFSET+3] = (currentTime & 0x000000FF);
	currentTime++;
	return fid;
}
__device__ inline void updateFileSize(int fid, int len) {
	volume[META_OFFSET+fid*META_SIZE+SIZE_OFFSET]   = (len & 0x0000FF00) >> 8;
	volume[META_OFFSET+fid*META_SIZE+SIZE_OFFSET+1] = (len & 0x000000FF);
}
__device__ inline void updateFileModifyTime(int fid) {
	volume[META_OFFSET+fid*META_SIZE+MODIFYT_OFFSET]   = (currentTime & 0xFF000000) >> 24;
	volume[META_OFFSET+fid*META_SIZE+MODIFYT_OFFSET+1] = (currentTime & 0x00FF0000) >> 16;
	volume[META_OFFSET+fid*META_SIZE+MODIFYT_OFFSET+2] = (currentTime & 0x0000FF00) >> 8;
	volume[META_OFFSET+fid*META_SIZE+MODIFYT_OFFSET+3] = (currentTime & 0x000000FF);
	currentTime++;
}
__device__ inline void updateData(int addr, uchar data) {
	volume[addr] = data;
}
__device__ inline u32 getData(int addr) {
	return volume[addr];
}
__device__ inline u32 getFileDataAddress(int fid) {
	return DATA_OFFSET+fid*DATA_SIZE;
}
__device__ u32 getFid(const char *name) {
	u32 fid = DNE;
	for (u32 i = 0; i < FILE_ENTRIES; i++) {
