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

Change-Id: Iaba16b6cc2929fecc6baa44c48552df5c7d63555

tests/catch/include/hip_test_process.hh

@@ -66,7 +66,12 @@ class SpawnProc {
     auto dir = fs::path(TestContext::get().currentPath());
     dir /= exeName;
     exeName = dir.string();
-
+    // On Windows, fs::exists returns false without extension.
+    if (TestContext::get().isWindows()) {
+      if(fs::path(exeName).extension().empty()) {
+        exeName += ".exe";
+      }
+    }
     INFO("Testing that exe exists: " << exeName);
     REQUIRE(fs::exists(exeName));
 
@@ -92,13 +97,13 @@ class SpawnProc {
       execCmd += " > ";
       execCmd += tmpFileName;
     }
-
     auto res = std::system(execCmd.c_str());
 
     if (captureOutput) {
       std::ifstream t(tmpFileName.c_str());
       resultStr =
           std::string((std::istreambuf_iterator<char>(t)), std::istreambuf_iterator<char>());
+      t.close();
     }
 #if HT_LINUX
     return WEXITSTATUS(res);

tests/catch/multiproc/hipMallocConcurrencyMproc.cc

@@ -103,6 +103,14 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
   HIP_CHECK(hipSetDevice(gpu));
   HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
   HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+  HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
+
+  if (!concurOnOneGPU && (prevAvl < curAvl || prevTot != curTot)) {
+    //In concurrent calls on one GPU, we cannot verify leaking in this way
+    printf("%s : Memory allocation mismatch observed."
+        "Possible memory leak.\n", __func__);
+    TestPassed &= false;
+  }
 
   unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
 
@@ -122,9 +130,17 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
     TestPassed = false;
   }
 
+  HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
   HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
   HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
 
