From 32ff766dbd8416b45ae042463c8f65145b75806f Mon Sep 17 00:00:00 2001
From: Christoph Lehner <christoph@lhnr.de>
Date: Sun, 13 Sep 2020 14:02:53 -0400
Subject: [PATCH] fix evict scheme, slab alloc

---
 Grid/allocator/MemoryManagerCache.cc | 6 ++++--
 Grid/threads/SlabAllocator.cc        | 8 ++++++++
 2 files changed, 12 insertions(+), 2 deletions(-)

diff --git a/Grid/allocator/MemoryManagerCache.cc b/Grid/allocator/MemoryManagerCache.cc
index 5dd7575e..7d4581d7 100644
--- a/Grid/allocator/MemoryManagerCache.cc
+++ b/Grid/allocator/MemoryManagerCache.cc
@@ -227,12 +227,13 @@ uint64_t MemoryManager::AcceleratorViewOpen(uint64_t CpuPtr,size_t bytes,ViewMod
   // Find if present, otherwise get or force an empty
   ////////////////////////////////////////////////////////////////////////////
   if ( EntryPresent(CpuPtr)==0 ){
-    EvictVictims(bytes);
     EntryCreate(CpuPtr,bytes,mode,hint);
   }
 
   auto AccCacheIterator = EntryLookup(CpuPtr);
   auto & AccCache = AccCacheIterator->second;
+  if (!AccCache.AccPtr)
+    EvictVictims(bytes);
   
   assert((mode==AcceleratorRead)||(mode==AcceleratorWrite)||(mode==AcceleratorWriteDiscard));
 
@@ -361,12 +362,13 @@ uint64_t MemoryManager::CpuViewOpen(uint64_t CpuPtr,size_t bytes,ViewMode mode,V
   // Find if present, otherwise get or force an empty
   ////////////////////////////////////////////////////////////////////////////
   if ( EntryPresent(CpuPtr)==0 ){
-    EvictVictims(bytes);
     EntryCreate(CpuPtr,bytes,mode,transient);
   }
 
   auto AccCacheIterator = EntryLookup(CpuPtr);
   auto & AccCache = AccCacheIterator->second;
+  if (!AccCache.AccPtr)
+    EvictVictims(bytes);
   
   assert((mode==CpuRead)||(mode==CpuWrite));
   assert(AccCache.accLock==0);  // Programming error
diff --git a/Grid/threads/SlabAllocator.cc b/Grid/threads/SlabAllocator.cc
index da863687..5590f835 100644
--- a/Grid/threads/SlabAllocator.cc
+++ b/Grid/threads/SlabAllocator.cc
@@ -36,6 +36,9 @@ NAMESPACE_BEGIN(Grid);
 #define GRID_DEVICE_HEAP_SLAB_THRESHOLD (1024*1024)
 #define GRID_DEVICE_HEAP_SLAB_SIZE (2*1024*1024)
 
+size_t currentDeviceAlloc = 0;
+std::unordered_map<void*,size_t> ptr_size;
+
 void *acceleratorAllocDeviceCUDA(size_t bytes) {
   void *ptr=NULL;
   auto err = cudaMalloc((void **)&ptr,bytes);
@@ -43,11 +46,16 @@ void *acceleratorAllocDeviceCUDA(size_t bytes) {
     ptr = (void *) NULL;
     printf(" cudaMalloc failed for %d %s \n",bytes,cudaGetErrorString(err));
   }
+  currentDeviceAlloc += bytes;
+  ptr_size[ptr] = bytes;
+  std::cout << "Current device alloc: " << currentDeviceAlloc << std::endl;
   return ptr;
 }
 
 void acceleratorFreeDeviceCUDA(void *ptr) {
   cudaFree(ptr);
+  currentDeviceAlloc -= ptr_size[ptr];
+  std::cout << "Current device alloc: " << currentDeviceAlloc << std::endl;
 }
 
 struct grid_device_heap_slab_t {