Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe

# Conflicts:
#	tests/src/runtimeApi/stream/hipStreamSync2.cpp


[ROCm/clr commit: f7919a8321]

projects/clr/hipamd/src/hip_event.cpp

@@ -31,12 +31,9 @@ THE SOFTWARE.
 
 
 ihipEvent_t::ihipEvent_t(unsigned flags)
+    : _criticalData(this)
 {
-    _state  = hipEventStatusCreated;
-    _stream = NULL;
     _flags  = flags;
-    _timestamp  = 0;
-    _type   = hipEventTypeIndependent;
 };
 
 
@@ -45,56 +42,45 @@ ihipEvent_t::ihipEvent_t(unsigned flags)
 void ihipEvent_t::attachToCompletionFuture(const hc::completion_future *cf, 
                                            hipStream_t stream, ihipEventType_t eventType)
 {
-    _marker = *cf;
+    LockedAccessor_EventCrit_t crit(_criticalData);
-    _type   = eventType;
+    crit->_eventData.marker(*cf);
-    _stream = stream;
+    crit->_eventData._type   = eventType;
-    _state  = hipEventStatusRecording;
+    crit->_eventData._stream = stream;
+    crit->_eventData._state  = hipEventStatusRecording;
 }
 
 
 
-void ihipEvent_t::refereshEventStatus()
+std::pair<hipEventStatus_t, uint64_t> 
+ihipEvent_t::refreshEventStatus()
 {
-    bool isReady0 = locked_isReady();
+    auto ecd = locked_copyCrit();
-    bool isReady1;
+    if (ecd._state == hipEventStatusRecording) {
-    int val = 0;
+        bool isReady1 = ecd._stream->locked_eventIsReady(this);
-    if (_state == hipEventStatusRecording) {
+        if (isReady1) {
-        // TODO - use completion-future functions to obtain ticks and timestamps:
+            LockedAccessor_EventCrit_t eCrit(_criticalData);
-        hsa_signal_t *sig  = static_cast<hsa_signal_t*> (_marker.get_native_handle());
-        isReady1 = locked_isReady();
-        if (sig) {
-            val = hsa_signal_load_acquire(*sig);
-            if (val == 0) {
 
-                if ((_type == hipEventTypeIndependent) || (_type == hipEventTypeStopCommand)) {
+            if ((eCrit->_eventData._type == hipEventTypeIndependent) || 
-                    _timestamp =  _marker.get_end_tick();
+                (eCrit->_eventData._type == hipEventTypeStopCommand)) {
-                } else if (_type == hipEventTypeStartCommand) {
+                eCrit->_eventData._timestamp  =  eCrit->_eventData.marker().get_end_tick();
-                    _timestamp =  _marker.get_begin_tick();
+            } else if (eCrit->_eventData._type == hipEventTypeStartCommand) {
-                } else {
+                eCrit->_eventData._timestamp =  eCrit->_eventData.marker().get_begin_tick();
-                    assert(0); // TODO - move to debug assert
+            } else {
-                    _timestamp =  0;
+                eCrit->_eventData._timestamp =  0;
-                }
+                assert(0); // TODO - move to debug assert
-
-                _state = hipEventStatusComplete;
             }
+
+            eCrit->_eventData._state = hipEventStatusComplete;
+
+            return std::pair<hipEventStatus_t, uint64_t> (eCrit->_eventData._state,  eCrit->_eventData._timestamp);
         }
     } 
 
-    if (_state != hipEventStatusComplete) {
+    // Not complete path here:
-        //printf (" not ready isReady0=%d val=%d isReady1=%d\n", isReady0, val, isReady1);
+    return std::pair<hipEventStatus_t, uint64_t> (ecd._state,  ecd._timestamp); 
-    }
 }
 
 
-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)
@@ -136,33 +122,44 @@ hipError_t hipEventCreate(hipEvent_t* event)
     return ihipLogStatus(ihipEventCreate(event, 0));
 }
 
+
 hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
 {
     HIP_INIT_SPECIAL_API(TRACE_SYNC, event, stream);
 
-    if (event && event->_state != hipEventStatusUnitialized)   {
+    auto ecd = event->locked_copyCrit();
+
+    if (event && ecd._state != hipEventStatusUnitialized)   {
         stream = ihipSyncAndResolveStream(stream);
 
-        event->_stream = stream;
-
         if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
-
             // TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
-            
+            //
             // If default stream , then wait on all queues.
             ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
             ctx->locked_syncDefaultStream(true, true);
 
-            event->_timestamp = hc::get_system_ticks();
+            {
-            event->_state = hipEventStatusComplete;
+                LockedAccessor_EventCrit_t eCrit(event->criticalData());
+                eCrit->_eventData.marker(hc::completion_future()); // reset event
+                eCrit->_eventData._stream = stream;
+                eCrit->_eventData._timestamp = hc::get_system_ticks();
+                eCrit->_eventData._state = hipEventStatusComplete;
+            }
             return ihipLogStatus(hipSuccess);
         } else {
-            // Clear timestamps
-            event->_timestamp = 0;
-
             // Record the event in the stream:
-            stream->locked_recordEvent(event);
+            // Keep a copy outside the critical section so we lock stream first, then event - to avoid deadlock
-            event->_state  = hipEventStatusRecording;
+            hc::completion_future cf = stream->locked_recordEvent(event);
+
+            {
+                LockedAccessor_EventCrit_t eCrit(event->criticalData());
+                eCrit->_eventData.marker(cf);
+                eCrit->_eventData._stream = stream;
+                eCrit->_eventData._timestamp = 0;
+                eCrit->_eventData._state  = hipEventStatusRecording;
+            }
+
             return ihipLogStatus(hipSuccess);
         }
     } else {
@@ -170,15 +167,13 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
     }
 }
 
+
 hipError_t hipEventDestroy(hipEvent_t event)
 {
     HIP_INIT_API(event);
 
     if (event) {
-        event->_state  = hipEventStatusUnitialized;
-
         delete event;
-        event = NULL;
 
         return ihipLogStatus(hipSuccess);
     } else {
@@ -193,20 +188,24 @@ hipError_t hipEventSynchronize(hipEvent_t event)
     if (!(event->_flags & hipEventReleaseToSystem)) {
         tprintf(DB_WARN, "hipEventSynchronize on event without system-scope fence ; consider creating with hipEventReleaseToSystem\n");
     }
+    auto ecd = event->locked_copyCrit();
 
     if (event) {
-        if (event->_state == hipEventStatusUnitialized) {
+        if (ecd._state == hipEventStatusUnitialized) {
             return ihipLogStatus(hipErrorInvalidResourceHandle);
-        } else if (event->_state == hipEventStatusCreated ) {
+        } else if (ecd._state == hipEventStatusCreated ) {
             // Created but not actually recorded on any device:
             return ihipLogStatus(hipSuccess);
-        } else if (HIP_SYNC_NULL_STREAM && (event->_stream->isDefaultStream() )) {
+        } else if (HIP_SYNC_NULL_STREAM && (ecd._stream->isDefaultStream() )) {
             auto *ctx = ihipGetTlsDefaultCtx();
             // TODO-HIP_SYNC_NULL_STREAM - can remove this code
             ctx->locked_syncDefaultStream(true, true);
             return ihipLogStatus(hipSuccess);
         } else {
-            event->locked_waitComplete((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
+            ecd._stream->locked_eventWaitComplete(
+                             ecd.marker(), 
+                             (event->_flags & hipEventBlockingSync) ? 
+                             hc::hcWaitModeBlocked : hc::hcWaitModeActive);
 
             return ihipLogStatus(hipSuccess);
         }
@@ -223,44 +222,50 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop)
 
     *ms = 0.0f;
 
-    if ((start == nullptr) ||
+    if ((start == nullptr) || (stop == nullptr)) {
-        (start->_flags & hipEventDisableTiming) ||
-        (start->_state == hipEventStatusUnitialized) || (start->_state == hipEventStatusCreated) ||
-        (stop == nullptr) ||
-        (stop->_flags & hipEventDisableTiming) ||
-        ( stop->_state == hipEventStatusUnitialized) || ( stop->_state == hipEventStatusCreated)) {
-
-        // Both events must be at least recorded else return hipErrorInvalidResourceHandle
-
         status = hipErrorInvalidResourceHandle;
-
     } else {
-        // Refresh status, if still recording...
-        start->refereshEventStatus();
-        stop->refereshEventStatus();
     
-        if ((start->_state == hipEventStatusComplete) && (stop->_state == hipEventStatusComplete)) {
+        auto startEcd = start->locked_copyCrit();
-        // Common case, we have good information for both events.
+        auto stopEcd  = stop->locked_copyCrit();
 
-        int64_t tickDiff = (stop->timestamp() - start->timestamp());
+        if ((start->_flags & hipEventDisableTiming) ||
+            (startEcd._state == hipEventStatusUnitialized) || (startEcd._state == hipEventStatusCreated) ||
+            (stop->_flags & hipEventDisableTiming) ||
+            (stopEcd._state == hipEventStatusUnitialized) || (stopEcd._state == hipEventStatusCreated)) {
+
+            // Both events must be at least recorded else return hipErrorInvalidResourceHandle
+
+            status = hipErrorInvalidResourceHandle;
 
-        uint64_t freqHz;
-        hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz);
-        if (freqHz) {
-            *ms = ((double)(tickDiff) /  (double)(freqHz)) * 1000.0f;
-            status = hipSuccess;
         } else {
-            * ms = 0.0f;
+            // Refresh status, if still recording...
-            status = hipErrorInvalidValue;
+
-        }
+            auto startStatus = start->refreshEventStatus();  // pair < state, timestamp >
+            auto stopStatus  = stop->refreshEventStatus();   // pair < state, timestamp >
+
+            if ((startStatus.first == hipEventStatusComplete) && (stopStatus.first == hipEventStatusComplete)) {
+                // Common case, we have good information for both events.  'second" is the timestamp:
+                int64_t tickDiff = (stopStatus.second - startStatus.second);
+
+                uint64_t freqHz;
+                hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz);
+                if (freqHz) {
+                    *ms = ((double)(tickDiff) /  (double)(freqHz)) * 1000.0f;
+                    status = hipSuccess;
+                } else {
+                * ms = 0.0f;
+                status = hipErrorInvalidValue;
+            }
 
 
-        } else if ((start->_state == hipEventStatusRecording) ||
+            } else if ((startStatus.first == hipEventStatusRecording) ||
-                   (stop->_state  == hipEventStatusRecording)) {
+                       (stopStatus.first  == hipEventStatusRecording)) {
 
-            status = hipErrorNotReady;
+                status = hipErrorNotReady;
-        } else {
+            } else {
                 assert(0);
+            }
         }
     } 
 
@@ -275,7 +280,9 @@ hipError_t hipEventQuery(hipEvent_t event)
         tprintf(DB_WARN, "hipEventQuery on event without system-scope fence ; consider creating with hipEventReleaseToSystem\n");
     }
 
-    if ((event->_state == hipEventStatusRecording) && !event->locked_isReady()) {
+    auto ecd = event->locked_copyCrit();
+
+    if ((ecd._state == hipEventStatusRecording) && !ecd._stream->locked_eventIsReady(event)) {
         return ihipLogStatus(hipErrorNotReady);
     } else {
         return ihipLogStatus(hipSuccess);

projects/clr/hipamd/src/hip_hcc.cpp

@@ -339,12 +339,11 @@ 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_streamWaitEvent(hipEvent_t event)
+void ihipStream_t::locked_streamWaitEvent(ihipEventData_t &ecd)
 {
     LockedAccessor_StreamCrit_t crit(_criticalData);
 
-
+    crit->_av.create_blocking_marker(ecd.marker(), hc::accelerator_scope);
-    crit->_av.create_blocking_marker(event->marker(), hc::accelerator_scope);
 }
 
 
@@ -352,24 +351,28 @@ void ihipStream_t::locked_streamWaitEvent(hipEvent_t event)
 // 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);
+    LockedAccessor_StreamCrit_t scrit(_criticalData);
 
-    return (event->marker().is_ready());
+    LockedAccessor_EventCrit_t ecrit(event->criticalData());
+
+    return (ecrit->_eventData.marker().is_ready());
 }
 
-void ihipStream_t::locked_eventWaitComplete(hipEvent_t event, hc::hcWaitMode waitMode)
+// Waiting on event can cause HCC to reclaim stream resources - so need to lock the stream.
+void ihipStream_t::locked_eventWaitComplete(hc::completion_future &marker, hc::hcWaitMode waitMode)
 {
     LockedAccessor_StreamCrit_t crit(_criticalData);
 
-    event->marker().wait(waitMode);
+    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)
+hc::completion_future ihipStream_t::locked_recordEvent(hipEvent_t event)
 {
     // Lock the stream to prevent simultaneous access
     LockedAccessor_StreamCrit_t crit(_criticalData);
@@ -385,7 +388,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));
+    return crit->_av.create_marker(scopeFlag);
 };
 
 //=============================================================================

projects/clr/hipamd/src/hip_hcc_internal.h

@@ -137,6 +137,7 @@ extern std::vector<ProfTrigger> g_dbStopTriggers;
 class ihipStream_t;
 class ihipDevice_t;
 class ihipCtx_t;
+struct ihipEventData_t;
 
 // Color defs for debug messages:
 #define KNRM  "\x1B[0m"
@@ -152,10 +153,12 @@ extern const char *API_COLOR;
 extern const char *API_COLOR_END;
 
 
-// If set, thread-safety is enforced on all stream functions.
+// If set, thread-safety is enforced on all event/stream/ctx/device functions.
-// Stream functions will acquire a mutex before entering critical sections.
+// Can disable for performance or functional experiments - in this case
-#define STREAM_THREAD_SAFE 1
+// the code uses a dummy "no-op" mutex.
+#define EVENT_THREAD_SAFE 1
 
+#define STREAM_THREAD_SAFE 1
 
 #define CTX_THREAD_SAFE 1
 
@@ -390,6 +393,12 @@ class FakeMutex
     void unlock() { }
 };
 
+#if EVENT_THREAD_SAFE
+typedef std::mutex EventMutex;
+#else
+#warning "Stream thread-safe disabled"
+typedef FakeMutex EventMutex;
+#endif
 
 #if STREAM_THREAD_SAFE
 typedef std::mutex StreamMutex;
@@ -540,11 +549,11 @@ public:
 
     hc::accelerator_view* locked_getAv() { LockedAccessor_StreamCrit_t crit(_criticalData); return &(crit->_av); };
 
-    void                 locked_streamWaitEvent(hipEvent_t event);
+    void                 locked_streamWaitEvent(ihipEventData_t & event);
-    void                 locked_recordEvent(hipEvent_t event);
+    hc::completion_future    locked_recordEvent(hipEvent_t event);
 
     bool                 locked_eventIsReady(hipEvent_t event);
-    void                 locked_eventWaitComplete(hipEvent_t event, hc::hcWaitMode waitMode);
+    void                 locked_eventWaitComplete(hc::completion_future &marker, hc::hcWaitMode waitMode);
 
     ihipStreamCritical_t  &criticalData() { return _criticalData; };
 
@@ -628,32 +637,76 @@ enum ihipEventType_t {
     hipEventTypeStopCommand,
 };
 
+
+struct ihipEventData_t
+{
+    ihipEventData_t() {
+        _state  = hipEventStatusCreated;
+        _stream = NULL;
+        _timestamp  = 0;
+        _type   = hipEventTypeIndependent;
+    };
+
+    void marker(const hc::completion_future & marker)  { _marker = marker; };
+    hc::completion_future & marker() { return _marker; }
+    uint64_t timestamp() const { return _timestamp; } ;
+    ihipEventType_t type() const { return _type; };
+
+    ihipEventType_t       _type;
+    hipEventStatus_t      _state;
+    hipStream_t           _stream;  // Stream where the event is recorded.  Null stream is resolved to actual stream when recorded
+    uint64_t              _timestamp;  // store timestamp, may be set on host or by marker.
+private:
+    hc::completion_future _marker;
+};
+
+
+//=============================================================================
+//class ihipEventCriticalBase_t
+template <typename MUTEX_TYPE>
+class ihipEventCriticalBase_t : LockedBase<MUTEX_TYPE>
+{
+public:
+    ihipEventCriticalBase_t(const ihipEvent_t *parentEvent) : 
+        _parent(parentEvent)
+    {}
+    ~ihipEventCriticalBase_t() {};
+   
+     // Keep data in structure so it can be easily copied into snapshots
+     // (used to reduce lock contention and preserve correct lock order)
+    ihipEventData_t _eventData;
+
+private:
+    const ihipEvent_t                 *_parent;
+    friend class LockedAccessor<ihipEventCriticalBase_t>;
+};
+
+typedef ihipEventCriticalBase_t<EventMutex> ihipEventCritical_t;
+
+typedef LockedAccessor<ihipEventCritical_t> LockedAccessor_EventCrit_t;
+
 // internal hip event structure.
 class ihipEvent_t {
 public:
     ihipEvent_t(unsigned flags);
     void attachToCompletionFuture(const hc::completion_future *cf, hipStream_t stream, ihipEventType_t eventType);
-    void refereshEventStatus();
+    std::pair<hipEventStatus_t, uint64_t> refreshEventStatus(); // returns pair <state, timestamp>
-    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; } ;
+    // Return a copy of the critical state. The critical data is locked during the copy.
-    ihipEventType_t type() const { return _type; };
+	ihipEventData_t locked_copyCrit() {
+        LockedAccessor_EventCrit_t crit(_criticalData);
+        return _criticalData._eventData; 
+    };
+
+	ihipEventCritical_t &criticalData() { return _criticalData; };
 
 public:
-    hipEventStatus_t      _state;
-
-    hipStream_t           _stream;  // Stream where the event is recorded.  Null stream is resolved to actual stream when recorded
     unsigned              _flags;
 
-
 private:
-    hc::completion_future _marker;
+	ihipEventCritical_t       _criticalData;
-    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);
 } ;
 
@@ -671,7 +724,6 @@ public:
     };
 
     ~ihipDeviceCriticalBase_t()  {
-
     }
 
     // Contexts:

