 __device__ static void fallocMark(fallocContext* ctx, void* &mark, unsigned short &mark2) { mark = ctx->availableNodes; mark2 = ctx->availableNodes->freeOffset; }
 __device__ static bool fallocAtMark(fallocContext* ctx, void* mark, unsigned short mark2) { return (mark == ctx->availableNodes && mark2 == ctx->availableNodes->freeOffset); }
 
-
 #if __CUDA_ARCH__ > 100 // atomics only used with > sm_10 architecture
 //////////////////////
 // ATOMIC

    offset %= atom->bufferLength;
    return (void*)(atom->buffer + offset);
}

#endif

#endif // HOSTONLY

#ifdef TRACE
///////////////////////////////////////////////////////////////////////////////
// TRACE

const static int TRACEHEAP_SIZE = 2048;

typedef struct __align__(8) _cudaFallocTrace {
	volatile __int8* trace;
	fallocHeapChunk* lastChunk;
	int contextIndex;
	bool complete;
	struct _cudaFallocTrace* deviceTrace;
} fallocTrace;

#ifndef HOSTONLY

typedef struct {
	unsigned short magic;


#define CHUNKSIZE (sizeof(fallocHeapChunk)+HEAPCHUNK_SIZE)
#define CHUNKSIZEALIGN (CHUNKSIZE%16?CHUNKSIZE+16-(CHUNKSIZE%16):CHUNKSIZE)

#ifdef TRACE
///////////////////////////////////////////////////////////////////////////////
// TRACE
const static int TRACEHEAP_SIZE = 2048;

typedef struct __align__(8) _cudaFallocTrace {
	volatile __int8* trace;
	fallocHeapChunk* lastChunk;
	int contextIndex;
	bool complete;
	struct _cudaFallocTrace* deviceTrace;
} fallocTrace;
#endif


#ifndef HOSTONLY

