 const uint32_t INVALID			= 0;
 const uint32_t PAGENUMBERMASK	= 0x00003FFE;
 const uint32_t LASTTIMEMASK		= 0xFFFFC000;
+const uint32_t PIDMASK			= 0xC0000000;
 const uint32_t DNE				= 0xFFFFFFFF;
 
 // Lock & Unlock


int main() {
    int input_size = loadBinaryFile(DATAFILE, input, STORAGE_SIZE);
    cudaSetDevice(2);
    mykernel<<<1, 4, 16384>>>(input_size/4);
    cudaDeviceSynchronize();
    cudaDeviceReset();

	u32 leastEntry = DNE;
	u32 leastTime  = DNE;
	for (u32 i = 0; i < PAGE_ENTRIES; i++) {
		if (leastTime > getLastUsedTime(pageTable[i]) && threadIdx.x == getPid(pageTable[i])) {
			leastTime = getLastUsedTime(pageTable[i]);
			leastEntry = i;
		}
