#ifndef PA_H
#define PA_H

#include <cuda.h>

namespace MACU{
	typedef unsigned int uint32;
	typedef int int32;
	typedef unsigned long long uint64;

	struct __align__(8) BlockListHeadDesc{ 
		int32 first_avail:32;	
		uint64 tag:32;	
	};

	struct BlockList{
		BlockListHeadDesc head_desc;
                __device__ int32 pop();        
	};

	struct Superblock{
		BlockList avail_list;	
		__device__ void * get_block();		
	};
	__device__ int32 BlockList::pop(){
		BlockListHeadDesc new_desc,old_desc;
		while(1){
			new_desc=old_desc=head_desc;
			if(old_desc.first_avail==-1)
				return -1;	
			new_desc.tag++;
			if(*(uint64*)&old_desc==atomicCAS((uint64*)&head_desc,*(uint64*)&old_desc,*(uint64*)&new_desc))
				return 0;
		}
	}
	__device__ void * Superblock::get_block(){
		while(1){
			uint32 block_num=avail_list.pop();
			return NULL;
		}	
	}


}
#endif