Merge 'master' into 'amd-master'

Change-Id: I815c98287213e742673caafec084cbd5ed11b1c7

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"

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<hsa_signal_t*> (_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);

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));
 };
 
 //=============================================================================

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);

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

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.