diff options
Diffstat (limited to 'util/arena.cuh')
| -rw-r--r-- | util/arena.cuh | 15 |
1 files changed, 12 insertions, 3 deletions
diff --git a/util/arena.cuh b/util/arena.cuh index 618b426..afbc575 100644 --- a/util/arena.cuh +++ b/util/arena.cuh @@ -25,7 +25,7 @@ class Arena { ~Arena(); // Return a pointer to a newly allocated memory block of "bytes" bytes. - char* Allocate(size_t bytes); + __device__ char* Allocate(size_t bytes); // Allocate memory with the normal alignment guarantees provided by malloc. __device__ char* AllocateAligned(size_t bytes); @@ -40,13 +40,22 @@ class Arena { __device__ char* AllocateFallback(size_t bytes); __device__ char* AllocateNewBlock(size_t block_bytes); + struct ArenaNode { + char * block; + ArenaNode * next; + }; + // Allocation state char* alloc_ptr_; size_t alloc_bytes_remaining_; // Array of new[] allocated memory blocks //thrust::host_vector<char *> blocks_; - std::vector<char*> blocks_; + //std::vector<char*> blocks_; + + ArenaNode * head_; + ArenaNode * blocks_; + // Total memory usage of the arena. // @@ -55,7 +64,7 @@ class Arena { cuda::atomic<size_t> memory_usage_; }; -inline char* Arena::Allocate(size_t bytes) { +__device__ inline char* Arena::Allocate(size_t bytes) { // The semantics of what to return are a bit messy if we allow // 0-byte allocations, so we disallow them here (we don't need // them for our internal use). |