Create address tracker for am_alloc.
Tracks device where memory is allocated, pinned-host or device, and more. 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.
Αυτή η υποβολή περιλαμβάνεται σε:
Ben Sander
@@ -0,0 +1,219 @@
|
||||
|
||||
#include "hc_am.hpp"
|
||||
#include "hsa.h"
|
||||
|
||||
|
||||
#include "hcc_detail/AM.h" // TODO - Remove me.
|
||||
|
||||
#define DB_TRACKER 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 {
|
||||
void * _basePointer;
|
||||
void * _endPointer;
|
||||
AmMemoryRange(void *basePointer, size_t sizeBytes) :
|
||||
_basePointer(basePointer), _endPointer((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
|
||||
<< " isDeviceMem:" << ap._isDeviceMem << " allocFlags:" << ap._allocationFlags;
|
||||
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(void *hostPtr);
|
||||
|
||||
MapTrackerType::iterator end() { return _tracker.end(); };
|
||||
|
||||
std::ostream & print (std::ostream &os);
|
||||
private:
|
||||
MapTrackerType _tracker;
|
||||
//std::shared_timed_mutex _mut;
|
||||
};
|
||||
|
||||
|
||||
//---
|
||||
void AmPointerTracker::insert (void *pointer, const hc::AmPointerInfo &p)
|
||||
{
|
||||
// TODO-mutex - write lock.
|
||||
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)
|
||||
{
|
||||
// TODO-mutex - write lock.
|
||||
mprintf ("remove: %p\n", pointer);
|
||||
return _tracker.erase(AmMemoryRange(pointer,1));
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
AmPointerTracker::MapTrackerType::iterator AmPointerTracker::find (void *pointer)
|
||||
{
|
||||
// TODO-mutex- read lock
|
||||
auto iter = _tracker.find(AmMemoryRange(pointer,1));
|
||||
mprintf ("find: %p\n", pointer);
|
||||
return iter;
|
||||
}
|
||||
|
||||
|
||||
std::ostream & AmPointerTracker::print (std::ostream &os)
|
||||
{
|
||||
for (auto iter = _tracker.begin() ; iter != _tracker.end(); iter++) {
|
||||
os << " " << iter->first._basePointer << "..." << iter->first._endPointer << ":: ";
|
||||
os << iter->second << std::endl;
|
||||
}
|
||||
|
||||
return os;
|
||||
}
|
||||
|
||||
|
||||
//=========================================================================================================
|
||||
// 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(false/*isDevice*/, ptr/*hostPointer*/, ptr /*devicePointer*/, sizeBytes, acc, flags));
|
||||
} else {
|
||||
g_amPointerTracker.insert(ptr,
|
||||
hc::AmPointerInfo(true/*isDevice*/, NULL/*hostPointer*/, ptr /*devicePointer*/, sizeBytes, acc, flags));
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return ptr;
|
||||
};
|
||||
|
||||
|
||||
am_status_t AM_free(void* ptr)
|
||||
{
|
||||
am_status_t status = AM_SUCCESS;
|
||||
|
||||
if (ptr != NULL) {
|
||||
hsa_memory_free(ptr);
|
||||
|
||||
size_t 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_get_pointer_info(hc::AmPointerInfo *info, void *ptr)
|
||||
{
|
||||
auto infoI = g_amPointerTracker.find(ptr);
|
||||
if (infoI != g_amPointerTracker.end()) {
|
||||
*info = infoI->second;
|
||||
return AM_SUCCESS;
|
||||
} else {
|
||||
return AM_ERROR_MISC;
|
||||
}
|
||||
}
|
||||
|
||||
void AM_print_tracker()
|
||||
{
|
||||
g_amPointerTracker.print(std::cerr);
|
||||
}
|
||||
|
||||
|
||||
} // end namespace hc.
|
||||
+62
-9
@@ -31,6 +31,8 @@ THE SOFTWARE.
|
||||
#include <list>
|
||||
#include <sys/types.h>
|
||||
#include <unistd.h>
|
||||
#include <unordered_map>
|
||||
|
||||
#include <hc.hpp>
|
||||
#include <hc_am.hpp>
|
||||
|
||||
@@ -38,6 +40,9 @@ THE SOFTWARE.
|
||||
|
||||
#include "hsa_ext_amd.h"
|
||||
|
||||
|
||||
#include "hc_AM.cpp"
|
||||
|
||||
#define USE_PINNED_HOST (__hcc_workweek__ >= 1601)
|
||||
|
||||
#define USE_ASYNC_COPY 0
|
||||
@@ -466,7 +471,8 @@ void ihipInit()
|
||||
g_devices.reserve(accs.size());
|
||||
for (int i=0; i<accs.size(); i++) {
|
||||
if (! accs[i].get_is_emulated()) {
|
||||
g_devices.emplace_back(ihipDevice_t(g_devices.size(), accs[i]));
|
||||
int deviceId = g_devices.size();
|
||||
g_devices.emplace_back(ihipDevice_t(deviceId, accs[i]));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -1262,6 +1268,53 @@ hipError_t hipEventQuery(hipEvent_t event)
|
||||
// Memory
|
||||
//
|
||||
//
|
||||
//
|
||||
|
||||
//---
|
||||
/**
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice
|
||||
*/
|
||||
hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr)
|
||||
{
|
||||
std::call_once(hip_initialized, ihipInit);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
hc::AmPointerInfo amPointerInfo;
|
||||
am_status_t status = hc::AM_get_pointer_info(&amPointerInfo, ptr);
|
||||
if (status == AM_SUCCESS) {
|
||||
|
||||
attributes->memoryType = amPointerInfo._isDeviceMem ? hipMemoryTypeDevice: hipMemoryTypeHost;
|
||||
attributes->hostPointer = amPointerInfo._hostPointer;
|
||||
attributes->devicePointer = amPointerInfo._devicePointer;
|
||||
attributes->isManaged = 0;
|
||||
attributes->allocationFlags = amPointerInfo._allocationFlags;
|
||||
|
||||
|
||||
attributes->device = -1;
|
||||
e = hipErrorInvalidDevice;
|
||||
for (int i=0; i<g_devices.size(); i++) {
|
||||
if (g_devices[i]._acc == amPointerInfo._acc) {
|
||||
attributes->device = i;
|
||||
e = hipSuccess;
|
||||
break;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
attributes->memoryType = hipMemoryTypeDevice;
|
||||
attributes->hostPointer = 0;
|
||||
attributes->devicePointer = 0;
|
||||
attributes->device = -1;
|
||||
attributes->isManaged = 0;
|
||||
attributes->allocationFlags = 0;
|
||||
|
||||
e = hipErrorInvalidValue;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
|
||||
|
||||
// kernel for launching memcpy operations:
|
||||
template <typename T>
|
||||
@@ -1345,9 +1398,9 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
|
||||
hipError_t hip_status = hipSuccess;
|
||||
|
||||
const unsigned am_flags = 0;
|
||||
*ptr = hc::am_alloc(sizeBytes, ihipGetTlsDefaultDevice()->_acc, am_flags);
|
||||
*ptr = hc::AM_alloc(sizeBytes, ihipGetTlsDefaultDevice()->_acc, am_flags);
|
||||
|
||||
if (*ptr == NULL) {
|
||||
if (sizeBytes && (*ptr == NULL)) {
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
} else {
|
||||
hip_status = hipSuccess;
|
||||
@@ -1367,9 +1420,9 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
|
||||
|
||||
const unsigned am_flags = amHostPinned;
|
||||
|
||||
*ptr = hc::am_alloc(sizeBytes, ihipGetTlsDefaultDevice()->_acc, am_flags);
|
||||
*ptr = hc::AM_alloc(sizeBytes, ihipGetTlsDefaultDevice()->_acc, am_flags);
|
||||
hipError_t hip_status = hipSuccess;
|
||||
if (*ptr == NULL) {
|
||||
if (sizeBytes && (*ptr == NULL)) {
|
||||
hip_status = hipErrorMemoryAllocation;
|
||||
} else {
|
||||
hip_status = hipSuccess;
|
||||
@@ -1444,7 +1497,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
|
||||
|
||||
#else
|
||||
// TODO-hsart - what synchronization does hsa_copy provide?
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
hc::AM_copy(dst, src, sizeBytes);
|
||||
e = hipSuccess;
|
||||
#endif
|
||||
|
||||
@@ -1475,7 +1528,7 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
|
||||
|
||||
// TODO-hsart This routine needs to ensure that dst and src are mapped on the GPU.
|
||||
// This is a synchronous copy - remove and replace with code below when we have appropriate LOCK APIs.
|
||||
hc::am_copy(dst, src, sizeBytes);
|
||||
hc::AM_copy(dst, src, sizeBytes);
|
||||
|
||||
#if 0
|
||||
|
||||
@@ -1592,7 +1645,7 @@ hipError_t hipFree(void* ptr)
|
||||
ihipWaitAllStreams(ihipGetTlsDefaultDevice());
|
||||
|
||||
if (ptr) {
|
||||
hc::am_free(ptr);
|
||||
hc::AM_free(ptr);
|
||||
}
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
@@ -1606,7 +1659,7 @@ hipError_t hipFreeHost(void* ptr)
|
||||
if (ptr) {
|
||||
#if USE_PINNED_HOST
|
||||
tprintf (TRACE_MEM, " %s: %p\n", __func__, ptr);
|
||||
hc::am_free(ptr);
|
||||
hc::AM_free(ptr);
|
||||
#else
|
||||
free(ptr);
|
||||
#endif
|
||||
|
||||
Αναφορά σε νέο ζήτημα
Block a user