From f8f40e07bf77c1aadb6df5574015b840f78ff262 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 15 Feb 2016 21:16:00 -0600 Subject: [PATCH] Update before checkin to HCC. Add support for USE_AM_TRACKER=2 (HCC version). Add AM_ALLOC, AM_FREE indirection to ease swapping AM implementations. [ROCm/clr commit: 1ed431c0f66791f7b0df78cf57e5cc8f6bf83427] --- projects/clr/hipamd/include/hcc_detail/AM.h | 23 +++---- projects/clr/hipamd/src/hc_AM.cpp | 2 - projects/clr/hipamd/src/hip_hcc.cpp | 69 ++++++++++++++++----- 3 files changed, 66 insertions(+), 28 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/AM.h b/projects/clr/hipamd/include/hcc_detail/AM.h index 40d9ea8382..74542789af 100644 --- a/projects/clr/hipamd/include/hcc_detail/AM.h +++ b/projects/clr/hipamd/include/hcc_detail/AM.h @@ -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. **/ diff --git a/projects/clr/hipamd/src/hc_AM.cpp b/projects/clr/hipamd/src/hc_AM.cpp index 221322f4b0..272024cfe7 100644 --- a/projects/clr/hipamd/src/hc_AM.cpp +++ b/projects/clr/hipamd/src/hc_AM.cpp @@ -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 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 l (_mutex); auto iter = _tracker.find(AmMemoryRange(pointer,1)); mprintf ("find: %p\n", pointer); diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index d4a6857559..30a0d993e0 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -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 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 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);