Query tracked memory sizes.

Support more accurate hipMemGetInfo.  Add test to hipPointerAttrib.

hipamd/include/hcc_detail/AM.h

@@ -78,6 +78,7 @@ am_status_t AM_free(void*  ptr);
  */
 am_status_t AM_copy(void*  dst, const void*  src, size_t size);
 
+
 /**
  * Return information about tracked pointer.
  *
@@ -92,10 +93,14 @@ am_status_t AM_copy(void*  dst, const void*  src, size_t size);
  */
 am_status_t am_memtracker_getinfo(hc::AmPointerInfo *info, void *ptr);
 
+
+//TODO-doc
+am_status_t am_memtracker_add(void* ptr, size_t sizeBytes, hc::accelerator acc, bool isDeviceMem=false);
+
+
 //TODO-doc
 am_status_t am_memtracker_update(void* ptr, int appId, unsigned allocationFlags);
 
-am_status_t am_memtracker_add(void* ptr, size_t sizeBytes, hc::accelerator acc, bool isDeviceMem=false);
 
 /** 
  * Remove the pointer from the tracker structure.
@@ -109,16 +114,20 @@ am_status_t am_memtracker_remove(void* ptr);
 
 /**
  * Remove all memory allocations associated with specified accelerator.
+ *
+ * @returns Number of entries reset.
  */
 size_t am_memtracker_reset(hc::accelerator acc);
 
 /**
- * Prints the contents of the memory tracker table to stderr
+ * Prints info about the memory tracker table.
  *
  * Intended primarily for debug purposes.
  **/
 void am_memtracker_print();
 
+void am_memtracker_sizeinfo(hc::accelerator acc, size_t *deviceMemSize, size_t *hostMemSize, size_t *userMemSize);
+
 
 }; // namespace hc
 

hipamd/src/hc_AM.cpp

@@ -63,13 +63,15 @@ public:
     void insert(void *pointer, const hc::AmPointerInfo &p);
     int remove(void *pointer);
 
-    MapTrackerType::iterator find(void *hostPtr);
+    MapTrackerType::iterator find(void *hostPtr) ;
+
+    MapTrackerType::iterator readerLockBegin() { _mutex.lock(); return _tracker.begin(); } ;
+    MapTrackerType::iterator end() { return _tracker.end(); } ;
+    void readerUnlock() { _mutex.unlock(); };
 
-    MapTrackerType::iterator end() { return _tracker.end(); };
 
     size_t reset (hc::accelerator acc);
 
-    std::ostream & print (std::ostream &os);
 private:
     // TODO - use or remove.
     inline void writeLock(); 
@@ -115,6 +117,7 @@ AmPointerTracker::MapTrackerType::iterator  AmPointerTracker::find (void *pointe
 }
 
 
+#if 0
 //---
 std::ostream & AmPointerTracker::print (std::ostream &os)
 {
@@ -126,6 +129,7 @@ std::ostream & AmPointerTracker::print (std::ostream &os)
 
     return os;
 }
+#endif
 
 //---
 // Remove all tracked locations, and free the associated memory (if the range was originally allocated by AM).
@@ -326,7 +330,38 @@ am_status_t am_memtracker_remove(void* ptr)
 //---
 void am_memtracker_print()
 {
-    g_amPointerTracker.print(std::cerr);
+    std::ostream &os = std::cerr;
+
+    //g_amPointerTracker.print(std::cerr);
+    for (auto iter = g_amPointerTracker.readerLockBegin() ; iter != g_amPointerTracker.end(); iter++) {
+        os << "  " << iter->first._basePointer << "..." << iter->first._endPointer << "::  ";
+        os << iter->second << std::endl;
+    }
+
+    g_amPointerTracker.readerUnlock();
+}
+
+
+//---
+void am_memtracker_sizeinfo(hc::accelerator acc, size_t *deviceMemSize, size_t *hostMemSize, size_t *userMemSize)
+{
+    *deviceMemSize = *hostMemSize = *userMemSize = 0;
+    for (auto iter = g_amPointerTracker.readerLockBegin() ; iter != g_amPointerTracker.end(); iter++) {
+        if (iter->second._acc == acc) {
+            size_t sizeBytes = iter->second._sizeBytes;
+            if (iter->second._isAmManaged) {
+                if (iter->second._isInDeviceMem) {
+                    *deviceMemSize += sizeBytes;
+                } else {
+                    *hostMemSize += sizeBytes;
+                }
+            } else {
+                *userMemSize += sizeBytes;
+            }
+        }
+    }
+
+    g_amPointerTracker.readerUnlock();
 }
 
 

