AI/rocm-systems
2
0
Bifurcation 0
Code Tickets Demandes d'ajout Actions Paquets Projets Publications Wiki Activité

Remove HIP-local AM tracker (now in HCC)

Parcourir la source 
[ROCm/clr commit: b08e468c06]
Cette révision appartient à :
Ben Sander
2016-02-17 21:33:32 -06:00
Parent 98e608a5ce
révision 8ed32daefa
3 fichiers modifiés avec 10 ajouts et 494 suppressions
Télécharger le Fichier Patch Télécharger le Fichier des Différences Développer tous les fichiers Réduire tous les fichiers
projects/clr/hipamd/include/hcc_detail/AM.h
-157
Voir le fichier
@@ -1,157 +0,0 @@
#pragma once
#include <hc_am.hpp>
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
projects/clr/hipamd/src/hc_AM.cpp
-319
Voir le fichier
@@ -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 <map>
#include <iostream>
//#include <shared_mutex>
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<AmMemoryRange, hc::AmPointerInfo, AmMemoryRangeCompare> 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<std::mutex> 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<std::mutex> l (_mutex);
mprintf ("remove: %p\n", pointer);
return _tracker.erase(AmMemoryRange(pointer,1));
}
//---
AmPointerTracker::MapTrackerType::iterator AmPointerTracker::find (const void *pointer)
{
std::lock_guard<std::mutex> 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<std::mutex> 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<void*> (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<hsa_agent_t*> (acc.get_default_view().get_hsa_agent());
hsa_region_t *alloc_region;
if (flags & amHostPinned) {
alloc_region = static_cast<hsa_region_t*>(acc.get_hsa_am_system_region());
} else {
alloc_region = static_cast<hsa_region_t*>(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.
projects/clr/hipamd/src/hip_hcc.cpp
+10 -18
Voir le fichier
@@ -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);