SWDEV-299127 - Merge 'develop' into 'amd-staging'

Change-Id: I3ddbbc4724b90a55c57b85d2297ea1cfca711772

bin/hipcc.pl

@@ -52,6 +52,19 @@ if(scalar @ARGV == 0){
     exit(-1);
 }
 
+# retrieve --rocm-path hipcc option from command line.
+# We need to respect this over the env var ROCM_PATH for this compilation.
+sub get_rocm_path_option {
+  my $rocm_path="";
+  my @CLArgs = @ARGV;
+  foreach $arg (@CLArgs) {
+    if (index($arg,"--rocm-path=") != -1) {
+      ($rocm_path) = $arg=~ /=\s*(.*)\s*$/;
+    }
+  }
+  return $rocm_path;
+}
+
 $verbose = $ENV{'HIPCC_VERBOSE'} // 0;
 # Verbose: 0x1=commands, 0x2=paths, 0x4=hipcc args
 
@@ -88,12 +101,18 @@ sub delete_temp_dirs {
 }
 
 my $base_dir;
+my $rocmPath;
 BEGIN {
     $base_dir = dirname(Cwd::realpath(__FILE__) );
+    $rocmPath = get_rocm_path_option();
+    if ($rocmPath ne '') {
+      # --rocm-path takes precedence over ENV{ROCM_PATH}
+      $ENV{ROCM_PATH}=$rocmPath;
+    }
 }
 use lib "$base_dir/";
-use hipvars;
 
+use hipvars;
 $isWindows      =   $hipvars::isWindows;
 $HIP_RUNTIME    =   $hipvars::HIP_RUNTIME;
 $HIP_PLATFORM   =   $hipvars::HIP_PLATFORM;

tests/catch/unit/stream/hipStreamSynchronize.cc

@@ -33,6 +33,19 @@ TEST_CASE("Unit_hipStreamSynchronize_EmptyStream") {
   HIP_CHECK(hipStreamDestroy(stream));
 }
 
+#if !HT_NVIDIA
+/**
+ * @brief Check that synchronization of uninitialized stream sets its status to
+ * hipErrorContextIsDestroyed
+ *
+ * Test removed for Nvidia devices because it returns unexpected error
+ */
+TEST_CASE("Unit_hipStreamSynchronize_UninitializedStream") {
+  hipStream_t stream{reinterpret_cast<hipStream_t>(0xFFFF)};
+  HIP_CHECK_ERROR(hipStreamSynchronize(stream), hipErrorContextIsDestroyed);
+}
+#endif
+
 #if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
 
 /**

tests/catch/unit/stream/hipStreamValue.cc

@@ -499,7 +499,40 @@ TEST_CASE("Unit_hipStreamValue_Negative_InvalidMemory") {
   HIP_CHECK(hipStreamDestroy(stream));
 }
 
-TEMPLATE_TEST_CASE("Unit_hipStreamWaitValue_Negative_InvalidFlag", "", uint32_t, uint64_t) {
+TEMPLATE_TEST_CASE("Unit_hipStreamValue_Negative_UninitializedStream", "", uint32_t, uint64_t) {
+  if (!streamWaitValueSupported()) {
+    HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device.");
+    return;
+  }
+
+  hipStream_t stream{reinterpret_cast<hipStream_t>(0xFFFF)};
+
+  // Allocate Host Memory
+  auto hostPtr = std::make_unique<TestType>();
+
+  // Register Host Memory
+  HIP_CHECK(hipHostRegister(hostPtr.get(), sizeof(TestType), 0));
+
+  // Set dummy data
+  *hostPtr = 0x0;
+
+  const auto compareOp = hipStreamWaitValueGte;
+  const auto expectedError = hipErrorContextIsDestroyed;
+
+  // Stream handle negative tests
+  SECTION("Invalid Stream handle for hipStreamWriteValue") {
+    HIP_CHECK_ERROR(writeFunc<TestType>(stream, hostPtr.get(), 0, writeFlag), expectedError);
+  }
+
+  SECTION("Invalid Stream handle for hipStreamWaitValue") {
+    HIP_CHECK_ERROR(waitFunc<TestType>(stream, hostPtr.get(), 0, compareOp), expectedError);
+  }
+
+  // Cleanup
+  HIP_CHECK(hipHostUnregister(hostPtr.get()));
+}
+
+TEMPLATE_TEST_CASE("Unit_hipStreamValue_Negative_InvalidFlag", "", uint32_t, uint64_t) {
   if (!streamWaitValueSupported()) {
     HipTest::HIP_SKIP_TEST("hipStreamWaitValue not supported on this device.");
     return;
@@ -518,7 +551,6 @@ TEMPLATE_TEST_CASE("Unit_hipStreamWaitValue_Negative_InvalidFlag", "", uint32_t,
   // Set dummy data
   *hostPtr = 0x0;
 
-  /* EXSWCPHIPT-96 */
   HIP_CHECK_ERROR(waitFunc<TestType>(stream, hostPtr.get(), 0, -1), hipErrorInvalidValue);
 
   // Cleanup