+  if (!concurOnOneGPU && (curAvl < prevAvl || prevTot != curTot)) {
+    // In concurrent calls on one GPU, we cannot verify leaking in this way
+    UNSCOPED_INFO("validateMemoryOnGPU : Memory allocation mismatch observed."
+                  << "Possible memory leak.");
+    TestPassed = false;
+  }
+
   if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
     // In concurrent calls on one GPU, we cannot verify leaking in this way
     printf(
@@ -200,7 +216,7 @@ TEST_CASE("Unit_hipMalloc_ChildConcurrencyMultiGpu") {
 
     } else if (!pid) {  // Child process
       bool TestPassedChild = false;
-      TestPassedChild = validateMemoryOnGPU(gpu);
+      TestPassedChild = validateMemoryOnGPU(gpu, true);
 
       if (TestPassedChild) {
         exit(resSuccess);  // child exit with success status

tests/catch/unit/device/hipGetDeviceAttribute.cc

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2022 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
@@ -22,10 +22,12 @@ THE SOFTWARE.
 // Test the device info API extensions for HIP
 
 #include <hip_test_common.hh>
+#include <string.h>
+#ifdef __linux__
+#include <unistd.h>
+#endif
 #include <iostream>
 
-
-
 static hipError_t test_hipDeviceGetAttribute(int deviceId,
                                       hipDeviceAttribute_t attr,
                                       int expectedValue = -1) {
@@ -234,15 +236,78 @@ TEST_CASE("Unit_hipGetDeviceAttribute_CheckAttrValues") {
 /**
  * Validate the hipDeviceAttributeFineGrainSupport property in AMD.
  */
+#ifdef __linux__
 #if HT_AMD
+#define COMMAND_LEN 256
+#define BUFFER_LEN 512
+
+static bool isRocmPathSet() {
+  FILE *fpipe;
+  char const *command = "echo $ROCM_PATH";
+  fpipe = popen(command, "r");
+
+  if (fpipe == nullptr) {
+    printf("Unable to create command\n");
+    return false;
+  }
+  char command_op[BUFFER_LEN];
+  if (fgets(command_op, BUFFER_LEN, fpipe)) {
+    size_t len = strlen(command_op);
+    if (len > 1) {  // This is because fgets always adds newline character
+      pclose(fpipe);
+      return true;
+    }
+  }
+  pclose(fpipe);
+  return false;
+}
+
 // This is AMD specific property test
 TEST_CASE("Unit_hipGetDeviceAttribute_CheckFineGrainSupport") {
   int deviceId;
   int deviceCount = 0;
+  FILE *fpipe;
+  char command[COMMAND_LEN] = "";
+  const char *rocmpath = nullptr;
+  if (isRocmPathSet()) {
+    // For STG2 testing where /opt/rocm path is not present
+    rocmpath = "$ROCM_PATH/bin/rocminfo";
+  } else {
+    // Check if the rocminfo tool exists
+    rocmpath = "/opt/rocm/bin/rocminfo";
+  }
+  snprintf(command, COMMAND_LEN, "%s", rocmpath);
+  strncat(command, " | grep -i \"Segment:\\|Uuid:\"", COMMAND_LEN);
+  // Execute the rocminfo command and extract the segment info
+  fpipe = popen(command, "r");
+  if (fpipe == nullptr) {
+    printf("Unable to create command file\n");
+    return;
+  }
   HIP_CHECK(hipGetDeviceCount(&deviceCount));
-  REQUIRE(deviceCount > 0);
-  // Check hipDeviceAttributeFineGrainSupport for all available device
-  // in system.
+  assert(deviceCount > 0);
+  int *fine_grained_val = new int[deviceCount];
+  assert(fine_grained_val != nullptr);
+  bool *gpuFound = new bool[deviceCount];
+  assert(gpuFound != nullptr);
+  for (int i = 0; i < deviceCount; i++) {
+    gpuFound[i] = false;
+    fine_grained_val[i] = 0;  // Initialize to 0
+  }
+  char command_op[BUFFER_LEN];
+  int count = -1;
+  // Extract each segment flags
+  while (fgets(command_op, BUFFER_LEN, fpipe)) {
+    std::string rocminfo_line(command_op);
+    if ((std::string::npos != rocminfo_line.find("GPU-")) &&
+        (std::string::npos != rocminfo_line.find("Uuid:"))) {
+      count++;
+      gpuFound[count] = true;
+    } else if (gpuFound[count] &&
+    (std::string::npos != rocminfo_line.find("FLAGS: FINE GRAINED"))) {
+      fine_grained_val[count] = 1;
+    }
+  }
   for (int dev = 0; dev < deviceCount; dev++) {
     HIP_CHECK(hipSetDevice(dev));
     HIP_CHECK(hipGetDevice(&deviceId));
@@ -250,21 +315,15 @@ TEST_CASE("Unit_hipGetDeviceAttribute_CheckFineGrainSupport") {
     HIP_CHECK(hipGetDeviceProperties(&props, deviceId));
     int value = 0;
     HIP_CHECK(hipDeviceGetAttribute(&value,
-        hipDeviceAttributeFineGrainSupport, deviceId));
-    std::string gpu_arch_name(props.gcnArchName);
-    if (std::string::npos != gpu_arch_name.find("gfx90a")) {
-      // Current GPU is gfx90a architecture
-      REQUIRE(value == 1);
-    } else if (std::string::npos != gpu_arch_name.find("gfx906")) {
-      // Current GPU is gfx906 architecture
-      REQUIRE(value == 0);
-    } else if (std::string::npos != gpu_arch_name.find("gfx908")) {
-      // Current GPU is gfx908 architecture
-      REQUIRE(value == 0);
-    }
+              hipDeviceAttributeFineGrainSupport, deviceId));
+    REQUIRE(value == fine_grained_val[dev]);
   }
+  // Validate hipDeviceAttributeFineGrainSupport
+  delete[] fine_grained_val;
+  delete[] gpuFound;
 }
 #endif
+#endif
 /**
  * Validates negative scenarios for hipDeviceGetAttribute
  * scenario1: pi = nullptr

tests/catch/unit/event/CMakeLists.txt

@@ -5,6 +5,7 @@ set(TEST_SRC
     Unit_hipEventElapsedTime.cc
     Unit_hipEventRecord.cc
     Unit_hipEventIpc.cc
+    hipEventDestroy.cc
 )
 
 # The test used wait mechanism and doesnt play well with all arch of nvidia

tests/catch/unit/event/hipEventDestroy.cc

@@ -0,0 +1,115 @@
+/*
+Copyright (c) 2022 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.
+*/
+
+#include <chrono>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+#include <hip_test_common.hh>
+#include "hip/hip_runtime_api.h"
+
+#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
+
+static constexpr size_t vectorSize{1024};
+
+/**
+ * @brief Launches vectorAdd kernel with a delay
+ */
+static inline void launchVectorAdd(float*& A_h, float*& B_h, float*& C_h,
+                                   std::chrono::milliseconds delay, hipStream_t stream = nullptr) {
+  float* A_d{nullptr};
+  float* B_d{nullptr};
+  float* C_d{nullptr};
+  HipTest::initArraysForHost(&A_h, &B_h, &C_h, vectorSize, true);
+  HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d), A_h, 0));
+  HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&B_d), B_h, 0));
+  HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&C_d), C_h, 0));
+  HipTest::runKernelForDuration(delay, stream);
+  HipTest::vectorADD<<<1, 1, 0, stream>>>(A_d, B_d, C_d, vectorSize);
+}
+
+
+/**
+ * @brief Check that destroying an event before the kernel has finished running causes no errors.
+ *
+ */
+TEST_CASE("Unit_hipEventDestroy_Unfinished") {
+  hipEvent_t event;
+
+  HIP_CHECK(hipEventCreate(&event));
+
+  float *A_h, *B_h, *C_h;
+  launchVectorAdd(A_h, B_h, C_h, std::chrono::milliseconds(1000));
+
+  HIP_CHECK(hipEventRecord(event));
+  HIP_CHECK_ERROR(hipEventQuery(event), hipErrorNotReady);
+  HIP_CHECK(hipEventDestroy(event));
+
+  HIP_CHECK(hipDeviceSynchronize());
+  HipTest::checkVectorADD(A_h, B_h, C_h, vectorSize);
+  HipTest::freeArraysForHost(A_h, B_h, C_h, true);
+}
+
+/**
+ * @brief Check that destroying an event enqueued to a stream causes no errors.
+ *
+ */
+TEST_CASE("Unit_hipEventDestroy_WithWaitingStream") {
+  hipEvent_t event;
+  HIP_CHECK(hipEventCreate(&event));
+
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  float *A_h, *B_h, *C_h;
+  launchVectorAdd(A_h, B_h, C_h, std::chrono::milliseconds(1000), stream);
+
+  HIP_CHECK(hipEventRecord(event, stream));
+  HIP_CHECK_ERROR(hipEventQuery(event), hipErrorNotReady);
+  HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady);
+  HIP_CHECK(hipEventDestroy(event));
+  HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady);
+  HIP_CHECK(hipStreamSynchronize(stream));
+  HIP_CHECK(hipStreamDestroy(stream));
+
+  HipTest::checkVectorADD(A_h, B_h, C_h, vectorSize);
+  HipTest::freeArraysForHost(A_h, B_h, C_h, true);
+}
+
+TEST_CASE("Unit_hipEventDestroy_Negative") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-103");
+  return;
+#endif
+
+  SECTION("Invalid Event") {
+    hipEvent_t event{nullptr};
+    HIP_CHECK_ERROR(hipEventDestroy(event), hipErrorInvalidResourceHandle);
+  }
+
+  SECTION("Destroy twice") {
+    hipEvent_t event;
+    HIP_CHECK(hipEventCreate(&event));
+    HIP_CHECK(hipEventDestroy(event));
+    HIP_CHECK_ERROR(hipEventDestroy(event), hipErrorContextIsDestroyed);
+  }
+}
+#endif

