AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Update before checkin to HCC.

Browse Source 
Add support for USE_AM_TRACKER=2 (HCC version).
Add AM_ALLOC, AM_FREE indirection to ease swapping AM implementations.


[ROCm/clr commit: 1ed431c0f6]
This commit is contained in:
Ben Sander
2016-02-15 21:16:00 -06:00
parent 93c07bc3d1
commit f8f40e07bf
3 changed files with 66 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hcc_detail/AM.h
+12 -11
View File
@@ -13,7 +13,7 @@ typedef int am_status_t;
namespace hc {
// This is the data that is maintained for each pointer:
// 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.
@@ -45,7 +45,7 @@ namespace hc {
/**
* Allocates a block of @p size bytes of memory on the specified @p acc.
* 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.
*
@@ -53,7 +53,7 @@ namespace hc {
*
* Flags must be 0.
*
* @returns : On success, pointer to the newly allocated memory is returned.
* @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.
@@ -63,17 +63,18 @@ namespace hc {
auto_voidp AM_alloc(size_t size, hc::accelerator acc, unsigned flags);
/**
* Frees a block of memory previously allocated with am_alloc.
* 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);
/**
* Copies @p size bytes of memory from @p src to @ dst. The memory areas (src+size and dst+size) must not overlap.
* Copy @p size bytes of memory from @p src to @ dst. The memory areas (src+size and dst+size) must not overlap.
*
* @returns AM_SUCCESS on error or AM_ERROR_MISC if an error occurs.
* @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);
@@ -96,7 +97,7 @@ am_status_t am_memtracker_getinfo(hc::AmPointerInfo *info, const void *ptr);
/**
* Adds a pointer to the memory tracker.
* Add a pointer to the memory tracker.
*
* @return AM_SUCCESS
* @see am_memtracker_getinfo
@@ -105,7 +106,7 @@ am_status_t am_memtracker_add(void* ptr, size_t sizeBytes, hc::accelerator acc,
/*
* Updates infor for an existing pointer in the memory tracker.
* 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.
@@ -116,7 +117,7 @@ am_status_t am_memtracker_update(const void* ptr, int appId, unsigned allocation
/**
* Remove the pointer from the tracker structure.
* Remove @ptr from the tracker structure.
*
* @p ptr may be anywhere in a tracked memory range.
*
@@ -136,7 +137,7 @@ am_status_t am_memtracker_remove(void* ptr);
size_t am_memtracker_reset(hc::accelerator acc);
/**
* Prints the entries in the memory tracker table.
* Print the entries in the memory tracker table.
*
* Intended primarily for debug purposes.
* @see am_memtracker_getinfo
@@ -145,7 +146,7 @@ void am_memtracker_print();
/**
* Returns total sizes of device, host, and user memory allocated by the application
* Return total sizes of device, host, and user memory allocated by the application
*
* User memory is registered with am_tracker_add.
**/
projects/clr/hipamd/src/hc_AM.cpp
-2
View File
@@ -93,7 +93,6 @@ void AmPointerTracker::insert (void *pointer, const hc::AmPointerInfo &p)
// Return 1 if removed or 0 if not found.
int AmPointerTracker::remove (void *pointer)
{
// TODO-mutex - write lock.
std::lock_guard<std::mutex> l (_mutex);
mprintf ("remove: %p\n", pointer);
return _tracker.erase(AmMemoryRange(pointer,1));
@@ -103,7 +102,6 @@ int AmPointerTracker::remove (void *pointer)
//---
AmPointerTracker::MapTrackerType::iterator AmPointerTracker::find (const void *pointer)
{
// TODO-mutex- read lock
std::lock_guard<std::mutex> l (_mutex);
auto iter = _tracker.find(AmMemoryRange(pointer,1));
mprintf ("find: %p\n", pointer);
projects/clr/hipamd/src/hip_hcc.cpp
+54 -15
View File
@@ -40,10 +40,18 @@ THE SOFTWARE.
#include "hsa_ext_amd.h"
#include "hc_AM.cpp"
#define USE_ASYNC_COPY 1
#define USE_AM_TRACKER 1 /* use new AM memory tracker features */
#define USE_AM_TRACKER 2 /* >0 = use new AM memory tracker features. 1= use HIP impl, 2=use HCC impl */
#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
@@ -1504,7 +1512,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
if (device) {
const unsigned am_flags = 0;
*ptr = hc::AM_alloc(sizeBytes, device->_acc, am_flags);
*ptr = AM_ALLOC(sizeBytes, device->_acc, am_flags);
if (sizeBytes && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
@@ -1531,7 +1539,7 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
auto device = ihipGetTlsDefaultDevice();
if (device) {
*ptr = hc::AM_alloc(sizeBytes, device->_acc, am_flags);
*ptr = AM_ALLOC(sizeBytes, device->_acc, am_flags);
if (sizeBytes && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
} else {
@@ -1577,7 +1585,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] = hc::AM_alloc(_bufferSize, device->_acc, amHostPinned);
_pinnedStagingBuffer[i] = AM_ALLOC(_bufferSize, device->_acc, amHostPinned);
if (_pinnedStagingBuffer[i] == NULL) {
throw;
}
@@ -1590,7 +1598,7 @@ StagingBuffer::~StagingBuffer()
{
for (int i=0; i<_numBuffers; i++) {
if (_pinnedStagingBuffer[i]) {
hc::AM_free(_pinnedStagingBuffer[i]);
AM_FREE(_pinnedStagingBuffer[i]);
_pinnedStagingBuffer[i] = NULL;
}
hsa_signal_destroy(_completion_signal[i]);
@@ -1695,8 +1703,7 @@ void StagingBuffer::CopyDeviceToHost(void* dst, const void* src, size_t sizeByte
#if USE_AM_TRACKER
// TODO - add mutex to limit in/out:
void ihipAsyncCopy(ihipDevice_t *device, void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind)
void ihipSyncCopy(ihipDevice_t *device, void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind)
{
hc::AmPointerInfo dstPtrInfo, srcPtrInfo;
@@ -1725,14 +1732,16 @@ void ihipAsyncCopy(ihipDevice_t *device, void* dst, const void* src, size_t size
std::lock_guard<std::mutex> l (device->_copy_lock[0]);
device->_staging_buffer[0]->CopyHostToDevice(dst, src, sizeBytes);
} else {
hc::AM_copy(dst, src, sizeBytes);
// TODO - remove, slow path.
hc::am_copy(dst, src, sizeBytes);
}
} else if ((kind == hipMemcpyDeviceToHost) && (dstNotTracked)) {
if (useStagingBuffer) {
std::lock_guard<std::mutex> l (device->_copy_lock[1]);
device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes);
} else {
hc::AM_copy(dst, src, sizeBytes);
// TODO - remove, slow path.
hc::am_copy(dst, src, sizeBytes);
}
} else if (kind == hipMemcpyHostToHost) {
memcpy(dst, src, sizeBytes); // TODO - not async.
@@ -1757,6 +1766,36 @@ void ihipAsyncCopy(ihipDevice_t *device, void* dst, const void* src, size_t size
#endif
#if 0 // USE_AM_TRACKER
void ihipAsyncCopy(ihipDevice_t *device, void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind)
{
bool useStagingBuffer = true; // TODO - remove when new copy bakes a bit.
hipStatus_t e = hipSuccess;
// TODO - check kind is not default.
if (kind == hipMemcpyDefault) {
e = hipErrorInvalidMemoryDirection;
} else {
// Let HSA runtime handle it:
// TODO - need buffer pool for the signals:
device->_copy_lock[1].lock();
hsa_signal_store_relaxed(device->_copy_signal, 1);
hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, src, sizeBytes, device->_hsa_agent, 0, NULL, device->_copy_signal);
if (hsa_status == HSA_STATUS_SUCCESS) {
hsa_signal_wait_relaxed(device->_copy_signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
}
device->_copy_lock[1].unlock();
}
}
#endif
//---
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind)
@@ -1775,7 +1814,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
ihipDevice_t *device = &g_devices[stream->_device_index];
ihipAsyncCopy(device, dst, src, sizeBytes, kind);
ihipSyncCopy(device, dst, src, sizeBytes, kind);
} else {
e = hipErrorInvalidResourceHandle;
@@ -1784,7 +1823,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
@@ -1815,7 +1854,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
@@ -1938,7 +1977,7 @@ hipError_t hipFree(void* ptr)
ihipWaitAllStreams(ihipGetTlsDefaultDevice());
if (ptr) {
hc::AM_free(ptr);
AM_FREE(ptr);
}
return ihipLogStatus(hipSuccess);
@@ -1952,7 +1991,7 @@ hipError_t hipFreeHost(void* ptr)
if (ptr) {
tprintf (TRACE_MEM, " %s: %p\n", __func__, ptr);
hc::AM_free(ptr);
AM_FREE(ptr);
}
return ihipLogStatus(hipSuccess);