hipamd/src/hip_hcc.cpp

@@ -1286,6 +1286,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr)
 
     hipError_t e = hipSuccess;
 
+#if USE_AM_TRACKER
     hc::AmPointerInfo amPointerInfo;
     am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr);
     if (status == AM_SUCCESS) {
@@ -1309,11 +1310,15 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr)
 
         e = hipErrorInvalidValue;
     }
+#else
+    e = hipErrorInvalidValue;
+#endif
 
     return ihipLogStatus(e);
 }
 
 
+#if USE_AM_TRACKER
 // TODO - test this function:
 /**
  * @returns #hipSuccess, 
@@ -1342,6 +1347,7 @@ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsi
 
     return ihipLogStatus(e);
 }
+#endif
 
 
 
@@ -1438,7 +1444,9 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
         if (sizeBytes && (*ptr == NULL)) {
             hip_status = hipErrorMemoryAllocation;
         } else {
+#ifdef USE_AM_TRACKER
             hc::am_memtracker_update(*ptr, device->_device_index, 0);
+#endif
         }
     } else {
         hip_status = hipErrorMemoryAllocation;
@@ -1462,7 +1470,9 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
         if (sizeBytes && (*ptr == NULL)) {
             hip_status = hipErrorMemoryAllocation;
         } else {
+#ifdef USE_AM_TRACKER
             hc::am_memtracker_update(*ptr, device->_device_index, 0);
+#endif
         }
 
         tprintf (TRACE_MEM, "  %s: pinned ptr=%p\n", __func__, *ptr);
@@ -1627,10 +1637,10 @@ hipError_t hipMemset(void* dst, int  value, size_t sizeBytes )
 
 
 /*
- * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue (if free != NULL due to bug)
- * @bug - on hcc free always returns 50% of peak regardless of current allocations.  hipMemGetInfo returns hipErrorInvalidValue to indicate this.
+ * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue (if free != NULL due to bug)S
+ * @warning On HCC, the free memory only accounts for memory allocated by this process and may be optimistic.
  */
-hipError_t hipMemGetInfo  (   size_t *    free, size_t *    total    )
+hipError_t hipMemGetInfo  (size_t *free, size_t *total)
 {
     std::call_once(hip_initialized, ihipInit);
 
@@ -1643,17 +1653,22 @@ hipError_t hipMemGetInfo  (   size_t *    free, size_t *    total    )
         }
 
         if (free) {
-            *free =  hipDevice->_props.totalGlobalMem  * 0.5; // TODO
+#if USE_AM_TRACKER
+            // TODO - replace with kernel-level for reporting free memory:
+            size_t deviceMemSize, hostMemSize, userMemSize;
+            hc::am_memtracker_sizeinfo(hipDevice->_acc, &deviceMemSize, &hostMemSize, &userMemSize);
+            *free =  hipDevice->_props.totalGlobalMem - deviceMemSize;
+#else
+            *free =  hipDevice->_props.totalGlobalMem * 0.5; // TODO
             e=hipErrorInvalidValue;
+#endif
         }
 
     } else {
         e = hipErrorInvalidDevice;
     }
 
-    // TODO-runtime  - when we fix the 50% bug.
-    //return ihipLogStatus(hipErrorSuccess);
-    return ihipLogStatus(hipErrorInvalidValue);
+    return ihipLogStatus(e);
 }
 
 