tests/catch/unit/memory/hipArrayCreate.cc

@@ -52,26 +52,28 @@ static constexpr auto ARRAY_LOOP{100};
 
 static void ArrayCreate_DiffSizes(int gpu) {
   HIP_CHECK_THREAD(hipSetDevice(gpu));
-  std::pair<size_t, size_t> size =
-      GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H));
-  std::array<HIP_ARRAY, ARRAY_LOOP> array;
-  size_t pavail, avail;
-  HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
-  HIP_ARRAY_DESCRIPTOR desc;
-  desc.NumChannels = 1;
-  desc.Width = std::get<0>(size);
-  desc.Height = std::get<1>(size);
-  desc.Format = HIP_AD_FORMAT_FLOAT;
+  //Use of GENERATE in thead function causes random failures with multithread condition.
+  std::vector<std::pair<size_t, size_t>> runs {std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)};
+  for (const auto& size : runs) {
+    std::array<HIP_ARRAY, ARRAY_LOOP> array;
+    size_t pavail, avail;
+    HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
+    HIP_ARRAY_DESCRIPTOR desc;
+    desc.NumChannels = 1;
+    desc.Width = std::get<0>(size);
+    desc.Height = std::get<1>(size);
+    desc.Format = HIP_AD_FORMAT_FLOAT;
 
-  for (int i = 0; i < ARRAY_LOOP; i++) {
-    HIP_CHECK_THREAD(hipArrayCreate(&array[i], &desc));
-  }
-  for (int i = 0; i < ARRAY_LOOP; i++) {
-    HIP_CHECK_THREAD(hipArrayDestroy(array[i]));
-  }
+    for (int i = 0; i < ARRAY_LOOP; i++) {
+      HIP_CHECK_THREAD(hipArrayCreate(&array[i], &desc));
+    }
+    for (int i = 0; i < ARRAY_LOOP; i++) {
+      HIP_CHECK_THREAD(hipArrayDestroy(array[i]));
+    }
 
-  HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
-  REQUIRE_THREAD(pavail == avail);
+    HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
+    REQUIRE_THREAD(pavail == avail);
+  }
 }
 
 /* This testcase verifies hipArrayCreate API for small and big chunks data*/

