 		fid = createNewFile(name);
 	}
 	return fid;
+}
+__device__ void sortD() {
+	for (int i = 0; i < fileCt; i++) {
+		volume[LS_D_OFFSET+i*2] 	= (i & 0x0000FF00) >> 8;
+		volume[LS_D_OFFSET+i*2+1] 	= i & 0x000000FF;
+	}
+}
+__device__ void printD() {
+	for (int i = 0; i < fileCt; i++) {
+		int fid = ((int)(volume[LS_D_OFFSET+i*2]) << 8) | ((int)(volume[LS_D_OFFSET+i*2+1]));
+		printf("%s\n", volume[META_OFFSET+fid*META_SIZE]);
+	}
+}
+__device__ void sortS() {
+	for (int i = 0; i < fileCt; i++) {
+		volume[LS_S_OFFSET+i*2] 	= i & 0xFFFF0000;
+		volume[LS_S_OFFSET+i*2+1] 	= i & 0x0000FFFF;
+	}
+}
+__device__ void printS() {
+	for (int i = 0; i < fileCt; i++) {
+		int fid = ((int)(volume[LS_S_OFFSET+i*2]) << 8) | ((int)(volume[LS_S_OFFSET+i*2+1]));
+		printf("%s\n", volume[META_OFFSET+fid*META_SIZE]);
+	}
+}
+__device__ void deleteFileByName(const char *name) {
+	int target = DNE;
+	for (int i = 0; i < fileCt; i++) {
+		if(isNameMatched(i, name)) {
+			target = i;
+			break;
+		}
+	}
+	fileCt--;
+	if (target != fileCt) {
+		for (int i = 0; i < META_SIZE; i++) {
+			volume[META_OFFSET+target*META_SIZE+i] = volume[META_OFFSET+fileCt*META_SIZE+i];
+		}
+		for (int i = 0; i < volume[META_OFFSET+fileCt*META_OFFSET+SIZE_OFFSET]; i++) {
+			volume[DATA_OFFSET+target*DATA_SIZE+i] = volume[DATA_OFFSET+fileCt*DATA_SIZE+i];
+		}
+	}
 }
 // ******************************************************************
 

const int RM_RF     = 6;

// Declare variables
__device__ __managed__ uchar *volume;

// Function
// ******************************************************************


// ******************************************************************
// 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++) {