hipamd/tests/src/hipPointerAttrib.cpp

@@ -115,6 +115,11 @@ void testSimple()
     HIPCHECK ( hipMallocHost(&A_Pinned_h, Nbytes) );
     A_OSAlloc_h = (char*)malloc(Nbytes);
 
+    size_t free, total;
+    HIPCHECK(hipMemGetInfo(&free, &total));
+    printf ("hipMemGetInfo: free=%zu (%4.2f) Nbytes=%lu total=%zu (%4.2f)\n", free, (float)(free/1024.0/1024.0), Nbytes, total, (float)(total/1024.0/1024.0));
+    HIPASSERT(free + Nbytes <= total);
+
 
     hipPointerAttribute_t attribs;
     hipPointerAttribute_t attribs2;
@@ -244,6 +249,10 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize)
 
     //---
     //Populate with device and host allocations.
+    size_t totalDeviceAllocated[numDevices];
+    for (int i =0; i<numDevices; i++) {
+        totalDeviceAllocated[i] = 0;
+    }
     for (int i=0; i<numAllocs; i++) {
         bool isDevice = rand() & 0x1;
         reference[i]._sizeBytes = zrand(maxSize-minSize) + minSize;
@@ -254,6 +263,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize)
 
         void * ptr;
         if (isDevice) {
+            totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes;
             HIPCHECK(hipMalloc(&ptr, reference[i]._sizeBytes));
             reference[i]._attrib.memoryType    = hipMemoryTypeDevice;
             reference[i]._attrib.devicePointer = ptr;
@@ -277,6 +287,15 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize)
 #endif
 
 
+    for (int i =0; i<numDevices; i++) {
+        size_t free, total;
+        HIPCHECK(hipMemGetInfo(&free, &total));
+        printf ("  device#%d: hipMemGetInfo: free=%zu (%4.2fMB) clusterAllocTotalDevice=%lu (%4.2fMB) total=%zu (%4.2fMB)\n", 
+                i, free, (float)(free/1024.0/1024.0), totalDeviceAllocated[i], (float)(totalDeviceAllocated[i])/1024.0/1024.0, total, (float)(total/1024.0/1024.0));
+        HIPASSERT(free + totalDeviceAllocated[i] <= total);
+    }
+
+
     // Now look up each pointer we inserted and verify we can find it:
     for (int i=0; i<numAllocs; i++) {
         SuperPointerAttribute &ref = reference[i];
@@ -296,6 +315,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize)
 
 
 
+
 #ifdef __HIP_PLATFORM_HCC__
     if (p_verbose & 0x2) {
         printf ("Tracker after cleanup:\n");
@@ -370,18 +390,24 @@ void thread_noise_generator(int iters, size_t numBuffers, Dir addDir, Dir remove
         auto acc = hc::accelerator();
 
         if (addDir == Up) {
-            for (char *p = basePtr; p<basePtr + maxSize; p+=bufferSize) 
-            {
+            for (char *p = basePtr; p<basePtr + maxSize; p+=bufferSize) {
+                hc::am_memtracker_add(p, bufferSize, acc, false);
+            }
+        } else if (addDir == Down) {
+            for (char *p = basePtr+maxSize-bufferSize; p>=0; p-=bufferSize) {
                 hc::am_memtracker_add(p, bufferSize, acc, false);
             }
         }
 
         if (removeDir == Up) {
-            for (char *p = basePtr; p<basePtr + maxSize; p+=bufferSize) 
-            {
+            for (char *p = basePtr; p<basePtr + maxSize; p+=bufferSize) {
                 hc::am_memtracker_remove(p);
             }
-        };
+        } else if (removeDir == Down) {
+            for (char *p = basePtr+maxSize-bufferSize; p>=0; p-=bufferSize)  {
+                hc::am_memtracker_remove(p);
+            }
+        }
     }
 }