// Copyright (c) 2011 The LevelDB Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file. See the AUTHORS file for names of contributors.

#include "util/arena.cuh"

namespace leveldb {

static const int kBlockSize = 4096;

__device__ Arena::Arena()
    : alloc_ptr_(nullptr), alloc_bytes_remaining_(0), memory_usage_(0),
      head_(nullptr), blocks_(nullptr) {}

__device__ Arena::~Arena() {
  ArenaNode * current = this->head_;
  while (current != nullptr) {
    ArenaNode * next = current->next;
    cudaFree(current->block);
    cudaFree(current);
    current = next;
  }
}

__device__ char* Arena::AllocateFallback(size_t bytes) {
  if (bytes > kBlockSize / 4) {
    // Object is more than a quarter of our block size.  Allocate it separately
    // to avoid wasting too much space in leftover bytes.
    char* result = AllocateNewBlock(bytes);
    return result;
  }

  // We waste the remaining space in the current block.
  alloc_ptr_ = AllocateNewBlock(kBlockSize);
  alloc_bytes_remaining_ = kBlockSize;

  char* result = alloc_ptr_;
  alloc_ptr_ += bytes;
  alloc_bytes_remaining_ -= bytes;
  return result;
}

__device__ char* Arena::AllocateAligned(size_t bytes) {
  const int align = (sizeof(void*) > 8) ? sizeof(void*) : 8;
  static_assert((align & (align - 1)) == 0,
                "Pointer size should be a power of 2");
  size_t current_mod = reinterpret_cast<uintptr_t>(alloc_ptr_) & (align - 1);
  size_t slop = (current_mod == 0 ? 0 : align - current_mod);
  size_t needed = bytes + slop;
  char* result;
  if (needed <= alloc_bytes_remaining_) {
    result = alloc_ptr_ + slop;
    alloc_ptr_ += needed;
    alloc_bytes_remaining_ -= needed;
  } else {
    // AllocateFallback always returned aligned memory
    result = AllocateFallback(bytes);
  }
  assert((reinterpret_cast<uintptr_t>(result) & (align - 1)) == 0);
  return result;
}

__device__ char* Arena::AllocateNewBlock(size_t block_bytes) {
  char* result = nullptr;
  cuda::atomic<ArenaNode*>* block_alloc = nullptr;
  cudaMalloc((void **)&result, sizeof(char) * block_bytes);
  cudaMalloc((void **)&block_alloc, sizeof(cuda::atomic<ArenaNode*>));
  while (true) {
      ArenaNode * current_ = this->blocks_.load(cuda::memory_order_acquire);
      if (current_ == nullptr) {
      // cudaMalloc((void**)&this->blocks_, sizeof(ArenaNode));
          ArenaNode * current_end = this->blocks_.load();
          if (!this->blocks_.compare_exchange_weak(current_end, reinterpret_cast<ArenaNode*>(block_alloc)))
            continue;
          if (!this->head_.compare_exchange_weak(current_end, reinterpret_cast<ArenaNode*>(block_alloc)))
             continue;
          break;
      }
      ArenaNode * except_next = current_->next;
      if (except_next != nullptr)
          continue;
      if (!this->blocks_.compare_exchange_weak(current_, reinterpret_cast<ArenaNode*>(block_alloc)))
          continue;
      current_->block = result;
      current_->next = nullptr;
      break;
  }
  memory_usage_.fetch_add(block_bytes + sizeof(char*),
                          cuda::memory_order_relaxed);
  return result;
}

}  // namespace leveldb