tests/catch/unit/memory/hipMalloc3DArray.cc

@@ -49,23 +49,26 @@ static constexpr auto ARRAY_LOOP{100};
  */
 static void Malloc3DArray_DiffSizes(int gpu) {
   HIP_CHECK_THREAD(hipSetDevice(gpu));
-  const int size = GENERATE(ARRAY_SIZE, BIG_ARRAY_SIZE);
-  int width{size}, height{size}, depth{size};
-  hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
-  std::array<hipArray_t, ARRAY_LOOP> arr;
-  size_t pavail, avail;
-  HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
+  //Use of GENERATE in thead function causes random failures with multithread condition.
+  std::vector<size_t> runs {ARRAY_SIZE, BIG_ARRAY_SIZE};
+  for (const auto& size : runs) {
+    size_t width{size}, height{size}, depth{size};
+    hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
+    std::array<hipArray_t, ARRAY_LOOP> arr;
+    size_t pavail, avail;
+    HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
 
-  for (int i = 0; i < ARRAY_LOOP; i++) {
-    HIP_CHECK_THREAD(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, height, depth),
+    for (int i = 0; i < ARRAY_LOOP; i++) {
+      HIP_CHECK_THREAD(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, height, depth),
                                       hipArrayDefault));
-  }
-  for (int i = 0; i < ARRAY_LOOP; i++) {
-    HIP_CHECK_THREAD(hipFreeArray(arr[i]));
-  }
+    }
+    for (int i = 0; i < ARRAY_LOOP; i++) {
+      HIP_CHECK_THREAD(hipFreeArray(arr[i]));
+    }
 
