diff --git a/bin/hipcc.pl b/bin/hipcc.pl index 76163d97b6..da9559be9e 100755 --- a/bin/hipcc.pl +++ b/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; diff --git a/tests/catch/unit/stream/hipStreamSynchronize.cc b/tests/catch/unit/stream/hipStreamSynchronize.cc index 0555620b1f..a31546a52a 100644 --- a/tests/catch/unit/stream/hipStreamSynchronize.cc +++ b/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(0xFFFF)}; + HIP_CHECK_ERROR(hipStreamSynchronize(stream), hipErrorContextIsDestroyed); +} +#endif + #if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */ /** diff --git a/tests/catch/unit/stream/hipStreamValue.cc b/tests/catch/unit/stream/hipStreamValue.cc index 164384acee..7a55360eff 100644 --- a/tests/catch/unit/stream/hipStreamValue.cc +++ b/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(0xFFFF)}; + + // Allocate Host Memory + auto hostPtr = std::make_unique(); + + // 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(stream, hostPtr.get(), 0, writeFlag), expectedError); + } + + SECTION("Invalid Stream handle for hipStreamWaitValue") { + HIP_CHECK_ERROR(waitFunc(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(stream, hostPtr.get(), 0, -1), hipErrorInvalidValue); // Cleanup