From 6db7569dc4b7ccd404aab93e84dba3a2bdbe5185 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 24 Aug 2017 15:05:45 +0530 Subject: [PATCH 1/3] Add a basic build+install script Change-Id: Ida0702057fda3cbbe9b4674e07aecacb1004fc90 --- hipamd/install.sh | 48 +++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 48 insertions(+) create mode 100755 hipamd/install.sh diff --git a/hipamd/install.sh b/hipamd/install.sh new file mode 100755 index 0000000000..f8ad640798 --- /dev/null +++ b/hipamd/install.sh @@ -0,0 +1,48 @@ +#!/bin/bash + +BUILD_ROOT="$( mktemp -d )" +SRC_ROOT="$( cd "$( dirname "${BASH_SOURCE[0]}" )" && pwd )" +WORKING_DIR=$PWD +DASH_JAY="-j $(getconf _NPROCESSORS_ONLN)" + +err() { + echo "${1-Died}." >&2 +} + +die() { + err "$1" + exit 1 +} + +pushd () { + command pushd "$@" > /dev/null +} + +popd () { + command popd "$@" > /dev/null +} + +function setupENV() +{ + sudo apt-get update + sudo apt-get install dpkg-dev rpm doxygen libelf-dev +} + +function buildHIP() +{ + pushd $BUILD_ROOT + cmake $SRC_ROOT -DCMAKE_BUILD_TYPE=Release -DCOMPILE_HIP_ATP_MARKER=1 + make $DASH_JAY + make package + rename -v 's/([a-z0-9_.\-]).deb/$1-amd64.deb/' *.deb;rename -v 's/([a-z0-9_.\-]).rpm/$1.x86_64.rpm/' *.rpm + cp hip_*.deb $WORKING_DIR + sudo dpkg -i hip_base*.deb hip_hcc*.deb hip_sample*.deb hip_doc*.deb + popd + rm -rf $BUILD_ROOT +} + +echo "Preparing build environment" +setupENV || die "setupENV failed" +echo "Building and installing HIP packages" +buildHIP || die "buildHIP failed" +echo "Finished building HIP packages" From 765d7f3c7b83ff1d8ade77f78f770af41ca199f0 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 25 Aug 2017 08:46:34 +0530 Subject: [PATCH 2/3] Null check on input pointer arguments --- hipamd/src/hip_memory.cpp | 77 +++++++++++++++++++-------------------- 1 file changed, 38 insertions(+), 39 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 6db3d82376..4b3a4fcb12 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -150,45 +150,45 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, const void HIP_INIT_API(attributes, ptr); hipError_t e = hipSuccess; - - hc::accelerator acc; -#if (__hcc_workweek__ >= 17332) - hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); -#else - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); -#endif - am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); - if (status == AM_SUCCESS) { - - attributes->memoryType = amPointerInfo._isInDeviceMem ? hipMemoryTypeDevice: hipMemoryTypeHost; - attributes->hostPointer = amPointerInfo._hostPointer; - attributes->devicePointer = amPointerInfo._devicePointer; - attributes->isManaged = 0; - if(attributes->memoryType == hipMemoryTypeHost){ - attributes->hostPointer = (void*)ptr; - } - if(attributes->memoryType == hipMemoryTypeDevice){ - attributes->devicePointer = (void*)ptr; - } - attributes->allocationFlags = amPointerInfo._appAllocationFlags; - attributes->device = amPointerInfo._appId; - - if (attributes->device < 0) { - e = hipErrorInvalidDevice; - } - - + if((attributes == nullptr) || (ptr == nullptr)) { + e = hipErrorInvalidValue; } else { - attributes->memoryType = hipMemoryTypeDevice; - attributes->hostPointer = 0; - attributes->devicePointer = 0; - attributes->device = -1; - attributes->isManaged = 0; - attributes->allocationFlags = 0; + hc::accelerator acc; +#if (__hcc_workweek__ >= 17332) + hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); +#else + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); +#endif + am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr); + if (status == AM_SUCCESS) { - e = hipErrorUnknown; // TODO - should be hipErrorInvalidValue ? + attributes->memoryType = amPointerInfo._isInDeviceMem ? hipMemoryTypeDevice: hipMemoryTypeHost; + attributes->hostPointer = amPointerInfo._hostPointer; + attributes->devicePointer = amPointerInfo._devicePointer; + attributes->isManaged = 0; + if(attributes->memoryType == hipMemoryTypeHost){ + attributes->hostPointer = (void*)ptr; + } + if(attributes->memoryType == hipMemoryTypeDevice){ + attributes->devicePointer = (void*)ptr; + } + attributes->allocationFlags = amPointerInfo._appAllocationFlags; + attributes->device = amPointerInfo._appId; + + if (attributes->device < 0) { + e = hipErrorInvalidDevice; + } + } else { + attributes->memoryType = hipMemoryTypeDevice; + attributes->hostPointer = 0; + attributes->devicePointer = 0; + attributes->device = -1; + attributes->isManaged = 0; + attributes->allocationFlags = 0; + + e = hipErrorUnknown; // TODO - should be hipErrorInvalidValue ? + } } - return ihipLogStatus(e); } @@ -199,13 +199,12 @@ hipError_t hipHostGetDevicePointer(void **devicePointer, void *hostPointer, unsi hipError_t e = hipSuccess; - *devicePointer = NULL; - // Flags must be 0: - if (flags != 0) { + if ((flags != 0) || (devicePointer == nullptr) || (hostPointer == nullptr)){ e = hipErrorInvalidValue; } else { hc::accelerator acc; + *devicePointer = NULL; #if (__hcc_workweek__ >= 17332) hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0); #else From bc9ba7cd81e9d9ae77c9e21c69d55e1df95e8211 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 26 Aug 2017 14:39:14 +0000 Subject: [PATCH 3/3] Lock streams when waiting on event completion or querying event safety. --- hipamd/src/hip_event.cpp | 23 ++++++++++++++++------- hipamd/src/hip_hcc.cpp | 26 +++++++++++++++++++++++--- hipamd/src/hip_hcc_internal.h | 15 ++++++++++++--- hipamd/src/hip_stream.cpp | 2 +- 4 files changed, 52 insertions(+), 14 deletions(-) diff --git a/hipamd/src/hip_event.cpp b/hipamd/src/hip_event.cpp index 3a8f1ab611..d1ee37a45e 100644 --- a/hipamd/src/hip_event.cpp +++ b/hipamd/src/hip_event.cpp @@ -55,13 +55,13 @@ void ihipEvent_t::attachToCompletionFuture(const hc::completion_future *cf, void ihipEvent_t::refereshEventStatus() { - bool isReady0 = _marker.is_ready(); + bool isReady0 = locked_isReady(); bool isReady1; int val = 0; if (_state == hipEventStatusRecording) { // TODO - use completion-future functions to obtain ticks and timestamps: hsa_signal_t *sig = static_cast (_marker.get_native_handle()); - isReady1 = _marker.is_ready(); + isReady1 = locked_isReady(); if (sig) { val = hsa_signal_load_acquire(*sig); if (val == 0) { @@ -86,6 +86,17 @@ void ihipEvent_t::refereshEventStatus() } +bool ihipEvent_t::locked_isReady() +{ + return _stream->locked_eventIsReady(this); +} + +void ihipEvent_t::locked_waitComplete(hc::hcWaitMode waitMode) +{ + return _stream->locked_eventWaitComplete(this, waitMode); +} + + hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) { hipError_t e = hipSuccess; @@ -127,7 +138,7 @@ hipError_t hipEventCreate(hipEvent_t* event) hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { - HIP_INIT_SPECIAL_API(TRACE_QUERY, event, stream); + HIP_INIT_SPECIAL_API(TRACE_SYNC, event, stream); if (event && event->_state != hipEventStatusUnitialized) { stream = ihipSyncAndResolveStream(stream); @@ -192,9 +203,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) ctx->locked_syncDefaultStream(true, true); return ihipLogStatus(hipSuccess); } else { - event->_marker.wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); - - assert (event->_marker.is_ready()); + event->locked_waitComplete((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); return ihipLogStatus(hipSuccess); } @@ -259,7 +268,7 @@ hipError_t hipEventQuery(hipEvent_t event) { HIP_INIT_SPECIAL_API(TRACE_QUERY, event); - if ((event->_state == hipEventStatusRecording) && (!event->_marker.is_ready())) { + if ((event->_state == hipEventStatusRecording) && !event->locked_isReady()) { return ihipLogStatus(hipErrorNotReady); } else { return ihipLogStatus(hipSuccess); diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index f7082e7e9e..dc72714e3e 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -328,14 +328,34 @@ void ihipStream_t::locked_wait() // Causes current stream to wait for specified event to complete: // Note this does not provide any kind of host serialization. -void ihipStream_t::locked_waitEvent(hipEvent_t event) +void ihipStream_t::locked_streamWaitEvent(hipEvent_t event) { LockedAccessor_StreamCrit_t crit(_criticalData); - crit->_av.create_blocking_marker(event->_marker, hc::accelerator_scope); + crit->_av.create_blocking_marker(event->marker(), hc::accelerator_scope); } + +// Causes current stream to wait for specified event to complete: +// Note this does not provide any kind of host serialization. +bool ihipStream_t::locked_eventIsReady(hipEvent_t event) +{ + // Event query that returns "Complete" may cause HCC to manipulate + // internal queue state so lock the stream's queue here. + LockedAccessor_StreamCrit_t crit(_criticalData); + + return (event->marker().is_ready()); +} + +void ihipStream_t::locked_eventWaitComplete(hipEvent_t event, hc::hcWaitMode waitMode) +{ + LockedAccessor_StreamCrit_t crit(_criticalData); + + event->marker().wait(waitMode); +} + + // Create a marker in this stream. // Save state in the event so it can track the status of the event. void ihipStream_t::locked_recordEvent(hipEvent_t event) @@ -354,7 +374,7 @@ void ihipStream_t::locked_recordEvent(hipEvent_t event) scopeFlag = HIP_EVENT_SYS_RELEASE ? hc::system_scope : hc::accelerator_scope; } - event->_marker = crit->_av.create_marker(scopeFlag); + event->marker(crit->_av.create_marker(scopeFlag)); }; //============================================================================= diff --git a/hipamd/src/hip_hcc_internal.h b/hipamd/src/hip_hcc_internal.h index af5e7a121b..56ca37b3e2 100644 --- a/hipamd/src/hip_hcc_internal.h +++ b/hipamd/src/hip_hcc_internal.h @@ -517,9 +517,12 @@ public: hc::accelerator_view* locked_getAv() { LockedAccessor_StreamCrit_t crit(_criticalData); return &(crit->_av); }; - void locked_waitEvent(hipEvent_t event); + void locked_streamWaitEvent(hipEvent_t event); void locked_recordEvent(hipEvent_t event); + bool locked_eventIsReady(hipEvent_t event); + void locked_eventWaitComplete(hipEvent_t event, hc::hcWaitMode waitMode); + ihipStreamCritical_t &criticalData() { return _criticalData; }; //--- @@ -608,18 +611,24 @@ public: ihipEvent_t(unsigned flags); void attachToCompletionFuture(const hc::completion_future *cf, hipStream_t stream, ihipEventType_t eventType); void refereshEventStatus(); + hc::completion_future & marker() { return _marker; } + void marker(hc::completion_future cf) { _marker = cf; }; + + bool locked_isReady(); + void locked_waitComplete(hc::hcWaitMode waitMode); + uint64_t timestamp() const { return _timestamp; } ; ihipEventType_t type() const { return _type; }; public: hipEventStatus_t _state; - hipStream_t _stream; // Stream where the event is recorded, or NULL if all streams. + hipStream_t _stream; // Stream where the event is recorded. Null stream is resolved to actual stream when recorded unsigned _flags; - hc::completion_future _marker; private: + hc::completion_future _marker; ihipEventType_t _type; uint64_t _timestamp; // store timestamp, may be set on host or by marker. friend hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream); diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index f0e9283201..51aeb01412 100644 --- a/hipamd/src/hip_stream.cpp +++ b/hipamd/src/hip_stream.cpp @@ -96,7 +96,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int if (stream != hipStreamNull) { // This will user create_blocking_marker to wait on the specified queue. - stream->locked_waitEvent(event); + stream->locked_streamWaitEvent(event); } else { // TODO-hcc Convert to use create_blocking_marker(...) functionality.