-  HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
-  REQUIRE_THREAD(pavail == avail);
+    HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
+    REQUIRE_THREAD(pavail == avail);
+  }
 }
 
 TEST_CASE("Unit_hipMalloc3DArray_DiffSizes") {

tests/catch/unit/memory/hipMallocArray.cc

@@ -51,23 +51,25 @@ static constexpr int ARRAY_LOOP{100};
  */
 static void MallocArray_DiffSizes(int gpu) {
   HIP_CHECK_THREAD(hipSetDevice(gpu));
-  std::pair<size_t, size_t> size =
-      GENERATE(std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H));
-  hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
-  std::array<hipArray_t, ARRAY_LOOP> A_d;
-  size_t pavail, avail;
-  HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
+  //Use of GENERATE in thead function causes random failures with multithread condition.
+  std::vector<std::pair<size_t, size_t>> runs {std::make_pair(NUM_W, NUM_H), std::make_pair(BIGNUM_W, BIGNUM_H)};
+  for (const auto& size : runs) {
+    hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
+    std::array<hipArray_t, ARRAY_LOOP> A_d;
+    size_t pavail, avail;
+    HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr));
 
-  for (int i = 0; i < ARRAY_LOOP; i++) {
-    HIP_CHECK_THREAD(
-        hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault));
-  }
-  for (int i = 0; i < ARRAY_LOOP; i++) {
-    HIP_CHECK_THREAD(hipFreeArray(A_d[i]));
-  }
+    for (int i = 0; i < ARRAY_LOOP; i++) {
+      HIP_CHECK_THREAD(
+           hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault));
+    }
+    for (int i = 0; i < ARRAY_LOOP; i++) {
+      HIP_CHECK_THREAD(hipFreeArray(A_d[i]));
+    }
 
-  HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
-  REQUIRE_THREAD(pavail == avail);
+    HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr));
+    REQUIRE_THREAD(pavail == avail);
+  }
 }
 
 TEST_CASE("Unit_hipMallocArray_DiffSizes") {

tests/catch/unit/memory/hipMallocConcurrency.cc

@@ -203,6 +203,14 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
   HIPCHECK(hipSetDevice(gpu));
   HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot));
   HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+  HIPCHECK(hipMemGetInfo(&curAvl, &curTot));
+
+  if (!concurOnOneGPU && (prevAvl < curAvl || prevTot != curTot)) {
+    //In concurrent calls on one GPU, we cannot verify leaking in this way
+    printf("%s : Memory allocation mismatch observed."
+        "Possible memory leak.\n", __func__);
+    TestPassed &= false;
+  }
 
   unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
 
@@ -221,9 +229,17 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) {
     TestPassed = false;
   }
 
+  HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot));
   HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
   HIPCHECK(hipMemGetInfo(&curAvl, &curTot));
 
