diff --git a/hipamd/include/hcc_detail/AM.h b/hipamd/include/hcc_detail/AM.h index d41fed317a..c183844869 100644 --- a/hipamd/include/hcc_detail/AM.h +++ b/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 diff --git a/hipamd/src/hc_AM.cpp b/hipamd/src/hc_AM.cpp index 92310164c0..3a6d116261 100644 --- a/hipamd/src/hc_AM.cpp +++ b/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(); } diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index e63186692c..fe273aa21c 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/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); } diff --git a/hipamd/tests/src/hipPointerAttrib.cpp b/hipamd/tests/src/hipPointerAttrib.cpp index 93d503af65..1418997274 100644 --- a/hipamd/tests/src/hipPointerAttrib.cpp +++ b/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=0; p-=bufferSize) { hc::am_memtracker_add(p, bufferSize, acc, false); } } if (removeDir == Up) { - for (char *p = basePtr; p=0; p-=bufferSize) { + hc::am_memtracker_remove(p); + } + } } }