// 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) {}

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;
  cudaMalloc((void **)&result, sizeof(char) * block_bytes);
  if (this->blocks_ == nullptr) {
    cudaMalloc((void**)&this->blocks_, sizeof(ArenaNode));
    // First alloc
    this->head_ = this->blocks_;
  } else {
    cudaMalloc((void**)&this->blocks_->next, sizeof(ArenaNode));
    this->blocks_ = this->blocks_->next;
  }
  this->blocks_->block = result;
  this->blocks_->next = nullptr;
  memory_usage_.fetch_add(block_bytes + sizeof(char*),
                          cuda::memory_order_relaxed);
  return result;
}

}  // namespace leveldb