+  if (!concurOnOneGPU && (curAvl < prevAvl || prevTot != curTot)) {
+    // In concurrent calls on one GPU, we cannot verify leaking in this way
+    UNSCOPED_INFO("validateMemoryOnGPU : Memory allocation mismatch observed."
+                  << "Possible memory leak.");
+    TestPassed = false;
+  }
+
   if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
     // In concurrent calls on one GPU, we cannot verify leaking in this way
     UNSCOPED_INFO(
@@ -291,7 +307,7 @@ static bool regressAllocInLoopMthread(int gpu) {
  * Thread func to regress alloc and check data consistency
  */
 static void threadFunc(int gpu) {
-  g_thTestPassed = regressAllocInLoopMthread(gpu) && validateMemoryOnGpuMThread(gpu);
+  g_thTestPassed = regressAllocInLoopMthread(gpu) && validateMemoryOnGpuMThread(gpu, true);
 
   UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" << g_thTestPassed.load());
 }

tests/catch/unit/stream/CMakeLists.txt

@@ -7,6 +7,7 @@ set(TEST_SRC
     hipStreamAddCallback.cc
     hipStreamCreateWithFlags.cc
     hipStreamCreateWithPriority.cc
+    hipStreamDestroy.cc
     hipStreamGetCUMask.cc
     hipAPIStreamDisable.cc
     streamCommon.cc
@@ -25,6 +26,7 @@ set(TEST_SRC
     hipStreamAddCallback.cc
     hipStreamCreateWithFlags.cc
     hipStreamCreateWithPriority.cc
+    hipStreamDestroy.cc
     hipAPIStreamDisable.cc
     # hipStreamAttachMemAsync.cc   # Disabling it on nvidia due to issue in function definition of hipStreamAttachMemAsync
                                    # Fixing would break ABI, to be re-enabled when the fix is made.

tests/catch/unit/stream/hipStreamCreateWithFlags.cc

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2022 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
@@ -17,23 +17,69 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
-/**
-Testcase Scenarios :
-1) Validates functionality of hipStreamCreateWithFlags when stream = nullptr.
-2) Validates functionality of hipStreamCreateWithFlags when flag = 0xffffffff.
-*/
-
+#include <chrono>
 #include <hip_test_common.hh>
 
+namespace hipStreamCreateWithFlagsTests {
 
-TEST_CASE("Unit_hipStreamCreateWithFlags_ArgValidation") {
-  // stream = nullptr test
-  SECTION("stream is nullptr") {
-    REQUIRE(hipStreamCreateWithFlags(nullptr, hipStreamDefault) != hipSuccess);
-  }
-  // flag value invalid test
-  SECTION("flag value invalid") {
-    hipStream_t stream;
-    REQUIRE(hipStreamCreateWithFlags(&stream, 0xffffffff) != hipSuccess);
-  }
+TEST_CASE("Unit_hipStreamCreateWithFlags_Negative_NullStream") {
+  HIP_CHECK_ERROR(hipStreamCreateWithFlags(nullptr, hipStreamDefault), hipErrorInvalidValue);
 }
+
+TEST_CASE("Unit_hipStreamCreateWithFlags_Negative_InvalidFlag") {
+  hipStream_t stream{};
+  unsigned int flag = 0xFF;
+  REQUIRE(flag != hipStreamDefault);
+  REQUIRE(flag != hipStreamNonBlocking);
+  HIP_CHECK_ERROR(hipStreamCreateWithFlags(&stream, flag), hipErrorInvalidValue);
+}
+
+// create a stream and check the properties are correctly set
+TEST_CASE("Unit_hipStreamCreateWithFlags_Default") {
+  const unsigned int flagUnderTest = GENERATE(hipStreamDefault, hipStreamNonBlocking);
+  hipStream_t stream{};
+  HIP_CHECK(hipStreamCreateWithFlags(&stream, flagUnderTest));
+
+  unsigned int flag{};
+  HIP_CHECK(hipStreamGetFlags(stream, &flag));
+  REQUIRE(flag == flagUnderTest);
+
+  int priority{};
+  HIP_CHECK(hipStreamGetPriority(stream, &priority));
+  // zero is considered default priority
+  REQUIRE(priority == 0);
+
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+
+// a stream will default to blocking the null stream, but will not block the null stream when
+// created with hipStreamNonBlocking
+#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
+TEST_CASE("Unit_hipStreamCreateWithFlags_DefaultStreamInteraction") {
+  const hipStream_t defaultStream = GENERATE(static_cast<hipStream_t>(nullptr), hipStreamPerThread);
+  const unsigned int flagUnderTest = GENERATE(hipStreamDefault, hipStreamNonBlocking);
+  const hipError_t expectedError = (flagUnderTest == hipStreamDefault) && (defaultStream == nullptr)
+      ? hipErrorNotReady
+      : hipSuccess;
+  CAPTURE(defaultStream, flagUnderTest, expectedError, hipGetErrorString(expectedError));
+
+  hipStream_t stream{};
+  HIP_CHECK(hipStreamCreateWithFlags(&stream, flagUnderTest));
+
+  constexpr auto delay = std::chrono::milliseconds(500);
+
+  SECTION("default stream waiting for created stream") {
+    HipTest::runKernelForDuration(delay, stream);
+    REQUIRE(hipStreamQuery(defaultStream) == expectedError);
+  }
+  SECTION("created stream waiting for default stream") {
+    HipTest::runKernelForDuration(delay, defaultStream);
+    REQUIRE(hipStreamQuery(stream) == expectedError);
+  }
+
+  HIP_CHECK(hipDeviceSynchronize());
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+#endif
+
+}  // namespace hipStreamCreateWithFlagsTests

tests/catch/unit/stream/hipStreamDestroy.cc

@@ -0,0 +1,97 @@
+/*
+Copyright (c) 2022 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.
+*/
+#include <chrono>
+#include <hip_test_common.hh>
+
+namespace hipStreamDestroyTests {
+
+TEST_CASE("Unit_hipStreamDestroy_Default") {
+  hipStream_t stream{};
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+
+TEST_CASE("Unit_hipStreamDestroy_Negative_DoubleDestroy") {
+  hipStream_t stream{};
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK_ERROR(hipStreamDestroy(stream), hipErrorContextIsDestroyed);
+}
+
+TEST_CASE("Unit_hipStreamDestroy_Negative_NullStream") {
+  HIP_CHECK_ERROR(hipStreamDestroy(nullptr), hipErrorInvalidResourceHandle);
+}
+
+template <size_t numDataPoints> void checkDataSet(int* deviceData) {
+  HIP_CHECK(hipStreamSynchronize(nullptr));
+  std::array<int, numDataPoints> hostData{};
+  HIP_CHECK(
+      hipMemcpy(hostData.data(), deviceData, sizeof(int) * numDataPoints, hipMemcpyDeviceToHost));
+  REQUIRE(std::all_of(std::begin(hostData), std::end(hostData), [](int x) { return x == 1; }));
+}
+
+__global__ void setToOne(int* x, size_t size) {
+  unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx < size) {
+    x[idx] = 1;
+  }
+}
+
+TEST_CASE("Unit_hipStreamDestroy_WithFinishedWork") {
+  hipStream_t stream{};
+  HIP_CHECK(hipStreamCreate(&stream));
+
+  constexpr int numDataPoints = 10;
+  int* deviceData{};
+  HIP_CHECK(hipMalloc(&deviceData, sizeof(int) * numDataPoints));
+  HIP_CHECK(hipMemset(deviceData, 0, sizeof(int) * numDataPoints));
+
+  setToOne<<<1, numDataPoints, 0, stream>>>(deviceData, numDataPoints);
+  checkDataSet<numDataPoints>(deviceData);
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipFree(deviceData));
+}
+
+// hipStreamDestroy should return immediately then clean up the resources when the stream is empty
+// of work
+#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
+TEST_CASE("Unit_hipStreamDestroy_WithPendingWork") {
+#if HT_AMD
+  HipTest::HIP_SKIP_TEST(
+      "EXSWCPHIPT-44 - expected hipStreamDestroy to return immediately then release the resources "
+      "when the queued jobs are finished");
+  return;
+#endif
+  hipStream_t stream{};
+  HIP_CHECK(hipStreamCreate(&stream));
+  constexpr int numDataPoints = 10;
+  int* deviceData{};
+  HIP_CHECK(hipMalloc(&deviceData, sizeof(int) * numDataPoints));
+  HIP_CHECK(hipMemset(deviceData, 0, sizeof(int) * numDataPoints));
+
+  HipTest::runKernelForDuration(std::chrono::milliseconds(500), stream);
+  setToOne<<<1, numDataPoints, 0, stream>>>(deviceData, numDataPoints);
+  HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorNotReady);
+  HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady);
+  HIP_CHECK(hipStreamDestroy(stream));
+  checkDataSet<numDataPoints>(deviceData);
+  HIP_CHECK(hipFree(deviceData));
+}
+#endif
+}  // namespace hipStreamDestroyTests