From d30033359c0b97b28b50f54eeea6e820156f7578 Mon Sep 17 00:00:00 2001 From: KunoiSayami Date: Sun, 21 Nov 2021 23:22:35 +0800 Subject: fix(skiplist): Fix some device function not declear correctly Signed-off-by: KunoiSayami --- util/arena.cuh | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) (limited to 'util/arena.cuh') 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 blocks_; - std::vector blocks_; + //std::vector blocks_; + + ArenaNode * head_; + ArenaNode * blocks_; + // Total memory usage of the arena. // @@ -55,7 +64,7 @@ class Arena { cuda::atomic 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). -- cgit v1.3.1