From 8ed32daefab2725286b591c118b474eba54ca056 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 17 Feb 2016 21:33:32 -0600 Subject: [PATCH] Remove HIP-local AM tracker (now in HCC) [ROCm/clr commit: b08e468c0661df5d2076d15fce654a222025da91] --- projects/clr/hipamd/include/hcc_detail/AM.h | 157 ---------- projects/clr/hipamd/src/hc_AM.cpp | 319 -------------------- projects/clr/hipamd/src/hip_hcc.cpp | 28 +- 3 files changed, 10 insertions(+), 494 deletions(-) delete mode 100644 projects/clr/hipamd/include/hcc_detail/AM.h delete mode 100644 projects/clr/hipamd/src/hc_AM.cpp diff --git a/projects/clr/hipamd/include/hcc_detail/AM.h b/projects/clr/hipamd/include/hcc_detail/AM.h deleted file mode 100644 index 74542789af..0000000000 --- a/projects/clr/hipamd/include/hcc_detail/AM.h +++ /dev/null @@ -1,157 +0,0 @@ -#pragma once - -#include - -typedef int am_status_t; -#define AM_SUCCESS 0 -// TODO - provide better mapping of HSA error conditions to HC error codes. -#define AM_ERROR_MISC -1 /** Misellaneous error */ - -// Flags for am_alloc API: -#define amHostPinned 0x1 - - -namespace hc { - -// Info for each pointer in the memtry tracker: -struct AmPointerInfo { - void * _hostPointer; ///< Host pointer. If host access is not allowed, NULL. - void * _devicePointer; ///< Device pointer. - size_t _sizeBytes; ///< Size of allocation. - hc::accelerator _acc; ///< Device / Accelerator to use. - bool _isInDeviceMem; ///< Memory is physically resident on a device (if false, memory is located on host) - bool _isAmManaged; ///< Memory was allocated by AM and should be freed when am_reset is called. - - int _appId; ///< App-specific storage. (Used by HIP to store deviceID.) - unsigned _appAllocationFlags; ///< App-specific allocation flags. (Used by HIP to store allocation flags.) - - AmPointerInfo() {}; - - AmPointerInfo(void *hostPointer, void *devicePointer, size_t sizeBytes, hc::accelerator acc, bool isInDeviceMem, bool isAmManaged) : - _hostPointer(hostPointer), - _devicePointer(devicePointer), - _sizeBytes(sizeBytes), - _acc(acc), - _isInDeviceMem(isInDeviceMem), - _isAmManaged(isAmManaged), - _appId(-1), - _appAllocationFlags(0) {}; -}; -} - - - -namespace hc { - - -/** - * Allocate a block of @p size bytes of memory on the specified @p acc. - * - * The contents of the newly allocated block of memory are not initialized. - * - * If @p size == 0, 0 is returned. - * - * Flags must be 0. - * - * @return : On success, pointer to the newly allocated memory is returned. - * The pointer is typecast to the desired return type. - * - * If an error occurred trying to allocate the requested memory, 0 is returned. - * - * @see am_free, am_copy - */ -auto_voidp AM_alloc(size_t size, hc::accelerator acc, unsigned flags); - -/** - * Free a block of memory previously allocated with am_alloc. - * - * @return AM_SUCCESS - * @see am_alloc, am_copy - */ -am_status_t AM_free(void* ptr); - - -/** - * Copy @p size bytes of memory from @p src to @ dst. The memory areas (src+size and dst+size) must not overlap. - * - * @return AM_SUCCESS on error or AM_ERROR_MISC if an error occurs. - * @see am_alloc, am_free - */ -am_status_t AM_copy(void* dst, const void* src, size_t size); - - - -/** - * Return information about tracked pointer. - * - * AM tracks pointers when they are allocated or added to tracker with am_track_pointer. - * The tracker tracks the base pointer as well as the size of the allocation, and will - * find the information for a pointer anywhere in the tracked range. - * - * @returns AM_ERROR_MISC if pointer is not currently being tracked. - * @returns AM_SUCCESS if pointer is tracked and writes info to @p info. - * - * @see AM_memtracker_add, - */ -am_status_t am_memtracker_getinfo(hc::AmPointerInfo *info, const void *ptr); - - -/** - * Add a pointer to the memory tracker. - * - * @return AM_SUCCESS - * @see am_memtracker_getinfo - */ -am_status_t am_memtracker_add(void* ptr, size_t sizeBytes, hc::accelerator acc, bool isDeviceMem=false); - - -/* - * Update info for an existing pointer in the memory tracker. - * - * @returns AM_ERROR_MISC if pointer is not found in tracker. - * @returns AM_SUCCESS if pointer is not found in tracker. - * - * @see am_memtracker_getinfo, am_memtracker_add - */ -am_status_t am_memtracker_update(const void* ptr, int appId, unsigned allocationFlags); - - -/** - * Remove @ptr from the tracker structure. - * - * @p ptr may be anywhere in a tracked memory range. - * - * @returns AM_ERROR_MISC if pointer is not found in tracker. - * @returns AM_SUCCESS if pointer is not found in tracker. - * - * @see am_memtracker_getinfo, am_memtracker_add - */ -am_status_t am_memtracker_remove(void* ptr); - -/** - * Remove all memory allocations associated with specified accelerator from the memory tracker. - * - * @returns Number of entries reset. - * @see am_memtracker_getinfo - */ -size_t am_memtracker_reset(hc::accelerator acc); - -/** - * Print the entries in the memory tracker table. - * - * Intended primarily for debug purposes. - * @see am_memtracker_getinfo - **/ -void am_memtracker_print(); - - -/** - * Return total sizes of device, host, and user memory allocated by the application - * - * User memory is registered with am_tracker_add. - **/ -void am_memtracker_sizeinfo(hc::accelerator acc, size_t *deviceMemSize, size_t *hostMemSize, size_t *userMemSize); - - -}; // namespace hc - diff --git a/projects/clr/hipamd/src/hc_AM.cpp b/projects/clr/hipamd/src/hc_AM.cpp deleted file mode 100644 index 272024cfe7..0000000000 --- a/projects/clr/hipamd/src/hc_AM.cpp +++ /dev/null @@ -1,319 +0,0 @@ - -#include "hc_am.hpp" -#include "hsa.h" - - -#include "hcc_detail/AM.h" // TODO - Remove me. - -#define DB_TRACKER 0 -#define MUTEX_LOCK 1 - -#if DB_TRACKER -#define mprintf( ...) {\ - fprintf (stderr, __VA_ARGS__);\ - }; -#else -#define mprintf( ...) -#endif - -//========================================================================================================= -// Pointer Tracker Structures: -//========================================================================================================= -#include -#include -//#include - -struct AmMemoryRange { - const void * _basePointer; - const void * _endPointer; - AmMemoryRange(const void *basePointer, size_t sizeBytes) : - _basePointer(basePointer), _endPointer((const unsigned char*)basePointer + sizeBytes - 1) {}; -}; - -// Functor to compare ranges: -struct AmMemoryRangeCompare { - // Return true is LHS range is less than RHS - used to order the - bool operator()(const AmMemoryRange &lhs, const AmMemoryRange &rhs) const - { - return lhs._endPointer < rhs._basePointer; - } - -}; - - -std::ostream &operator<<(std::ostream &os, const hc::AmPointerInfo &ap) -{ - os << "hostPointer:" << ap._hostPointer << " devicePointer:"<< ap._devicePointer << " sizeBytes:" << ap._sizeBytes - << " isInDeviceMem:" << ap._isInDeviceMem << " isAmManaged:" << ap._isAmManaged - << " appId:" << ap._appId << " appAllocFlags:" << ap._appAllocationFlags; - return os; -} - - -//------------------------------------------------------------------------------------------------- -// This structure tracks information for each pointer. -// Uses memory-range-based lookups - so pointers that exist anywhere in the range of hostPtr + size -// will find the associated AmPointerInfo. -// The insertions and lookups use a self-balancing binary tree and should support O(logN) lookup speed. -// The structure is thread-safe - writers obtain a mutex before modifying the tree. Multiple simulatenous readers are supported. -class AmPointerTracker { -typedef std::map MapTrackerType; -public: - - void insert(void *pointer, const hc::AmPointerInfo &p); - int remove(void *pointer); - - MapTrackerType::iterator find(const void *hostPtr) ; - - MapTrackerType::iterator readerLockBegin() { _mutex.lock(); return _tracker.begin(); } ; - MapTrackerType::iterator end() { return _tracker.end(); } ; - void readerUnlock() { _mutex.unlock(); }; - - - size_t reset (hc::accelerator acc); - -private: - MapTrackerType _tracker; - std::mutex _mutex; - //std::shared_timed_mutex _mut; -}; - - -//--- -void AmPointerTracker::insert (void *pointer, const hc::AmPointerInfo &p) -{ - std::lock_guard l (_mutex); - - mprintf ("insert: %p + %zu\n", pointer, p._sizeBytes); - _tracker.insert(std::make_pair(AmMemoryRange(pointer, p._sizeBytes), p)); -} - - -//--- -// Return 1 if removed or 0 if not found. -int AmPointerTracker::remove (void *pointer) -{ - std::lock_guard l (_mutex); - mprintf ("remove: %p\n", pointer); - return _tracker.erase(AmMemoryRange(pointer,1)); -} - - -//--- -AmPointerTracker::MapTrackerType::iterator AmPointerTracker::find (const void *pointer) -{ - std::lock_guard l (_mutex); - auto iter = _tracker.find(AmMemoryRange(pointer,1)); - mprintf ("find: %p\n", pointer); - return iter; -} - - -//--- -// Remove all tracked locations, and free the associated memory (if the range was originally allocated by AM). -// Returns count of ranges removed. -size_t AmPointerTracker::reset (hc::accelerator acc) -{ - std::lock_guard l (_mutex); - mprintf ("reset: \n"); - - size_t count = 0; - // relies on C++11 (erase returns iterator) - for (auto iter = _tracker.begin() ; iter != _tracker.end(); ) { - if (iter->second._acc == acc) { - if (iter->second._isAmManaged) { - hsa_memory_free(const_cast (iter->first._basePointer)); - } - count++; - - iter = _tracker.erase(iter); - } else { - iter++; - } - } - - return count; -} - - -//========================================================================================================= -// Global var defs: -//========================================================================================================= -AmPointerTracker g_amPointerTracker; // Track all am pointer allocations. - - -//========================================================================================================= -// API Definitions. -//========================================================================================================= -// -// - -namespace hc { - -// Allocate accelerator memory, return NULL if memory could not be allocated: -auto_voidp AM_alloc(size_t sizeBytes, hc::accelerator acc, unsigned flags) -{ - - void *ptr = NULL; - - if (sizeBytes != 0 ) { - if (acc.is_hsa_accelerator()) { - hsa_agent_t *hsa_agent = static_cast (acc.get_default_view().get_hsa_agent()); - hsa_region_t *alloc_region; - if (flags & amHostPinned) { - alloc_region = static_cast(acc.get_hsa_am_system_region()); - } else { - alloc_region = static_cast(acc.get_hsa_am_region()); - } - - if (alloc_region->handle != -1) { - - hsa_status_t s1 = hsa_memory_allocate(*alloc_region, sizeBytes, &ptr); - hsa_status_t s2 = hsa_memory_assign_agent(ptr, *hsa_agent, HSA_ACCESS_PERMISSION_RW); - - if ((s1 != HSA_STATUS_SUCCESS) || (s2 != HSA_STATUS_SUCCESS)) { - ptr = NULL; - } else { - if (flags & amHostPinned) { - g_amPointerTracker.insert(ptr, - hc::AmPointerInfo(ptr/*hostPointer*/, ptr /*devicePointer*/, sizeBytes, acc, false/*isDevice*/, true /*isAMManaged*/)); - } else { - g_amPointerTracker.insert(ptr, - hc::AmPointerInfo(NULL/*hostPointer*/, ptr /*devicePointer*/, sizeBytes, acc, true/*isDevice*/, true /*isAMManaged*/)); - } - } - } - } - } - - return ptr; -}; - - -am_status_t AM_free(void* ptr) -{ - am_status_t status = AM_SUCCESS; - - if (ptr != NULL) { - // See also tracker::reset which can free memory. - hsa_memory_free(ptr); - - int numRemoved = g_amPointerTracker.remove(ptr) ; - if (numRemoved == 0) { - status = AM_ERROR_MISC; - } - } - return status; -} - - - -am_status_t AM_copy(void* dst, const void* src, size_t sizeBytes) -{ - am_status_t am_status = AM_ERROR_MISC; - hsa_status_t err = hsa_memory_copy(dst, src, sizeBytes); - - if (err == HSA_STATUS_SUCCESS) { - am_status = AM_SUCCESS; - } else { - am_status = AM_ERROR_MISC; - } - - return am_status; -} - - -am_status_t am_memtracker_getinfo(hc::AmPointerInfo *info, const void *ptr) -{ - auto infoI = g_amPointerTracker.find(ptr); - if (infoI != g_amPointerTracker.end()) { - *info = infoI->second; - return AM_SUCCESS; - } else { - return AM_ERROR_MISC; - } -} - -am_status_t am_memtracker_add(void* ptr, size_t sizeBytes, hc::accelerator acc, bool isDeviceMem) -{ - if (isDeviceMem) { - g_amPointerTracker.insert(ptr, hc::AmPointerInfo(ptr/*hostPointer*/, ptr /*devicePointer*/, sizeBytes, acc, true/*isDevice*/, false /*isAMManaged*/)); - } else { - g_amPointerTracker.insert(ptr, hc::AmPointerInfo(NULL/*hostPointer*/, ptr /*devicePointer*/, sizeBytes, acc, false/*isDevice*/, false /*isAMManaged*/)); - } - - return AM_SUCCESS; -} - - -am_status_t am_memtracker_update(const void* ptr, int appId, unsigned allocationFlags) -{ - auto iter = g_amPointerTracker.find(ptr); - if (iter != g_amPointerTracker.end()) { - iter->second._appId = appId; - iter->second._appAllocationFlags = allocationFlags; - return AM_SUCCESS; - } else { - return AM_ERROR_MISC; - } -} - - -am_status_t am_memtracker_remove(void* ptr) -{ - am_status_t status = AM_SUCCESS; - - int numRemoved = g_amPointerTracker.remove(ptr) ; - if (numRemoved == 0) { - status = AM_ERROR_MISC; - } - - return status; -} - -//--- -void am_memtracker_print() -{ - 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(); -} - - -//--- -size_t am_memtracker_reset(hc::accelerator acc) -{ - return g_amPointerTracker.reset(acc); -} - - -} // end namespace hc. diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 83083022c2..7a32f91747 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -42,21 +42,13 @@ THE SOFTWARE. #define USE_AM_TRACKER 0 /* >0 = use new AM memory tracker features. 2=use HCC impl */ -#define USE_ROCR_V2 0 +#define USE_ROCR_V2 0 /* use the ROCR v2 async copy API with dst and src agents */ -#if ((USE_AM_TRACKER!=0) && (USE_AM_TRACKER!=2)) -#error (USE_AM_TRACKER must be 0 or 2) +#if (USE_AM_TRACKER) and (__hcc_workweek__ < 16074) +#error (USE_AM_TRACKER requries HCC version of 16074 or newer) #endif -#if USE_AM_TRACKER==1 -#include "hc_AM.cpp" -#define AM_ALLOC hc::AM_alloc -#define AM_FREE hc::AM_free -#else -#define AM_ALLOC hc::am_alloc -#define AM_FREE hc::am_free -#endif #define INLINE static inline @@ -247,9 +239,9 @@ ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, unsig { _signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1); - auto s = this; #if 0 + auto s = this; std::for_each(_signalPool.begin(), _signalPool.end(), [s](ihipSignal_t &iter) { printf (" stream:%p allocated hsa_signal=%lu\n", s, (iter._hsa_signal.handle)); @@ -1642,7 +1634,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) if (device) { const unsigned am_flags = 0; - *ptr = AM_ALLOC(sizeBytes, device->_acc, am_flags); + *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); if (sizeBytes && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; @@ -1669,7 +1661,7 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes) auto device = ihipGetTlsDefaultDevice(); if (device) { - *ptr = AM_ALLOC(sizeBytes, device->_acc, am_flags); + *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); if (sizeBytes && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; } else { @@ -1715,7 +1707,7 @@ StagingBuffer::StagingBuffer(ihipDevice_t *device, size_t bufferSize, int numBuf for (int i=0; i<_numBuffers; i++) { // TODO - experiment with alignment here. - _pinnedStagingBuffer[i] = AM_ALLOC(_bufferSize, device->_acc, amHostPinned); + _pinnedStagingBuffer[i] = hc::am_alloc(_bufferSize, device->_acc, amHostPinned); if (_pinnedStagingBuffer[i] == NULL) { throw; } @@ -1728,7 +1720,7 @@ StagingBuffer::~StagingBuffer() { for (int i=0; i<_numBuffers; i++) { if (_pinnedStagingBuffer[i]) { - AM_FREE(_pinnedStagingBuffer[i]); + hc::am_free(_pinnedStagingBuffer[i]); _pinnedStagingBuffer[i] = NULL; } hsa_signal_destroy(_completion_signal[i]); @@ -2112,7 +2104,7 @@ hipError_t hipFree(void* ptr) ihipWaitAllStreams(ihipGetTlsDefaultDevice()); if (ptr) { - AM_FREE(ptr); + hc::am_free(ptr); } return ihipLogStatus(hipSuccess); @@ -2126,7 +2118,7 @@ hipError_t hipFreeHost(void* ptr) if (ptr) { tprintf (TRACE_MEM, " %s: %p\n", __func__, ptr); - AM_FREE(ptr); + hc::am_free(ptr); } return ihipLogStatus(hipSuccess);