projects/clr/hipamd/src/hip_stream.cpp

@@ -93,18 +93,23 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
 
     hipError_t e = hipSuccess;
 
+    auto ecd = event->locked_copyCrit();
+
     if (event == nullptr) {
         e = hipErrorInvalidResourceHandle;
 
-    } else if (event->_state != hipEventStatusUnitialized) {
+    } else if ((ecd._state != hipEventStatusUnitialized) && 
+               (ecd._state != hipEventStatusCreated)) {
 
         if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) {
             // conservative wait on host for the specified event to complete:
-            event->locked_waitComplete((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
+            // return _stream->locked_eventWaitComplete(this, waitMode);
+            //
+            ecd._stream->locked_eventWaitComplete(ecd.marker(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
         } else {
             stream = ihipSyncAndResolveStream(stream);
             // This will use create_blocking_marker to wait on the specified queue.
-            stream->locked_streamWaitEvent(event);
+            stream->locked_streamWaitEvent(ecd);
         }
 
     } // else event not recorded, return immediately and don't create marker.

projects/clr/hipamd/tests/README.md

@@ -75,6 +75,18 @@ RUN: %t CMAKE_TEST_NAME <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc
 ```
 
 
+#### CMAKECMD command
+
+The supported syntax for the CMAKECMD command is:
+```
+CMAKECMD: <cmake_command> <options_to_cmake_command>
+```
+cmake_command: refers to any of the commands supported by ```cmake -E``` as specified in the [cmake documentation](https://cmake.org/cmake/help/latest/manual/cmake.1.html#command-line-tool-mode). Note that the commands are limited by the version of cmake the user is running.
+options_to_cmake_command: refers to the arguments supported by the specific cmake_command. The arguments are parsed by HIT to replace special markers. The markers supported by HIT are:
+%S: Refers to the source directory containing the current source file.
+%B: Refers to the build directory for the current cmake project i.e. CMAKE_CURRENT_BINARY_DIR.
+
+
 ### Running tests:
 ```
 ctest

projects/clr/hipamd/tests/hit/HIT.cmake

@@ -155,6 +155,20 @@ macro(HIT_ADD_FILES _dir _label _parent)
             endif()
         endforeach()
 
+        # Run cmake commands
+        execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --cmakeCMDs ${file}
+            OUTPUT_VARIABLE _contents
+            ERROR_QUIET
+            WORKING_DIRECTORY ${_dir}
+            OUTPUT_STRIP_TRAILING_WHITESPACE)
+        string(REGEX REPLACE "\n" ";" _contents "${_contents}")
+        string(REGEX REPLACE "%S" ${_dir} _contents "${_contents}")
+        string(REGEX REPLACE "%B" ${CMAKE_CURRENT_BINARY_DIR} _contents "${_contents}")
+        foreach(_cmd ${_contents})
+            string(REGEX REPLACE " " ";" _cmd "${_cmd}")
+            execute_process(COMMAND ${CMAKE_COMMAND} -E ${_cmd})
+        endforeach()
+
         # Add tests
         execute_process(COMMAND ${HIP_SRC_PATH}/tests/hit/parser --runCMDs ${file}
             OUTPUT_VARIABLE _contents

projects/clr/hipamd/tests/hit/parser

@@ -8,7 +8,7 @@ use File::Spec;
 sub parse_file {
     my $file = shift;
     (my $exe = $file) =~ s/\.[^.]+$//g;
-    my (@buildCMDs, @runCMDs, @runNamedCMDs);
+    my (@buildCMDs, @runCMDs, @runNamedCMDs, @cmakeCMDs);
     if (open (SOURCE, '<:encoding(UTF-8)', "$file")) {
         while (<SOURCE>) {
             my $line=$_;
@@ -36,10 +36,17 @@ sub parse_file {
                 $line =~ s/\R//g;               # Remove line endings
                 push @runNamedCMDs, $line;
             }
+            # Look for CMAKECMD instructions
+            if ($line =~ /^ \* CMAKECMD:/) {
+                $line =~ s/^ \* CMAKECMD: //g;  # Remove " * CMAKECMD: "
+                # Substitute %S -> srcdir and %B -> builddir happens in cmake
+                $line =~ s/\R//g;               # Remove line endings
+                push @cmakeCMDs, $line;
+            }
         }
         close(SOURCE);
     }
-    return (\@buildCMDs, \@runCMDs, \@runNamedCMDs);
+    return (\@buildCMDs, \@runCMDs, \@runNamedCMDs, \@cmakeCMDs);
 }
 
 # Exit if no arguments specified
@@ -53,8 +60,9 @@ my @options = ();
 my $retBuildCMDs = 0;
 my $retRunCMDs = 0;
 my $retRunNamedCMDs = 0;
+my $retCmakeCMDs = 0;
 foreach $arg (@ARGV) {
-    if ($retBuildCMDs or $retRunCMDs or $retRunNamedCMDs) {
+    if ($retBuildCMDs or $retRunCMDs or $retRunNamedCMDs or $retCmakeCMDs) {
         push (@options, $arg);
     }
     if ($arg eq '--buildCMDs') {
@@ -66,18 +74,21 @@ foreach $arg (@ARGV) {
     if ($arg eq '--runNamedCMDs') {
         $retRunNamedCMDs = 1;
     }
+    if ($arg eq '--cmakeCMDs') {
+        $retCmakeCMDs = 1;
+    }
 }
 
 # Atleast one command needs to be specified
-if (($retBuildCMDs eq 0) and ($retRunCMDs eq 0) and ($retRunNamedCMDs eq 0)) {
+if (($retBuildCMDs eq 0) and ($retRunCMDs eq 0) and ($retRunNamedCMDs eq 0) and ($retCmakeCMDs eq 0)) {
-    die "Usage: $0 <--buildCMDs|--runCMDs|--runNamedCMDs> FILENAMEs\n";
+    die "Usage: $0 <--buildCMDs|--runCMDs|--runNamedCMDs|--cmakeCMDs> FILENAMEs\n";
 }
 
 # Iterate over input files
 foreach $file (@options) {
     # Convert absolute path to path relative to working directory
     my $relfile = File::Spec->abs2rel($file);
-    my ($buildCMDs, $runCMDs, $runNamedCMDs) = parse_file("$relfile");
+    my ($buildCMDs, $runCMDs, $runNamedCMDs, $cmakeCMDs) = parse_file("$relfile");
     if ($retBuildCMDs) {
         # print "BuildCMDs:\n";
         print "$_\n" for @$buildCMDs;
@@ -90,6 +101,10 @@ foreach $file (@options) {
         # print "RunNamedCMDs:\n";
         print "$_\n" for @$runNamedCMDs;
     }
+    if ($retCmakeCMDs) {
+        # print "CmakeCMDs:\n";
+        print "$_\n" for @$cmakeCMDs;
+    }
 }
 
 # vim: ts=4:sw=4:expandtab:smartindent

projects/clr/hipamd/tests/src/runtimeApi/error/hipPeekAtLastError.cpp

@@ -0,0 +1,38 @@
+/*
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipGetDevice(int *device);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+        hipSetDevice(-1);
+        if(hipPeekAtLastError() != hipSuccess)
+          passed();
+}

projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp

@@ -0,0 +1,80 @@
+/*
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t sizeBytes);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    hipDevice_t device;
+    size_t Nbytes = N*sizeof(int);
+    int numDevices = 0;
+    int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d;
+    int *A_h, *B_h, *C_h ;
+
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+     if(numDevices > 1)
+     {
+        HIPCHECK(hipSetDevice(0));
+        unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+        HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+        HIPCHECK(hipSetDevice(1));
+        HIPCHECK(hipMalloc(&X_d,Nbytes));
+        HIPCHECK(hipMalloc(&Y_d,Nbytes));
+        HIPCHECK(hipMalloc(&Z_d,Nbytes));
+
+        
+        HIPCHECK(hipSetDevice(0));
+        HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+        HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+        hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N);
+        HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+        HIPCHECK(hipDeviceSynchronize());
+        HipTest::checkVectorADD(A_h, B_h, C_h, N);
+    
+         
+        HIPCHECK(hipSetDevice(1));
+        HIPCHECK(hipMemcpyDtoD(X_d, A_d,  Nbytes)); 
+        HIPCHECK(hipMemcpyDtoD(Y_d, B_d,  Nbytes));
+
+        hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N);
+        HIPCHECK(hipMemcpyDtoH(C_h, Z_d, Nbytes));
+        HIPCHECK(hipDeviceSynchronize());
+        HipTest::checkVectorADD(A_h, B_h, C_h, N);
+
+        HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+        HIPCHECK(hipFree(X_d));
+        HIPCHECK(hipFree(Y_d));
+        HIPCHECK(hipFree(Z_d));
+      }
+     
+        passed();
+     
+}
+

projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp

@@ -0,0 +1,84 @@
+/*
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t sizeBytes);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    hipDevice_t device;
+    size_t Nbytes = N*sizeof(int);
+    int numDevices = 0;
+    int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d;
+    int *A_h, *B_h, *C_h ;
+    hipStream_t s;
+
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+     if(numDevices > 1)
+     {
+        HIPCHECK(hipSetDevice(0));
+        unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+        HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+        HIPCHECK(hipSetDevice(1));
+        HIPCHECK(hipMalloc(&X_d,Nbytes));
+        HIPCHECK(hipMalloc(&Y_d,Nbytes));
+        HIPCHECK(hipMalloc(&Z_d,Nbytes));
+
+        
+        HIPCHECK(hipSetDevice(0));
+        HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+        HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+        hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N);
+        HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+        HIPCHECK(hipDeviceSynchronize());
+        HipTest::checkVectorADD(A_h, B_h, C_h, N);
+    
+        HIPCHECK(hipStreamCreate(&s));    
+        HIPCHECK(hipSetDevice(1));
+        HIPCHECK(hipMemcpyDtoDAsync(X_d, A_d, Nbytes, s)); 
+        HIPCHECK(hipMemcpyDtoDAsync(Y_d, B_d, Nbytes, s));
+
+        hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N);
+        HIPCHECK(hipMemcpyDtoHAsync(C_h, Z_d, Nbytes, s));
+        HIPCHECK(hipStreamSynchronize(s));
+        HIPCHECK(hipDeviceSynchronize());
+        
+        HipTest::checkVectorADD(A_h, B_h, C_h, N);
+        HIPCHECK(hipStreamDestroy(s));
+        HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+        HIPCHECK(hipFree(X_d));
+        HIPCHECK(hipFree(Y_d));
+        HIPCHECK(hipFree(Z_d));
+     }
+
+        passed();
+     
+
+}
+

projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeer.cpp

@@ -0,0 +1,80 @@
+/*
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t sizeBytes);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    hipDevice_t device;
+    size_t Nbytes = N*sizeof(int);
+    int numDevices = 0;
+    int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d;
+    int *A_h, *B_h, *C_h ;
+
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+     if(numDevices > 1)
+     {
+        HIPCHECK(hipSetDevice(0));
+        unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+        HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+        HIPCHECK(hipSetDevice(1));
+        HIPCHECK(hipMalloc(&X_d,Nbytes));
+        HIPCHECK(hipMalloc(&Y_d,Nbytes));
+        HIPCHECK(hipMalloc(&Z_d,Nbytes));
+        
+        HIPCHECK(hipSetDevice(0));
+        HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+        HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+        hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N);
+        HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+        HIPCHECK(hipDeviceSynchronize());
+        HipTest::checkVectorADD(A_h, B_h, C_h, N);
+             
+        HIPCHECK(hipSetDevice(1));
+        hipMemcpyPeer(X_d, 1, A_d, 0, Nbytes); //this call is eqv to hipMemcpy(hipMemcpyD2D) which goes via stg bufs.
+        hipMemcpyPeer(Y_d, 1, B_d, 0, Nbytes);
+
+        hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N);
+        HIPCHECK(hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost));
+        HIPCHECK(hipDeviceSynchronize());
+        HipTest::checkVectorADD(A_h, B_h, C_h, N);
+        
+        HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+        HIPCHECK(hipFree(X_d));
+        HIPCHECK(hipFree(Y_d));
+        HIPCHECK(hipFree(Z_d));
+     }
+        passed();
+
+
+     
+
+}
+

projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemcpyPeerAsync.cpp

@@ -0,0 +1,85 @@
+/*
+Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t sizeBytes);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    hipDevice_t device;
+    size_t Nbytes = N*sizeof(int);
+    int numDevices = 0;
+    int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d;
+    int *A_h, *B_h, *C_h ;
+    hipStream_t s;
+
+
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+     if(numDevices > 1)
+     {
+        HIPCHECK(hipSetDevice(0));
+        unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+        HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+        HIPCHECK(hipSetDevice(1));
+        HIPCHECK(hipMalloc(&X_d,Nbytes));
+        HIPCHECK(hipMalloc(&Y_d,Nbytes));
+        HIPCHECK(hipMalloc(&Z_d,Nbytes));
+
+        
+        HIPCHECK(hipSetDevice(0));
+        HIPCHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+        HIPCHECK ( hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+        hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d,B_d, C_d, N);
+        HIPCHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+        HIPCHECK (hipDeviceSynchronize());
+        HipTest::checkVectorADD(A_h, B_h, C_h, N);
+    
+        HIPCHECK(hipStreamCreate(&s));  
+        HIPCHECK(hipSetDevice(1));
+        HIPCHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, Nbytes, s));
+        HIPCHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, Nbytes, s));
+
+        hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, X_d,Y_d, Z_d, N);
+        HIPCHECK ( hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost));
+        HIPCHECK (hipDeviceSynchronize());
+        HIPCHECK (hipStreamSynchronize(s));
+        HipTest::checkVectorADD(A_h, B_h, C_h, N);
+        
+        HIPCHECK(hipStreamDestroy(s));
+        HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+        HIPCHECK(hipFree(X_d));
+        HIPCHECK(hipFree(Y_d));
+        HIPCHECK(hipFree(Z_d));
+     }
+
+        passed();
+    
+
+}
+

projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamGetFlags.cpp

@@ -0,0 +1,44 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+
+int main(int argc, char *argv[])
+{
+    hipStream_t stream;
+    unsigned int flags;
+    HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
+    HIPCHECK(hipStreamGetFlags(stream, &flags));
+    HIPASSERT(flags == 0); 
+    HIPCHECK(hipStreamDestroy(stream));
+
+    HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
+    HIPCHECK(hipStreamGetFlags(stream, &flags));
+    HIPASSERT(flags == 1); 
+    HIPCHECK(hipStreamDestroy(stream));
+
+    passed();
+}