Merge 'develop' into 'amd-staging'

Change-Id: Iec27b57de521a2639ca4495f0b4f8ce155ad81f2
2023-03-01 00:10:54 +00:00
@@ -141,7 +141,7 @@ void MemcpyDeviceToDeviceShell(F memcpy_func, const hipStream_t kernel_stream =
    HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
    if (!can_access_peer) {
      INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device);
-      REQUIRE(can_access_peer);
+      return;
    }
    HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0));
  }
@@ -87,6 +87,8 @@ TEST_CASE("Unit_hipClassKernel_Friend") {
                      0,
                      0,
                      result_ecd);
+  HIP_CHECK(hipStreamSynchronize(nullptr));
+  HIP_CHECK(hipFree(result_ecd));
 }

 // check sizeof empty class is 1
@@ -110,8 +110,6 @@ TEST_CASE("Unit_hipEventSynchronize_NoEventRecord_Positive") {

  // Record the end_event
  HIP_CHECK(hipEventRecord(end_event, NULL));
-  // End event has not been completed
-  HIP_CHECK_ERROR(hipEventQuery(end_event), hipErrorNotReady);

  // When hipEventSynchronized is called on event that has not been recorded,
  // the function returns immediately
@@ -83,6 +83,7 @@ set(TEST_SRC
  hipGraphDebugDotPrint.cc
  hipGraphCloneComplx.cc
  hipGraphNodeSetEnabled.cc
+  hipGraphNodeGetEnabled.cc
  hipGraphCreate.cc
  hipGraphDestroy.cc
  hipGraphExecDestroy.cc
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2023 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
@@ -31,14 +31,27 @@ Negative Testcase Scenarios for api hipGraphAddMemsetNode :
 10) Pass hipMemsetParams::dst as nullptr should return error code.
 11) Pass hipMemsetParams::element size other than 1, 2, or 4 and check api should return error code.
 12) Pass hipMemsetParams::height as zero and check api should return error code.
+Functional Scenarios for api hipGraphAddMemsetNode :
+1. Allocate a 2D array using hipMallocPitch. Initialize the allocated memory using hipGraphAddMemsetNode.
+   Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results
+2. Allocate a 1D array using hipMallocPitch. Initialize the allocated memory using hipGraphAddMemsetNode.
+   Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results..
+3. Allocate a 2D array using hipMalloc3D. Initialize the allocated memory using hipGraphAddMemsetNode.
+   Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results.
+4. Allocate a 1D array using hipMalloc3D. Initialize the allocated memory using hipGraphAddMemsetNode.
+   Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results.
+5. Allocate a 1D array using hipMalloc. Initialize the allocated memory using hipGraphAddMemsetNode.
+   Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results.
+6. Allocate memory using hipMallocManaged. Initialize the allocated memory using hipGraphAddMemsetNode.
+   Copy the values in device memory to host using hipGraphAddMemcpyNode. Verify the results.
 */

 #include <hip_test_common.hh>
-
 /**
 * Negative Test for API hipGraphAddMemsetNode
 */
-
+#define SIZE 1024
+static char memSetVal = 'a';
 TEST_CASE("Unit_hipGraphAddMemsetNode_Negative") {
  hipError_t ret;
  hipGraph_t graph;
@@ -121,3 +134,416 @@ TEST_CASE("Unit_hipGraphAddMemsetNode_Negative") {
  HIP_CHECK(hipFree(devData));
  HIP_CHECK(hipGraphDestroy(graph));
 }
+/*
+ * Allocate a 2D array using hipMallocPitch. Initialize the allocated memory
+ * using hipGraphAddMemsetNode. Copy the values in device memory to host using
+ * hipGraphAddMemcpyNode. Verify the results.
+*/
+TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocPitch_2D") {
+  size_t width = SIZE * sizeof(char), numW{SIZE},
+         numH{SIZE}, pitch_A;
+  char *A_d;
+
+  hipGraph_t graph;
+  std::vector<hipGraphNode_t> nodeDependencies;
+  // Host memory.
+  char* A_h = new char[numW * numH];
+  for (size_t i = 0; i < numW; i++) {
+    for (size_t j = 0; j < numH; j++) {
+      *(A_h + i * numH + j) = ' ';
+    }
+  }
+  // 2D Memory allocation hipMallocPitch
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width,
+                          numH));
+  // Create Graph
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memsetNode, memcpyNode;
+  // Add MemSet Node
+  hipMemsetParams memsetParams{};
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void *>(A_d);
+  memsetParams.value = memSetVal;
+  memsetParams.pitch = pitch_A;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = numW;
+  memsetParams.height = numH;
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+          &memsetParams));
+  nodeDependencies.push_back(memsetNode);
+  // Add MemCpy Node
+  hipMemcpy3DParms myparms{};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.srcPtr = make_hipPitchedPtr(A_d, pitch_A, numW, numH);
+  myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
+  myparms.extent = make_hipExtent(width, numH, 1);
+  myparms.kind = hipMemcpyDeviceToHost;
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
+                                            nodeDependencies.size(), &myparms));
+  nodeDependencies.clear();
+  // Create executable graph
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
+                                nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verfication
+  for (size_t i = 0; i < numW; i++) {
+    for (size_t j = 0; j < numH; j++) {
+      REQUIRE(*(A_h + i*numH + j) == memSetVal);
+    }
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  delete[] A_h;
+  HIP_CHECK(hipFree(A_d));
+}
+/*
+ * Allocate a 1D array using hipMallocPitch. Initialize the allocated memory using
+ * hipGraphAddMemsetNode. Copy the values in device memory to host using
+ * hipGraphAddMemcpyNode. Verify the results.
+*/
+TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocPitch_1D") {
+  size_t width = SIZE * sizeof(char), numW{SIZE}, pitch_A;
+  char *A_d;
+
+  // Initialize the host memory
+  std::vector<char> A_h(numW, ' ');
+
+  hipGraph_t graph;
+  std::vector<hipGraphNode_t> nodeDependencies;
+  // 1D Memory allocation hipMallocPitch
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width,
+                          1));
+  // Create Graph
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memsetNode, memcpyNode;
+  // Add MemSet Node
+  hipMemsetParams memsetParams{};
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void *>(A_d);
+  memsetParams.value = memSetVal;
+  memsetParams.pitch = pitch_A;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = numW;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+          &memsetParams));
+  nodeDependencies.push_back(memsetNode);
+  // Add MemCpy Node
+  hipMemcpy3DParms myparms{};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.srcPtr = make_hipPitchedPtr(A_d, pitch_A, numW, 1);
+  myparms.dstPtr = make_hipPitchedPtr(A_h.data(), width, numW, 1);
+  myparms.extent = make_hipExtent(width, 1, 1);
+  myparms.kind = hipMemcpyDeviceToHost;
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
+                                            nodeDependencies.size(), &myparms));
+  nodeDependencies.clear();
+
+  // Create executable graph
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
+                                nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verfication
+  for (size_t i = 0; i < numW; i++) {
+    REQUIRE(A_h[i] == memSetVal);
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  HIP_CHECK(hipFree(A_d));
+}
+/*
+ * Allocate a 2D array using hipMalloc3D. Initialize the allocated memory using
+ * hipGraphAddMemsetNode. Copy the values in device memory to host using
+ * hipGraphAddMemcpyNode. Verify the results.
+*/
+TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc3D_2D") {
+  size_t width = SIZE * sizeof(char);
+  size_t numW = SIZE, numH = SIZE;
+
+  // Host Memory
+  char* A_h = new char[numW * numH];
+  for (size_t i = 0; i < numW; i++) {
+    for (size_t j = 0; j < numH; j++) {
+      *(A_h + i * numH + j) = ' ';
+    }
+  }
+  hipGraph_t graph;
+  std::vector<hipGraphNode_t> nodeDependencies;
+
+  hipPitchedPtr A_d;
+  hipExtent extent3D = make_hipExtent(width, numH, 1);
+
+  // Allocate 3D memory.
+  HIPCHECK(hipMalloc3D(&A_d, extent3D));
+
+  // Create Graph
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memsetNode, memcpyNode;
+
+  // Add MemSet Node
+  hipMemsetParams memsetParams{};
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = A_d.ptr;
+  memsetParams.value = memSetVal;
+  memsetParams.pitch = A_d.pitch;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = numW;
+  memsetParams.height = numH;
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+          &memsetParams));
+  nodeDependencies.push_back(memsetNode);
+
+  // MemCpy params
+  hipMemcpy3DParms myparms{};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.srcPtr = A_d;
+  myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH);
+  myparms.extent = make_hipExtent(width, numH, 1);
+  myparms.kind = hipMemcpyDeviceToHost;
+
+  // Add MemCpy Node
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
+                                            nodeDependencies.size(), &myparms));
+  nodeDependencies.clear();
+
+  // Create executable graph
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
+                                nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verfication
+  for (size_t i = 0; i < numW; i++) {
+    for (size_t j = 0; j < numH; j++) {
+      REQUIRE(*(A_h + i*numH + j) == memSetVal);
+    }
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  delete[] A_h;
+  HIP_CHECK(hipFree(A_d.ptr));
+}
+/*
+ * Allocate a 1D array using hipMalloc3D. Initialize the allocated
+ * memory using hipGraphAddMemsetNode. Copy the values in device
+ * memory to host using hipGraphAddMemcpyNode. Verify the results.
+*/
+TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc3D_1D") {
+  size_t width = SIZE * sizeof(char);
+  size_t numW = SIZE;
+
+  // Initialize the host memory
+  std::vector<char> A_h(numW, ' ');
+
+  hipGraph_t graph;
+  std::vector<hipGraphNode_t> nodeDependencies;
+
+  hipPitchedPtr A_d;
+  hipExtent extent1D = make_hipExtent(width, 1, 1);
+
+  // Allocate 3D memory.
+  HIPCHECK(hipMalloc3D(&A_d, extent1D));
+
+  // Create Graph
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipGraphNode_t memsetNode, memcpyNode;
+
+  // Add MemSet Node
+  hipMemsetParams memsetParams{};
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = A_d.ptr;
+  memsetParams.value = memSetVal;
+  memsetParams.pitch = A_d.pitch;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = numW;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+          &memsetParams));
+  nodeDependencies.push_back(memsetNode);
+
+  // MemCpy params
+  hipMemcpy3DParms myparms{};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.srcPtr = A_d;
+  myparms.dstPtr = make_hipPitchedPtr(A_h.data(), width, numW, 1);
+  myparms.extent = make_hipExtent(width, 1, 1);
+  myparms.kind = hipMemcpyDeviceToHost;
+
+  // Add MemCpy Node
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
+                                            nodeDependencies.size(), &myparms));
+  nodeDependencies.clear();
+
+  // Create executable graph
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
+                                nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verfication
+  for (size_t i = 0; i < numW; i++) {
+     REQUIRE(A_h[i] == memSetVal);
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph))
+  HIP_CHECK(hipFree(A_d.ptr));
+}
+/*
+ * Allocate a 1D array using hipMalloc. Initialize the allocated memory using
+ * hipGraphAddMemsetNode. Copy the values in device memory to host using
+ * hipGraphAddMemcpyNode. Verify the results.
+*/
+TEST_CASE("Unit_hipGraphAddMemsetNode_hipMalloc_1D") {
+  char *A_d;
+  size_t NumW = SIZE;
+  size_t Nbytes1D = SIZE * sizeof(char);
+
+  // Initialize the host memory
+  std::vector<char> A_h(NumW, ' ');
+
+  // Allocate memory to Device pointer
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), Nbytes1D));
+
+  // Create the graph
+  hipGraph_t graph;
+  std::vector<hipGraphNode_t> nodeDependencies;
+  hipGraphNode_t memsetNode, memcpyNode;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // Add Memset node
+  hipMemsetParams memsetParams{};
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void *>(A_d);
+  memsetParams.value = memSetVal;
+  memsetParams.pitch = Nbytes1D;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = NumW;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+            &memsetParams));
+  nodeDependencies.push_back(memsetNode);
+  // Add MemCpy Node
+  hipPitchedPtr devPitchedPtr{A_d, Nbytes1D, NumW, 0};
+  hipPitchedPtr hostPitchedPtr{A_h.data(), Nbytes1D, NumW, 0};
+  hipMemcpy3DParms myparms{};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.srcPtr = devPitchedPtr;
+  myparms.dstPtr = hostPitchedPtr;
+  myparms.extent = make_hipExtent(Nbytes1D, 1, 1);
+  myparms.kind = hipMemcpyDeviceToHost;
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
+                                            nodeDependencies.size(), &myparms));
+  nodeDependencies.clear();
+  // Create executable graph
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
+                                nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verfication
+  for (size_t i = 0; i < NumW; i++) {
+     REQUIRE(A_h[i] == memSetVal);
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  HIP_CHECK(hipFree(A_d));
+}
+
+TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocManaged") {
+  int managed = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&managed,
+                                  hipDeviceAttributeManagedMemory, 0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+  if (managed != 1) {
+    WARN(
+      "GPU 0 doesn't support hipDeviceAttributeManagedMemory attribute"
+       "so defaulting to system memory.");
+  }
+  size_t Nbytes1D = SIZE * sizeof(char);
+  char *A_d;
+  // Initialize the host memory
+  std::vector<char> A_h(SIZE, ' ');
+  // Device Memory
+  HIP_CHECK(hipMallocManaged(&A_d, SIZE * sizeof(char)));
+  // Create the graph
+  hipGraph_t graph;
+  std::vector<hipGraphNode_t> nodeDependencies;
+  hipGraphNode_t memsetNode, memcpyNode;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+
+  // Add Memset node
+  hipMemsetParams memsetParams{};
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void *>(A_d);
+  memsetParams.value = memSetVal;
+  memsetParams.pitch = Nbytes1D;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = SIZE;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+            &memsetParams));
+  nodeDependencies.push_back(memsetNode);
+
+  // Add MemCpy Node
+  hipPitchedPtr devPitchedPtr{A_d, Nbytes1D, SIZE, 1};
+  hipPitchedPtr hostPitchedPtr{A_h.data(), Nbytes1D, SIZE, 1};
+
+  hipMemcpy3DParms myparms{};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.srcPtr = devPitchedPtr;
+  myparms.dstPtr = hostPitchedPtr;
+  myparms.extent = make_hipExtent(Nbytes1D, 1, 1);
+  myparms.kind = hipMemcpyDeviceToHost;
+  HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(),
+                                            nodeDependencies.size(), &myparms));
+  nodeDependencies.clear();
+
+  // Create executable graph
+  hipStream_t streamForGraph;
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr,
+                                nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verfication
+  for (size_t i = 0; i < SIZE; i++) {
+     REQUIRE(A_h[i] == memSetVal);
+  }
+
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+  HIP_CHECK(hipFree(A_d));
+}
@@ -282,6 +282,7 @@ void hipGraphExecMemcpyNodeSetParamsToSymbol_GlobalMem(bool useConstVar) {
  }

  HIP_CHECK(hipGraphLaunch(graphExec, 0));
+  HIP_CHECK(hipStreamSynchronize(0));

  // Validating the result
  for (int i = 0; i < SIZE; i++) {
@@ -0,0 +1,257 @@
+/*
+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 <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+#define N   1024
+
+/**
+ * Functional Test for API - hipGraphNodeGetEnabled
+ 1) Add MemCpy node to the graph and verify in graphExec it enabled status
+ 2) Add MemSet node to the graph and verify in graphExec it enabled status
+ 3) Add Kernel node to the graph and verify in graphExec it enabled status
+- Can check other node type, as only above 3 types are mentioned in the doc
+ 4) Add HostNode node to the graph and verify in graphExec it enabled status
+ 5) Add emptyNode node to the graph and verify in graphExec it enabled status
+ 6) Add ChildNode node to the graph and verify in graphExec it enabled status
+ 7) Add EventWait node to the graph and verify in graphExec it enabled status
+ 8) Add EventRecord node to the graph and verify in graphExec it enabled status
+ */
+
+static void callbackfunc(void *A_h) {
+  int *A = reinterpret_cast<int *>(A_h);
+  for (int i = 0; i < N; i++) {
+    A[i] = i;
+  }
+}
+
+TEST_CASE("Unit_hipGraphNodeGetEnabled_Functional_Basic") {
+  constexpr size_t Nbytes = N * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  hipGraph_t graph, childGraph;
+  hipGraphNode_t memcpy_A, memcpy_B, memsetNode, kNodeAdd;
+  hipGraphNode_t hostNode, emptyNode, childGraphNode, eventWait, eventRecord;
+  int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h;
+  hipGraphExec_t graphExec;
+  size_t NElem{N};
+  unsigned int isEnabled = 0;
+  hipError_t ret;
+
+  HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+  unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+
+  void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  hipKernelNodeParams kNodeParams{};
+  kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kNodeParams.gridDim = dim3(blocks);
+  kNodeParams.blockDim = dim3(threadsPerBlock);
+  kNodeParams.sharedMemBytes = 0;
+  kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
+  kNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kNodeAdd, graph, nullptr, 0, &kNodeParams));
+
+  hipMemsetParams memsetParams{};
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(B_d);
+  memsetParams.value = 7;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+  HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
+                                  &memsetParams));
+
+  HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, NULL, 0));
+
+  hipHostNodeParams hostParams = {0, 0};
+  hostParams.fn = callbackfunc;
+  hostParams.userData = A_h;
+  HIP_CHECK(hipGraphAddHostNode(&hostNode, graph, nullptr, 0, &hostParams));
+
+  hipEvent_t event;
+  HIP_CHECK(hipEventCreate(&event));
+  HIP_CHECK(hipGraphAddEventRecordNode(&eventRecord, graph, nullptr,
+                                       0, event));
+  HIP_CHECK(hipGraphAddEventWaitNode(&eventWait, graph, nullptr, 0, event));
+
+  HIP_CHECK(hipGraphCreate(&childGraph, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, childGraph, nullptr, 0,
+                                    B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+  // Adding child node to clonedGraph
+  HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode, graph,
+                                      nullptr, 0, childGraph));
+
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
+
+  SECTION("Create graphExec with Kernel node and verify its enabled status") {
+    HIP_CHECK(hipGraphNodeGetEnabled(graphExec, kNodeAdd, &isEnabled));
+    REQUIRE(1 == isEnabled);
+  }
+  SECTION("Create graphExec with MemCpy node and verify its enabled status") {
+    HIP_CHECK(hipGraphNodeGetEnabled(graphExec, memcpy_A, &isEnabled));
+    REQUIRE(1 == isEnabled);
+  }
+  SECTION("Create graphExec with MemSet node and verify its enabled status") {
+    HIP_CHECK(hipGraphNodeGetEnabled(graphExec, memsetNode, &isEnabled));
+    REQUIRE(1 == isEnabled);
+  }
+  SECTION("Create graphExec with hostNode and verify its enabled status") {
+    ret = hipGraphNodeGetEnabled(graphExec, hostNode, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Create graphExec with emptyNode and verify its enabled status") {
+    ret = hipGraphNodeGetEnabled(graphExec, emptyNode, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Create graphExec with ChildGraphNode & verify its enabled status") {
+    ret = hipGraphNodeGetEnabled(graphExec, childGraphNode, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Create graphExec with EventWait and verify its enabled status") {
+    ret = hipGraphNodeGetEnabled(graphExec, eventWait, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Create graphExec with EventRecord and verify its enabled status") {
+    ret = hipGraphNodeGetEnabled(graphExec, eventRecord, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(childGraph));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+}
+
+/**
+ Negative Test for API - hipGraphNodeGetEnabled
+ 1) Pass graphExec as nullptr
+ 2) Pass graphExec as uninitialized object
+ 3) Pass Node as nullptr
+ 4) Pass Node as uninitialized object
+ 5) Pass isEnabled as nullptr
+
+ Negative Functional Test for API - hipGraphNodeGetEnabled
+ 6) Pass hNode from different graph and verify
+ 7) Create graphExec and then add one more new node to the graph verify
+ 8) Pass hNode a deleted node from same graph where exec was created
+ 9) Create graphExec and then delete the graph and verify a node
+ 10) Create graphExec and then delete the graphExec and verify a node
+ */
+
+TEST_CASE("Unit_hipGraphNodeGetEnabled_Negative_Functional") {
+  constexpr size_t Nbytes = N * sizeof(int);
+  constexpr auto blocksPerCU = 6;  // to hide latency
+  constexpr auto threadsPerBlock = 256;
+  hipGraph_t graph, graph2;
+  hipGraphNode_t memcpy_A, memcpy_B, memcpy_A2, memcpy_C, kNodeAdd;
+  hipKernelNodeParams kNodeParams{};
+  int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h;
+  hipGraphExec_t graphExec, graphExec2;
+  size_t NElem{N};
+  unsigned int isEnabled;
+  hipError_t ret;
+
+  HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+  unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+
+  void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
+  kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  kNodeParams.gridDim = dim3(blocks);
+  kNodeParams.blockDim = dim3(threadsPerBlock);
+  kNodeParams.sharedMemBytes = 0;
+  kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
+  kNodeParams.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kNodeAdd, graph, nullptr, 0, &kNodeParams));
+
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kNodeAdd, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kNodeAdd, 1));
+
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
+  HIP_CHECK(hipGraphInstantiate(&graphExec2, graph, NULL, NULL, 0));
+
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
+                                    Nbytes, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipGraphAddDependencies(graph, &kNodeAdd, &memcpy_C, 1));
+
+  HIP_CHECK(hipGraphCreate(&graph2, 0));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A2, graph2, nullptr, 0, A_d, A_h,
+                                    Nbytes, hipMemcpyHostToDevice));
+
+  SECTION("Pass hGraphExec as nullptr") {
+    ret = hipGraphNodeGetEnabled(nullptr, memcpy_B, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass hGraphExec as uninitialized object") {
+    hipGraphExec_t graphExec_uninit{};
+    ret = hipGraphNodeGetEnabled(graphExec_uninit, memcpy_B, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass Node as nullptr") {
+    ret = hipGraphNodeGetEnabled(graphExec, nullptr, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass Node as uninitialized object") {
+    hipGraphNode_t node_uninit{};
+    ret = hipGraphNodeGetEnabled(graphExec, node_uninit, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass isEnabled as nullptr") {
+    ret = hipGraphNodeGetEnabled(graphExec, memcpy_B, nullptr);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass hNode from different graph and verify") {
+    ret = hipGraphNodeGetEnabled(graphExec, memcpy_A2, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Create graphExec and add one more new node to the graph & verify") {
+    ret = hipGraphNodeGetEnabled(graphExec, memcpy_C, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  SECTION("Pass hNode a deleted node from same graph where exec was created") {
+    HIP_CHECK(hipGraphDestroyNode(memcpy_A));
+    ret = hipGraphNodeGetEnabled(graphExec, memcpy_A, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  HIP_CHECK(hipGraphExecDestroy(graphExec2));
+  SECTION("Create graphExec and then delete the graphExec and verify a node") {
+    ret = hipGraphNodeGetEnabled(graphExec2, memcpy_B, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  HIP_CHECK(hipGraphDestroy(graph));
+  SECTION("Create graphExec and then delete the graph and verify a node") {
+    ret = hipGraphNodeGetEnabled(graphExec, memcpy_B, &isEnabled);
+    REQUIRE(hipErrorInvalidValue == ret);
+  }
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph2));
+}
@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2023 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
@@ -161,8 +161,17 @@ TEST_CASE("Unit_hipGraphUpload_Functional_multidevice_test") {
    SECTION("Pass a common stream for all device") {
      hipStream_t stream;
      HIP_CHECK(hipStreamCreate(&stream));
-
+      int currDevice = -1;
+      HIP_CHECK(hipGetDevice(&currDevice));
      for (int i = 0; i < numDevices; i++) {
+        if (i != currDevice) {
+          int can_access_peer = 0;
+          HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, currDevice, i));
+          if (!can_access_peer) {
+            INFO("Peer access cannot be enabled between devices " << currDevice << " " << i);
+            continue;
+          }
+        }
        HIP_CHECK(hipSetDevice(i));
        hipGraphUploadFunctional_with_hipStreamBeginCapture(stream);
        hipGraphUploadFunctional_with_stream(stream);
@@ -200,6 +209,53 @@ TEST_CASE("Unit_hipGraphUpload_Functional_multidevice_test") {
  }
 }

+/**
+ * Functional Test for API - hipGraphUpload
+ - Make graph by using hipStreamBeginCapture with a low priority stream.
+   Upload the graph into high priority stream and execute the graph and verify.
+ */
+
+TEST_CASE("Unit_hipGraphUpload_Functional_With_Priority_Stream") {
+  constexpr size_t N = 1024;
+  constexpr size_t Nbytes = N * sizeof(int);
+  hipGraph_t graph;
+  hipGraphExec_t graphExec;
+
+  int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h;
+  HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
+
+  hipStream_t stream1, stream2;
+  int minPriority = 0, maxPriority = 0;
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&minPriority, &maxPriority));
+  HIP_CHECK(hipStreamCreateWithPriority(&stream1, hipStreamDefault,
+                                        minPriority));
+  HIP_CHECK(hipStreamCreateWithPriority(&stream2, hipStreamDefault,
+                                        maxPriority));
+
+  HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
+  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream1));
+  HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream1));
+  HipTest::vectorADD<int><<<1, 1, 0, stream1>>>(A_d, B_d, C_d, N);
+  HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream1));
+  HIP_CHECK(hipStreamEndCapture(stream1, &graph));
+
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+
+  HIP_CHECK(hipGraphUpload(graphExec, stream2));
+
+  HIP_CHECK(hipGraphLaunch(graphExec, stream2));
+  HIP_CHECK(hipStreamSynchronize(stream2));
+
+  // Verify graph execution result
+  HipTest::checkVectorADD(A_h, B_h, C_h, N);
+
+  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(stream1));
+  HIP_CHECK(hipStreamDestroy(stream2));
+}
+
 /**
 * Negative Test for API - hipGraphUpload Argument Check
 1) Pass graphExec node as nullptr.
@@ -233,4 +289,3 @@ TEST_CASE("Unit_hipGraphUpload_Negative_Argument_Check") {
  }
  HIP_CHECK(hipStreamDestroy(stream));
 }
-
@@ -30,6 +30,23 @@ Negative Testcase Scenarios :
 7) Begin capture on a thread with mode other than hipStreamCaptureModeRelaxed
   and try to end capture from different thread. Expect to return
   hipErrorStreamCaptureWrongThread.
+8) Start stream capture on stream1 using mode hipStreamCaptureModeRelaxed.
+   In stream1 queue a memcpy operation, queue a kernel square of a number operation.
+   Launch a thread. In the thread, queue a memcpy operation. End the capture on
+   stream1 and return the captured graph. Wait for the thread in main function.
+   Create an executable graph and launch the graph on input data and validate the
+   output.
+9) Create 2 streams s1 and s2. Begin stream capture in s1, spawn a
+   captured fork stream on s2. Queue some operations
+   (like increment kernel) on both s1 and s2. End the stream capture
+   on s2 and verify the error returned by the End capture.
+10)Create 2 streams s1 and s2. Begin stream capture in s1 and spawn a captured
+   fork stream s2. In main thread, queue a memcpy operation on s1.
+   Launch a thread, queue a memcpy operation on s2. Perform hipEventRecord on
+   s2 and wait Event on S1.  Wait for the thread to complete. Queue operations
+   kernel addition(Cd = Ad + Bd) operation and memcpy(Ch <- Cd) in s1. End the
+   stream capture in s1. Create an executable graph and launch the graph on input
+   data and validate the output.
 */

 #include <hip_test_common.hh>
@@ -123,14 +140,37 @@ static void thread_func(hipStream_t stream, hipGraph_t graph) {
  HIP_ASSERT(hipErrorStreamCaptureWrongThread ==
             hipStreamEndCapture(stream, &graph));
 }
-
-TEST_CASE("Unit_hipStreamEndCapture_Thread_Negative") {
+static void StreamEndCaptureThreadNegative(float* A_d, float* A_h,
+                float* C_d, float* C_h, hipStreamCaptureMode mode) {
  hipStream_t stream{nullptr};
  hipGraph_t graph{nullptr};
  constexpr unsigned blocks = 512;
  constexpr unsigned threadsPerBlock = 256;
  constexpr size_t N = 100000;
  size_t Nbytes = N * sizeof(float);
+
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipStreamBeginCapture(stream, mode));
+  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
+
+  HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream));
+  hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                              dim3(threadsPerBlock), 0, stream, A_d, C_d, N);
+  HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
+
+  std::thread t(thread_func, stream, graph);
+  t.join();
+
+#if HT_AMD
+  HIP_CHECK(hipStreamEndCapture(stream, &graph));
+#endif
+  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipGraphDestroy(graph));
+}
+TEST_CASE("Unit_hipStreamEndCapture_Thread_Negative") {
+  constexpr size_t N = 100000;
+  size_t Nbytes = N * sizeof(float);
  float *A_d, *C_d;
  float *A_h, *C_h;

@@ -149,28 +189,249 @@ TEST_CASE("Unit_hipStreamEndCapture_Thread_Negative") {
  REQUIRE(A_d != nullptr);
  REQUIRE(C_d != nullptr);

-  HIP_CHECK(hipStreamCreate(&stream));
-  HIP_CHECK(hipGraphCreate(&graph, 0));
-  HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
-  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
-
-  HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream));
-  hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
-                              dim3(threadsPerBlock), 0, stream, A_d, C_d, N);
-  HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
-
-  std::thread t(thread_func, stream, graph);
-  t.join();
-
-#if HT_AMD
-  HIP_CHECK(hipStreamEndCapture(stream, &graph));
-#endif
-
+  SECTION("Capture Mode:hipStreamCaptureModeGlobal") {
+    StreamEndCaptureThreadNegative(A_d, A_h, C_d, C_h,
+                            hipStreamCaptureModeGlobal);
+  }
+  SECTION("Capture Mode:hipStreamCaptureModeThreadLocal") {
+    StreamEndCaptureThreadNegative(A_d, A_h, C_d, C_h,
+                       hipStreamCaptureModeThreadLocal);
+  }
  free(A_h);
  free(C_h);
  HIP_CHECK(hipFree(A_d));
  HIP_CHECK(hipFree(C_d));
+}
+// Thread function
+static void thread_func1(hipStream_t stream, hipGraph_t *graph,
+                           size_t Nbytes, float* A_d, float* B_h) {
+  HIP_CHECK(hipMemcpyAsync(B_h, A_d, Nbytes, hipMemcpyDeviceToHost, stream));
+  HIP_CHECK(hipStreamEndCapture(stream, graph));
+}
+/*
+ * Start stream capture on stream1 using mode hipStreamCaptureModeRelaxed.
+ * In stream1 queue a memcpy operation, queue a kernel square of a number operation.
+ * Launch a thread. In the thread, queue a memcpy operation. End the capture on
+ * stream1 and return the captured graph. Wait for the thread in main function.
+ * Create an executable graph and launch the graph on input data and validate the output.
+ * */
+TEST_CASE("Unit_hipStreamEndCapture_mode_hipStreamCaptureModeRelaxed") {
+  hipStream_t stream{nullptr}, streamForGraph{nullptr};
+  hipGraph_t graph{nullptr};
+  constexpr unsigned blocks = 512;
+  constexpr unsigned threadsPerBlock = 256;
+  constexpr size_t N = 10;
+  size_t Nbytes = N * sizeof(float);
+  // Device Pointers
+  float *A_d;
+  // Host Pointers
+  float *A_h, *B_h, *C_h;
+
+  // Memory allocation to Host pointers
+  A_h = reinterpret_cast<float*>(malloc(Nbytes));
+  B_h = reinterpret_cast<float*>(malloc(Nbytes));
+  C_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+  REQUIRE(B_h != nullptr);
+  REQUIRE(C_h != nullptr);
+
+  // Initialize the Host data
+  for (size_t i = 0; i < N; i++) {
+    A_h[i] =  1.0f + i;
+    C_h[i] = A_h[i];
+  }
+  // Memory allocation to Device pointers
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), Nbytes));
+  REQUIRE(A_d != nullptr);
+
+  HIP_CHECK(hipStreamCreate(&stream));
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeRelaxed));
+  // Copy data from Host to Device
+  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
+
+  hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                              dim3(threadsPerBlock), 0, stream, A_d, A_d, N);
+  // Thread Launch
+  std::thread t(thread_func1, stream, &graph, Nbytes, A_d, B_h);
+  t.join();
+
+  // Launch the graph
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Output verification
+  for (size_t i = 0; i < N; i++) {
+    C_h[i] = C_h[i] * C_h[i];
+    REQUIRE(B_h[i] == C_h[i]);
+  }
+
+  free(A_h);
+  free(B_h);
+  free(C_h);
+  HIP_CHECK(hipFree(A_d));
  HIP_CHECK(hipStreamDestroy(stream));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
 }

+static __global__ void increment(int* A_d) {
+  atomicAdd(A_d, 1);
+}
+/*
+ * Create 2 streams s1 and s2. Begin stream capture in s1, spawn a
+ * captured fork stream on s2. Queue some operations
+ * (like increment kernel) on both s1 and s2. End the stream capture
+ * on s2 and verify the error returned by the End capture.
+*/
+TEST_CASE("Unit_hipStreamEndCapture_chkError_on_wrongStream") {
+  int *A_d{nullptr}, *A_h{nullptr};
+  hipStream_t stream1{nullptr}, stream2{nullptr};
+  hipEvent_t forkStreamEvent{nullptr};
+  hipGraph_t graph{nullptr};
+  hipError_t err;
+  constexpr unsigned blocks = 512;
+  constexpr unsigned threadsPerBlock = 256;
+  size_t Nbytes = sizeof(int);
+
+  HIP_CHECK(hipStreamCreate(&stream1));
+  HIP_CHECK(hipStreamCreate(&stream2));
+  HIP_CHECK(hipEventCreate(&forkStreamEvent));
+
+  A_h = reinterpret_cast<int*>(malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+  // Initialize the Host data
+  *A_h = 0;
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), Nbytes));
+  REQUIRE(A_d != nullptr);
+
+  HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
+  HIP_CHECK(hipEventRecord(forkStreamEvent, stream1));
+  HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
+
+  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes,
+                                        hipMemcpyHostToDevice, stream1));
+
+  hipLaunchKernelGGL(increment, dim3(blocks),
+                              dim3(threadsPerBlock), 0, stream1, A_d);
+  hipLaunchKernelGGL(increment, dim3(blocks),
+                              dim3(threadsPerBlock), 0, stream2, A_d);
+
+  err = hipStreamEndCapture(stream2, &graph);
+  REQUIRE(err == hipErrorStreamCaptureUnmatched);
+
+  HIP_CHECK(hipStreamDestroy(stream1));
+  HIP_CHECK(hipStreamDestroy(stream2));
+  HIP_CHECK(hipEventDestroy(forkStreamEvent));
+  free(A_h);
+  HIP_CHECK(hipFree(A_d));
+}
+static void thread_func4(hipStream_t stream1, hipStream_t stream2,
+                hipEvent_t event, size_t Nbytes, int* B_d, int* B_h) {
+  HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream2));
+  HIP_CHECK(hipEventRecord(event, stream2));
+  HIP_CHECK(hipStreamWaitEvent(stream1, event, 0));
+}
+/*
+ * Create 2 streams s1 and s2. Begin stream capture in s1 and spawn a captured
+ * fork stream s2. In main thread, queue a memcpy operation on s1.
+ * Launch a thread, queue a memcpy operation on s2. Perform hipEventRecord on
+ * s2 and wait Event on S1.  Wait for the thread to complete. Queue operations
+ * kernel addition(Cd = Ad + Bd) operation and memcpy(Ch <- Cd) in s1. End the
+ * stream capture in s1. Create an executable graph and launch the graph on input
+ * data and validate the output.
+ * */
+TEST_CASE("Unit_hipStreamEndCapture_streamMerge_in_thread") {
+  // Device Pointers
+  int *A_d, *B_d, *C_d;
+  // Host Pointers
+  int *A_h, *B_h, *C_h, *D_h;
+  hipStream_t stream1{nullptr}, stream2{nullptr}, streamForGraph{nullptr};
+  hipEvent_t forkStreamEvent{nullptr}, event{nullptr};
+  hipGraph_t graph{nullptr};
+
+  constexpr unsigned blocks = 512;
+  constexpr unsigned threadsPerBlock = 256;
+  constexpr size_t N = 5;
+  size_t Nbytes = N * sizeof(int);
+
+  HIP_CHECK(hipStreamCreate(&stream1));
+  HIP_CHECK(hipStreamCreate(&stream2));
+  HIP_CHECK(hipStreamCreate(&streamForGraph));
+  HIP_CHECK(hipEventCreate(&forkStreamEvent));
+  HIP_CHECK(hipEventCreate(&event));
+  // Memory allocation to Host Pointers
+  A_h = reinterpret_cast<int*>(malloc(Nbytes));
+  B_h = reinterpret_cast<int*>(malloc(Nbytes));
+  C_h = reinterpret_cast<int*>(malloc(Nbytes));
+  D_h = reinterpret_cast<int*>(malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+  REQUIRE(B_h != nullptr);
+  REQUIRE(C_h != nullptr);
+  REQUIRE(D_h != nullptr);
+  // Initialize the Host data
+  for (size_t i = 0; i < N; i++) {
+    A_h[i] = 1 + i;
+    B_h[i] = 2 + i;
+    C_h[i] = 0;
+    D_h[i] = 0;
+  }
+  // Memory allocation to Device Pointers
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), Nbytes));
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&B_d), Nbytes));
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&C_d), Nbytes));
+  REQUIRE(A_d != nullptr);
+  REQUIRE(B_d != nullptr);
+  REQUIRE(C_d != nullptr);
+
+  // Begin Capture
+  HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
+
+  HIP_CHECK(hipEventRecord(forkStreamEvent, stream1));
+  HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
+
+  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes,
+                                hipMemcpyHostToDevice, stream1));
+  // Thread Launch
+  std::thread t(thread_func4, stream1, stream2, event, Nbytes, B_d, B_h);
+  t.join();
+  // Launch kernal
+  hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks),
+                              dim3(threadsPerBlock), 0, stream1, A_d,
+                              B_d, C_d, N);
+
+  HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes,
+                                        hipMemcpyDeviceToHost, stream1));
+  HIP_CHECK(hipStreamEndCapture(stream1, &graph));
+
+  // Launch graph
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
+  HIP_CHECK(hipStreamSynchronize(streamForGraph));
+
+  // Verify Output
+  for (size_t i = 0; i < N; i++) {
+    D_h[i] = A_h[i] + B_h[i];
+    REQUIRE(C_h[i] == D_h[i]);
+  }
+
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamDestroy(stream1));
+  HIP_CHECK(hipStreamDestroy(stream2));
+  HIP_CHECK(hipEventDestroy(forkStreamEvent));
+  HIP_CHECK(hipStreamDestroy(streamForGraph));
+
+  // Release the memory
+  free(A_h);
+  free(B_h);
+  free(C_h);
+  free(D_h);
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipFree(B_d));
+  HIP_CHECK(hipFree(C_d));
+}
@@ -427,43 +427,37 @@ TEST_CASE("Unit_hipGraphRetainUserObject_Functional_2") {
  unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);

  HIP_CHECK(hipGraphCreate(&graph, 0));
-  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h,
-                                   Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, Nbytes,
+                                    hipMemcpyHostToDevice));
  dependencies.push_back(memcpyNode);
-  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h,
-                                   Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, B_d, B_h, Nbytes,
+                                    hipMemcpyHostToDevice));
  dependencies.push_back(memcpyNode);

-  void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
-  kNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
+  void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void*>(&NElem)};
+  kNodeParams.func = reinterpret_cast<void*>(HipTest::vectorADD<int>);
  kNodeParams.gridDim = dim3(blocks);
  kNodeParams.blockDim = dim3(threadsPerBlock);
  kNodeParams.sharedMemBytes = 0;
  kNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
  kNodeParams.extra = nullptr;
-  HIP_CHECK(hipGraphAddKernelNode(&kNode, graph, dependencies.data(),
-                                  dependencies.size(), &kNodeParams));
+  HIP_CHECK(
+      hipGraphAddKernelNode(&kNode, graph, dependencies.data(), dependencies.size(), &kNodeParams));

  dependencies.clear();
  dependencies.push_back(kNode);
-  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(),
-                                    dependencies.size(), C_h, C_d,
-                                    Nbytes, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, dependencies.data(), dependencies.size(),
+                                    C_h, C_d, Nbytes, hipMemcpyDeviceToHost));

-  int refCount = 2;
-  int refCountRetain = 3;
-
-  float *object = new float();
+  float* object = new float();
  REQUIRE(object != nullptr);
  hipUserObject_t hObject;

-  HIP_CHECK(hipUserObjectCreate(&hObject, object,
-                                  destroyFloatObj,
-                                  refCount, hipUserObjectNoDestructorSync));
+  HIP_CHECK(
+      hipUserObjectCreate(&hObject, object, destroyFloatObj, 1, hipUserObjectNoDestructorSync));
  REQUIRE(hObject != nullptr);
-  HIP_CHECK(hipUserObjectRetain(hObject, refCountRetain));
-  HIP_CHECK(hipGraphRetainUserObject(graph, hObject, refCountRetain,
-                                       hipGraphUserObjectMove));
+  HIP_CHECK(hipGraphRetainUserObject(graph, hObject, 1,
+                                     hipGraphUserObjectMove));  // Pass ownership to hipGraph

  // Instantiate and launch the graph
  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
@@ -473,9 +467,6 @@ TEST_CASE("Unit_hipGraphRetainUserObject_Functional_2") {
  // Verify result
  HipTest::checkVectorADD<int>(A_h, B_h, C_h, N);

-  HIP_CHECK(hipUserObjectRelease(hObject, refCount + refCountRetain));
-  HIP_CHECK(hipGraphReleaseUserObject(graph, hObject, refCountRetain));
-
  HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
  HIP_CHECK(hipGraphExecDestroy(graphExec));
  HIP_CHECK(hipGraphDestroy(graph));
@@ -7,213 +7,127 @@
 # 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 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.
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.

 set(COMMON_SHARED_SRC DriverContext.cc)

 # Common Tests - Test independent of all platforms
+set(TEST_SRC
+    memset.cc
+    malloc.cc
+    hipMemcpy2DToArray.cc
+    hipMemcpy2DToArray_old.cc
+    hipMemcpy2DToArrayAsync.cc
+    hipMemcpy2DToArrayAsync_old.cc
+    hipMemcpy3D.cc
+    hipMemcpy3DAsync.cc
+    hipMemcpyParam2D.cc
+    hipMemcpyParam2DAsync.cc
+    hipMemcpy2D.cc
+    hipMemcpy2DAsync.cc
+    hipMemcpy2DFromArray.cc
+    hipMemcpy2DFromArray_old.cc
+    hipMemcpy2DFromArrayAsync.cc
+    hipMemcpy2DFromArrayAsync_old.cc
+    hipMemcpyAtoH.cc
+    hipMemcpyAtoH_old.cc
+    hipMemcpyHtoA.cc
+    hipMemcpyHtoA_old.cc
+    hipMemcpyAllApiNegative.cc
+    hipMemcpy_MultiThread.cc
+    hipHostRegister.cc
+    hipHostUnregister.cc
+    hipHostGetFlags.cc
+    hipHostGetDevicePointer.cc
+    hipMallocManaged_MultiScenario.cc
+    hipMemsetNegative.cc
+    hipMemset.cc
+    hipMemsetAsyncMultiThread.cc
+    hipMemset3D.cc
+    hipMemset2D.cc
+    hipHostMallocTests.cc
+    hipMemset3DFunctional.cc
+    hipMemset3DRegressMultiThread.cc
+    hipMallocManagedFlagsTst.cc
+    hipMemPrefetchAsyncExtTsts.cc
+    hipMemAdviseMmap.cc
+    hipMallocManaged.cc
+    hipMemRangeGetAttribute.cc
+    hipMemRangeGetAttribute_old.cc
+    hipMemcpyFromSymbol.cc
+    hipPtrGetAttribute.cc
+    hipMemPoolApi.cc
+    hipMemcpyPeer.cc
+    hipMemcpyPeer_old.cc
+    hipMemcpyPeerAsync.cc
+    hipMemcpyPeerAsync_old.cc
+    hipMemcpyWithStream_old.cc
+    hipMemcpyWithStream.cc
+    hipMemcpyWithStreamMultiThread.cc
+    hipMemsetAsyncAndKernel.cc
+    hipMemset2DAsyncMultiThreadAndKernel.cc
+    hipMallocConcurrency.cc
+    hipMemcpyDtoD.cc
+    hipMemcpyDtoDAsync.cc
+    hipHostMalloc.cc
+    hipMemcpy_old.cc
+    hipMemcpy_derivatives.cc
+    hipMemcpyAsync.cc
+    hipMemsetFunctional.cc
+    hipMalloc.cc
+    hipMallocPitch.cc
+    hipMallocArray.cc
+    hipMalloc3D.cc
+    hipMalloc3DArray.cc
+    hipArrayCreate.cc
+    hipArray3DCreate.cc
+    hipDrvMemcpy3D.cc
+    hipDrvMemcpy3DAsync.cc
+    hipPointerGetAttribute.cc
+    hipDrvPtrGetAttributes.cc
+    hipMemPrefetchAsync.cc
+    hipMemGetInfo.cc
+    hipFree.cc
+    hipMemcpySync.cc
+    hipMemsetSync.cc
+    hipMemsetAsync.cc
+    hipMemAdvise_old.cc
+    hipMemAdvise.cc
+    hipMemRangeGetAttributes.cc
+    hipStreamAttachMemAsync.cc
+    hipMemRangeGetAttributes_old.cc
+    hipMemGetAddressRange.cc)
+
 if(HIP_PLATFORM MATCHES "amd")
-set(TEST_SRC
-    memset.cc
-    malloc.cc
-    hipMemcpy2DToArray.cc
-    hipMemcpy2DToArray_old.cc
-    hipMemcpy2DToArrayAsync.cc
-    hipMemcpy2DToArrayAsync_old.cc
-    hipMemcpy3D.cc
-    hipMemcpy3DAsync.cc
-    hipMemcpyParam2D.cc
-    hipMemcpyParam2DAsync.cc
-    hipMemcpy2D.cc
-    hipMemcpy2DAsync.cc
-    hipMemcpy2DFromArray.cc
-    hipMemcpy2DFromArray_old.cc
-    hipMemcpy2DFromArrayAsync.cc
-    hipMemcpy2DFromArrayAsync_old.cc
-    hipMemcpyAtoH.cc
-    hipMemcpyAtoH_old.cc
-    hipMemcpyHtoA.cc
-    hipMemcpyHtoA_old.cc
-    hipMemcpyAllApiNegative.cc
-    hipMemcpy_MultiThread.cc
-    hipHostRegister.cc
-    hipHostUnregister.cc
-    hipMemPtrGetInfo.cc
-    hipPointerGetAttributes.cc
-    hipHostGetFlags.cc
-    hipHostGetDevicePointer.cc
-    hipMallocManaged_MultiScenario.cc
-    hipMemsetNegative.cc
-    hipMemset.cc
-    hipMemsetAsyncMultiThread.cc
-    hipMemset3D.cc
-    hipMemset2D.cc
-    hipHostMallocTests.cc
-    hipMemset3DFunctional.cc
-    hipMemset3DRegressMultiThread.cc
-    hipMallocManagedFlagsTst.cc
-    hipMemPrefetchAsyncExtTsts.cc
-    hipMemAdviseMmap.cc
-    hipMemCoherencyTst.cc
-    hipMallocManaged.cc
-    hipMemRangeGetAttribute.cc
-    hipMemRangeGetAttribute_old.cc
-    hipMemcpyFromSymbol.cc
-    hipPtrGetAttribute.cc
-    hipMemPoolApi.cc
-    hipMemcpyPeer.cc
-    hipMemcpyPeer_old.cc
-    hipMemcpyPeerAsync.cc
-    hipMemcpyPeerAsync_old.cc
-    hipMemcpyWithStream_old.cc
-    hipMemcpyWithStream.cc
-    hipMemcpyWithStreamMultiThread.cc
-    hipMemsetAsyncAndKernel.cc
-    hipMemset2DAsyncMultiThreadAndKernel.cc
-    hipMallocConcurrency.cc
-    hipMemcpyDtoD.cc
-    hipMemcpyDtoDAsync.cc
-    hipHostMalloc.cc
-    hipMemcpy_old.cc
-    hipMemcpy_derivatives.cc
-    hipMemcpyAsync.cc
-    hipMemsetFunctional.cc
-    hipMalloc.cc
-    hipExtMallocWithFlags.cc
-    hipMallocPitch.cc
-    hipMallocArray.cc
-    hipMalloc3D.cc
-    hipMalloc3DArray.cc
-    hipArrayCreate.cc
-    hipArray3DCreate.cc
-    hipDrvMemcpy3D.cc
-    hipDrvMemcpy3DAsync.cc
-    hipPointerGetAttribute.cc
-    hipDrvPtrGetAttributes.cc
-    hipMallocMngdMultiThread.cc
-    hipMemPrefetchAsync.cc
-    hipArray.cc
-    hipMemVmm.cc
-    hipMemGetInfo.cc
-    hipFree.cc
-    hipMemcpySync.cc
-    hipMemsetSync.cc
-    hipMemsetAsync.cc
-    hipMemAdvise_old.cc
-    hipMemAdvise.cc
-    hipMemRangeGetAttributes.cc
-    hipStreamAttachMemAsync.cc
-    hipMemRangeGetAttributes_old.cc
-    hipMemGetAddressRange.cc
-)
+  set(TEST_SRC
+      ${TEST_SRC}
+      hipMemPtrGetInfo.cc
+      hipPointerGetAttributes.cc
+      hipMemCoherencyTst.cc
+      hipExtMallocWithFlags.cc
+      hipMallocMngdMultiThread.cc
+      hipArray.cc
+      hipMemVmm.cc)
 else()
-set(TEST_SRC
-    memset.cc
-    malloc.cc
-    hipMemcpy2DToArray.cc
-    hipMemcpy2DToArray_old.cc
-    hipMemcpy2DToArrayAsync.cc
-    hipMemcpy2DToArrayAsync_old.cc
-    hipMemcpy3D.cc
-    hipMemcpy3DAsync.cc
-    hipMemcpyParam2D.cc
-    hipMemcpyParam2DAsync.cc
-    hipMemcpy2D.cc
-    hipMemcpy2DAsync.cc
-    hipMemcpy2DFromArray.cc
-    hipMemcpy2DFromArray_old.cc
-    hipMemcpy2DFromArrayAsync.cc
-    hipMemcpy2DFromArrayAsync_old.cc
-    hipMemcpyAtoH.cc
-    hipMemcpyAtoH_old.cc
-    hipMemcpyHtoA.cc
-    hipMemcpyHtoA_old.cc
-    hipMemcpyAllApiNegative.cc
-    hipMemcpy_MultiThread.cc
-    hipHostRegister.cc
-    hipHostUnregister.cc
-    hipHostGetFlags.cc
-    hipHostGetDevicePointer.cc
-    hipMallocManaged_MultiScenario.cc
-    hipMemsetNegative.cc
-    hipMemset.cc
-    hipMemsetAsyncMultiThread.cc
-    hipMemset3D.cc
-    hipMemset2D.cc
-    hipHostMallocTests.cc
-    hipMemset3DFunctional.cc
-    hipMemset3DRegressMultiThread.cc
-    hipMallocManagedFlagsTst.cc
-    hipMemPrefetchAsyncExtTsts.cc
-    hipMemAdviseMmap.cc
-    hipMallocManaged.cc
-    hipMemRangeGetAttribute.cc
-    hipMemRangeGetAttribute_old.cc
-    hipMemcpyFromSymbol.cc
-    hipPtrGetAttribute.cc
-    hipMemPoolApi.cc
-    hipMemcpyPeer.cc
-    hipMemcpyPeer_old.cc
-    hipMemcpyPeerAsync.cc
-    hipMemcpyPeerAsync_old.cc
-    hipMemcpyWithStream_old.cc
-    hipMemcpyWithStream.cc
-    hipMemcpyWithStreamMultiThread.cc
-    hipMemsetAsyncAndKernel.cc
-    hipMemset2DAsyncMultiThreadAndKernel.cc
-    hipMallocConcurrency.cc
-    hipMemcpyDtoD.cc
-    hipMemcpyDtoDAsync.cc
-    hipHostMalloc.cc
-    hipMemcpy_old.cc
-    hipMemcpy_derivatives.cc
-    hipMemcpyAsync.cc
-    hipMemsetFunctional.cc
-    hipMalloc.cc
-    hipMallocPitch.cc
-    hipMallocArray.cc
-    hipMalloc3D.cc
-    hipMalloc3DArray.cc
-    hipArrayCreate.cc
-    hipArray3DCreate.cc
-    hipDrvMemcpy3D.cc
-    hipDrvMemcpy3DAsync.cc
-    hipPointerGetAttribute.cc
-    hipDrvPtrGetAttributes.cc
-    hipMemPrefetchAsync.cc
-    hipMemGetInfo.cc
-    hipFree.cc
-    hipMemcpySync.cc
-    hipMemsetSync.cc
-    hipMemsetAsync.cc
-    hipMemAdvise_old.cc
-    hipMemAdvise.cc
-    hipMemRangeGetAttributes.cc
-    hipMemRangeGetAttributes_old.cc
-    hipGetSymbolSizeAddress.cc
-    hipStreamAttachMemAsync.cc
-    hipMemGetAddressRange.cc
-)
+  set(TEST_SRC ${TEST_SRC} hipGetSymbolSizeAddress.cc)
 endif()

-# skipped due to os related code in tests
-# need to work on them when all the tests are enabled
+# skipped due to os related code in tests need to work on them when all the
+# tests are enabled
 if(UNIX)
-    set(TEST_SRC ${TEST_SRC}
-                 hipHmmOvrSubscriptionTst.cc
-                 hipMemoryAllocateCoherent.cc)
+  set(TEST_SRC ${TEST_SRC} hipHmmOvrSubscriptionTst.cc
+               hipMemoryAllocateCoherent.cc)
 endif()

 hip_add_exe_to_target(NAME MemoryTest
-                      TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests
-                      COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
+  TEST_SRC ${TEST_SRC}
+  TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
@@ -122,7 +122,8 @@ void checkPointer(const SuperPointerAttribute& ref, int major, int minor, void*
 // we do this in the testMultiThreaded_1 test.
 void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
  Nbytes = N * sizeof(char);
-  printf("clusterAllocs numAllocs=%d size=%lu..%lu\n", numAllocs, minSize, maxSize);
+  printf("clusterAllocs numAllocs=%d size=%lu..%lu\n",
+         numAllocs, static_cast<unsigned long>(minSize), static_cast<unsigned long>(maxSize));
  const int Max_Devices = 256;
  std::vector<SuperPointerAttribute> reference(numAllocs);

@@ -174,7 +175,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize) {
        "  device#%d: hipMemGetInfo: "
        "free=%zu (%4.2fMB) totalDevice=%lu (%4.2fMB) total=%zu "
        "(%4.2fMB)\n",
-        i, free, (free / 1024.0 / 1024.0), totalDeviceAllocated[i],
+        i, free, (free / 1024.0 / 1024.0), static_cast<unsigned long>(totalDeviceAllocated[i]),
        (totalDeviceAllocated[i]) / 1024.0 / 1024.0, total, (total / 1024.0 / 1024.0));
    REQUIRE(free + totalDeviceAllocated[i] <= total);
  }
@@ -221,7 +222,7 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") {
  size_t free, total;
  HIP_CHECK(hipMemGetInfo(&free, &total));
  printf("hipMemGetInfo: free=%zu (%4.2f) Nbytes=%lu total=%zu (%4.2f)\n", free,
-         (free / 1024.0 / 1024.0), Nbytes, total, (total / 1024.0 / 1024.0));
+         (free / 1024.0 / 1024.0), static_cast<unsigned long>(Nbytes), total, (total / 1024.0 / 1024.0));
  REQUIRE(free + Nbytes <= total);


@@ -50,7 +50,7 @@ TEST_CASE("Unit_hipPtrGetAttribute_Simple") {
  size_t free, total;
  HIP_CHECK(hipMemGetInfo(&free, &total));
  printf("hipMemGetInfo: free=%zu (%4.2f) Nbytes=%lu total=%zu (%4.2f)\n", free,
-         (free / 1024.0 / 1024.0), Nbytes, total,
+         (free / 1024.0 / 1024.0), static_cast<unsigned long>(Nbytes), total,
         (total / 1024.0 / 1024.0));
  REQUIRE(free + Nbytes <= total);

@@ -1,6 +1,5 @@
 set(COMMON_SHARED_SRC streamCommon.cc)

-if(HIP_PLATFORM MATCHES "amd")
 set(TEST_SRC
    hipStreamCreate.cc
    hipStreamGetFlags.cc
@@ -10,43 +9,30 @@ set(TEST_SRC
    hipStreamCreateWithFlags.cc
    hipStreamCreateWithPriority.cc
    hipStreamDestroy.cc
-    hipStreamGetCUMask.cc
    hipAPIStreamDisable.cc
    hipStreamValue.cc
-    hipStreamWithCUMask.cc
-    hipStreamACb_MultiThread.cc
-    hipStreamSynchronize.cc
-    hipStreamQuery.cc
-    hipStreamWaitEvent.cc
-    hipDeviceGetStreamPriorityRange.cc
-    hipStreamACb_StrmSyncTiming.cc
-    )
-else()
-set(TEST_SRC
-    hipStreamCreate.cc
-    hipStreamGetFlags.cc
-    hipStreamGetPriority.cc
-    hipMultiStream.cc
-    hipStreamACb_MultiThread.cc
-    hipStreamAddCallback.cc
-    hipStreamCreateWithFlags.cc
-    hipStreamCreateWithPriority.cc
-    hipStreamDestroy.cc
-    hipAPIStreamDisable.cc
-    # hipStreamAttachMemAsync_old.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.
-    hipStreamValue.cc
    hipStreamSynchronize.cc
    hipStreamQuery.cc
    hipDeviceGetStreamPriorityRange.cc
    hipStreamACb_StrmSyncTiming.cc
-    )
+)

-    # set_source_files_properties(hipStreamAttachMemAsync_old.cc PROPERTIES COMPILE_FLAGS -std=c++17)
+if(HIP_PLATFORM MATCHES "amd")
+  set(TEST_SRC ${TEST_SRC} hipStreamGetCUMask.cc hipStreamWithCUMask.cc
+               hipStreamACb_MultiThread.cc hipStreamWaitEvent.cc)
+else()
+  set(TEST_SRC
+      ${TEST_SRC}
+      # hipStreamAttachMemAsync_old.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.
+      hipStreamACb_MultiThread.cc)
+
+  # set_source_files_properties(hipStreamAttachMemAsync_old.cc PROPERTIES
+  # COMPILE_FLAGS -std=c++17)
 endif()

-hip_add_exe_to_target(NAME StreamTest
-                      TEST_SRC ${TEST_SRC}
-                      TEST_TARGET_NAME build_tests
-                      COMPILE_OPTIONS -std=c++17
-                      COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
+hip_add_exe_to_target(
+  NAME StreamTest
+  TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests
+  COMPILE_OPTIONS -std=c++17 COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
@@ -5,8 +5,17 @@ set(TEST_SRC
  hipStreamPerThread_MultiThread.cc
  hipStreamPerThread_DeviceReset.cc
  hipStreamPerThrdTsts.cc
+  hipStreamPerThrdCompilerOptn.cc
 )

+if(HIP_PLATFORM MATCHES "amd")
+        set_source_files_properties(hipStreamPerThrdCompilerOptn.cc PROPERTIES COMPILE_OPTIONS "-fgpu-default-stream=per-thread")
+endif()
+
+if(HIP_PLATFORM MATCHES "nvidia")
+        set_source_files_properties(hipStreamPerThrdCompilerOptn.cc PROPERTIES COMPILE_OPTIONS "--default-stream=per-thread")
+endif()
+
 hip_add_exe_to_target(NAME StreamPerThreadTest
                      TEST_SRC ${TEST_SRC}
                      TEST_TARGET_NAME build_tests)
@@ -0,0 +1,793 @@
+/*
+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 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.
+*/
+
+/* This file will be compiled using the compiler option given below.
+   The functions will be calculating the time taken to execute under
+   the influence of the compiler flag:
+   -fgpu-default-stream=per-thread
+*/
+
+#include <ctime>
+#include <hip_test_common.hh>
+#include "hip/hip_cooperative_groups.h"
+
+namespace cg = cooperative_groups;
+
+namespace DefltStrmPT {
+  int64_t N = 1024 * 1024 * 100;
+  int64_t Sz = N * sizeof(int64_t);
+  int64_t *DevA, *HstA, *HstRes;
+  int64_t OneMB = 1024 * 1024;
+  int64_t OneMBSz = OneMB * sizeof(int64_t);
+  hipStream_t Strm;
+  int clockrate, CONST = 123;
+  size_t numH = 1024, numW = 1024;
+  size_t pitch_A, width = numW * sizeof(int64_t);
+  size_t sizeElements = width * numH;
+  size_t elements = numW * numH;
+}  // namespace DefltStrmPT
+
+__device__ int64_t globalInDStrmPT[1024 * 1024];
+__device__ int SigComplte = 0;
+
+// Kernel codes
+__global__ void DefltStrmPT_Square(int64_t *C_d, int64_t N) {
+    int64_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+    int64_t stride = hipBlockDim_x * hipGridDim_x;
+
+    for (int64_t i = offset; i < N; i += stride) {
+        C_d[i] = C_d[i] * C_d[i];
+    }
+}
+
+__global__ void Wait_Kernel3(int clockrate, uint64_t WaitSecs,
+                             int PassSignal = 0) {
+  uint64_t num_cycles = WaitSecs * clockrate * 1000;
+  uint64_t start = clock64(), cycles = 0;
+  while (cycles < num_cycles) {
+    cycles = clock64() - start;
+  }
+  if (PassSignal) {
+    SigComplte = 1;
+  }
+}
+
+__global__ void DefltStrmPT_Test_gws(uint* buf, uint bufSize,
+                                       int64_t* tmpBuf, int64_t* result) {
+    extern __shared__ int64_t tmp[];
+    uint offset = blockIdx.x * blockDim.x + threadIdx.x;
+    uint stride = gridDim.x  * blockDim.x;
+    cg::grid_group gg = cg::this_grid();
+
+    int64_t sum = 0;
+    for (uint i = offset; i < bufSize; i += stride) {
+        sum += buf[i];
+    }
+    tmp[threadIdx.x] = sum;
+
+    __syncthreads();
+
+    if (threadIdx.x == 0) {
+        sum = 0;
+        for (uint i = 0; i < blockDim.x; i++) {
+            sum += tmp[i];
+        }
+        tmpBuf[blockIdx.x] = sum;
+    }
+
+    gg.sync();
+
+    if (offset == 0) {
+        for (uint i = 1; i < gridDim.x; ++i) {
+          sum += tmpBuf[i];
+       }
+       *result = sum;
+    }
+}
+
+float DefaultPT2_Memcpy_MemSet(int CpyAsync, int MemSetAsync) {
+  bool IfTstPassed = true;
+  DefltStrmPT::HstA = reinterpret_cast<int64_t*> (malloc(DefltStrmPT::Sz));
+  DefltStrmPT::HstRes = reinterpret_cast<int64_t*> (malloc(DefltStrmPT::Sz));
+  HIP_CHECK(hipDeviceGetAttribute(&(DefltStrmPT::clockrate),
+            hipDeviceAttributeMemoryClockRate, 0));
+  HIP_CHECK(hipMalloc(&(DefltStrmPT::DevA), DefltStrmPT::Sz));
+  for (int64_t i = 0; i < DefltStrmPT::N; ++i) {
+    DefltStrmPT::HstA[i] = DefltStrmPT::CONST;
+  }
+  HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm)));
+  if (CpyAsync) {
+    HIP_CHECK(hipMemcpyAsync(DefltStrmPT::DevA, DefltStrmPT::HstA,
+                             DefltStrmPT::Sz, hipMemcpyHostToDevice,
+                             DefltStrmPT::Strm));
+    HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm));
+  } else {
+    HIP_CHECK(hipMemcpy(DefltStrmPT::DevA, DefltStrmPT::HstA,
+                        DefltStrmPT::Sz, hipMemcpyHostToDevice));
+  }
+  DefltStrmPT_Square<<<(DefltStrmPT::N/256 + 1), 256, 0, DefltStrmPT::Strm>>>
+                      (DefltStrmPT::DevA, DefltStrmPT::N);
+  HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm));
+  HIP_CHECK(hipMemcpy(DefltStrmPT::HstRes, DefltStrmPT::DevA,
+                      DefltStrmPT::Sz, hipMemcpyDeviceToHost));
+  // Verifying the result
+  for (int64_t i = 0; i < DefltStrmPT::N; ++i) {
+    if (DefltStrmPT::HstRes[i] !=
+       (DefltStrmPT::HstA[i] * DefltStrmPT::HstA[i])) {
+      IfTstPassed = false;
+    }
+  }
+  if (MemSetAsync) {
+    HIP_CHECK(hipMemsetAsync(DefltStrmPT::DevA, 0, DefltStrmPT::Sz,
+                             DefltStrmPT::Strm));
+    HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm));
+  } else {
+    HIP_CHECK(hipMemset(DefltStrmPT::DevA, 0,
+                        DefltStrmPT::Sz));
+  }
+  // Copying the device memory to host to check if Memset is successful
+  HIP_CHECK(hipMemcpy(DefltStrmPT::HstA, DefltStrmPT::DevA,
+                      DefltStrmPT::Sz, hipMemcpyDeviceToHost));
+  // verifying if memset was successful
+  for (int64_t i = 0; i < DefltStrmPT::N; ++i) {
+    if (DefltStrmPT::HstA[i] != 0) {
+      IfTstPassed = false;
+    }
+  }
+  HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+  HIP_CHECK(hipFree(DefltStrmPT::DevA));
+  free(DefltStrmPT::HstA);
+  free(DefltStrmPT::HstRes);
+  return IfTstPassed;
+}
+
+void DefaultPT2_Memset2D(int Async) {
+  constexpr int memsetval = 0x24;
+  constexpr size_t numH = 256;
+  constexpr size_t numW = 256;
+  size_t pitch_A;
+  size_t width = numW * sizeof(char);
+  size_t sizeElements = width * numH;
+  size_t elements = numW * numH;
+  char *A_d, *A_h;
+
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&A_d), &pitch_A, width,
+                          numH));
+  A_h = reinterpret_cast<char*>(malloc(sizeElements));
+  REQUIRE(A_h != nullptr);
+
+  for (size_t i = 0; i < elements; i++) {
+    A_h[i] = 1;
+  }
+
+  if (Async) {
+    hipStream_t Strm;
+    HIP_CHECK(hipStreamCreate(&Strm));
+    HIP_CHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, Strm));
+    HIP_CHECK(hipStreamSynchronize(Strm));
+    HIP_CHECK(hipStreamDestroy(Strm));
+  } else {
+    HIP_CHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH));
+  }
+  HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH,
+                       hipMemcpyDeviceToHost));
+
+  for (size_t i = 0; i < elements; i++) {
+    if (A_h[i] != memsetval) {
+      INFO("Memset2D mismatch at index:" << i << " computed:"
+                                     << A_h[i] << " memsetval:" << memsetval);
+      REQUIRE(false);
+    }
+  }
+
+  HIP_CHECK(hipFree(A_d));
+  free(A_h);
+}
+
+
+void PerThrdDefltStrm_Memset3D(int Async) {
+  constexpr int memsetval = 0x22;
+  constexpr size_t numH = 256;
+  constexpr size_t numW = 256;
+  constexpr size_t depth = 10;
+  size_t width = numW * sizeof(char);
+  size_t sizeElements = width * numH * depth;
+  size_t elements = numW * numH * depth;
+  char *A_h;
+
+  hipExtent extent = make_hipExtent(width, numH, depth);
+  hipPitchedPtr devPitchedPtr;
+
+  HIP_CHECK(hipMalloc3D(&devPitchedPtr, extent));
+  A_h = reinterpret_cast<char *>(malloc(sizeElements));
+  REQUIRE(A_h != nullptr);
+
+  for (size_t i = 0; i < elements; i++) {
+      A_h[i] = 1;
+  }
+
+  if (Async) {
+    hipStream_t Strm;
+    HIP_CHECK(hipStreamCreate(&Strm));
+    HIP_CHECK(hipMemset3DAsync(devPitchedPtr, memsetval, extent, Strm));
+    HIP_CHECK(hipStreamSynchronize(Strm));
+    HIP_CHECK(hipStreamDestroy(Strm));
+  } else {
+    HIP_CHECK(hipMemset3D(devPitchedPtr, memsetval, extent));
+  }
+  hipMemcpy3DParms myparms{};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH);
+  myparms.srcPtr = devPitchedPtr;
+  myparms.extent = extent;
+#if HT_NVIDIA
+  myparms.kind = hipMemcpyKindToCudaMemcpyKind(hipMemcpyDeviceToHost);
+#else
+  myparms.kind = hipMemcpyDeviceToHost;
+#endif
+  HIP_CHECK(hipMemcpy3D(&myparms));
+
+  for (size_t i = 0; i < elements; i++) {
+      if (A_h[i] != memsetval) {
+        INFO("Memset3D mismatch at index:" << i << " computed:"
+                                      << A_h[i] << " memsetval:" << memsetval);
+        REQUIRE(false);
+      }
+  }
+  HIP_CHECK(hipFree(devPitchedPtr.ptr));
+  free(A_h);
+}
+
+
+void DefaultPT2_StrmQuery() {
+  HIP_CHECK(hipDeviceGetAttribute(&(DefltStrmPT::clockrate),
+          hipDeviceAttributeMemoryClockRate, 0));
+  HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm)));
+  // StreamQuery with null stream
+  Wait_Kernel3<<<1, 1>>>(DefltStrmPT::clockrate, 3);
+  REQUIRE((hipErrorNotReady == hipStreamQuery(0)));
+  // StreamQuery with user created stream
+  Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 3);
+  REQUIRE((hipErrorNotReady == hipStreamQuery(DefltStrmPT::Strm)));
+  HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+}
+
+
+void DefaultPT2_StreamSync() {
+  HIP_CHECK(hipDeviceGetAttribute(&(DefltStrmPT::clockrate),
+          hipDeviceAttributeMemoryClockRate, 0));
+  HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm)));
+  // Calling hipStreamSync on user created stream object
+  Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 1);
+  HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm));
+  // Calling hipStreamSync on null stream
+  Wait_Kernel3<<<1, 1>>>(DefltStrmPT::clockrate, 1);
+  HIP_CHECK(hipStreamSynchronize(0));
+  HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+}
+
+
+void DefaultPT2_StrmWaitEvent() {
+  hipEvent_t evt;
+  hipStream_t Strm1;
+  HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm)));
+  HIP_CHECK(hipStreamCreate(&Strm1));
+  HIP_CHECK(hipEventCreate(&evt));
+  Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 3, 1);
+  HIP_CHECK(hipEventRecord(evt, DefltStrmPT::Strm));
+  HIP_CHECK(hipStreamWaitEvent(Strm1, evt, 0));
+  Wait_Kernel3<<<1, 1, 0, Strm1>>>(DefltStrmPT::clockrate, 1);
+  // By the time control reaches the below point SigComplte is expected
+  // to be still zero
+  if (SigComplte) {
+    REQUIRE(false);
+  }
+  HIP_CHECK(hipStreamSynchronize(Strm1));
+  HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+  HIP_CHECK(hipStreamDestroy(Strm1));
+  HIP_CHECK(hipEventDestroy(evt));
+}
+
+void DefaultPT2_EvtQuery() {
+  hipEvent_t evt, evt1;
+  hipError_t err;
+  HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm)));
+  HIP_CHECK(hipEventCreate(&evt));
+  HIP_CHECK(hipEventCreate(&evt1));
+  Wait_Kernel3<<<1, 1, 0, DefltStrmPT::Strm>>>(DefltStrmPT::clockrate, 3);
+  HIP_CHECK(hipEventRecord(evt, DefltStrmPT::Strm));
+  err = hipEventQuery(evt);
+  if (err != hipErrorNotReady) {
+    REQUIRE(false);
+  }
+  // Testing for Null or default stream
+  HIP_CHECK(hipEventRecord(evt1, 0));
+  std::chrono::time_point start = std::chrono::steady_clock::now();
+  int Got_hipSuccess = 0;  // 0 for no, 1 for yes
+  while (true) {
+    err = hipEventQuery(evt1);
+    if (err == hipSuccess) {
+      Got_hipSuccess = 1;
+      break;
+    }
+    if (std::chrono::steady_clock::now() - start > std::chrono::seconds(60)) {
+      break;
+    }
+  }
+  if (!Got_hipSuccess) {
+    REQUIRE(false);
+  }
+  HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+  HIP_CHECK(hipEventDestroy(evt));
+  HIP_CHECK(hipEventDestroy(evt1));
+}
+
+
+void Default_LaunchKernel(int NullStrm) {
+  DefltStrmPT::N =  DefltStrmPT::N/4;
+  DefltStrmPT::Sz = DefltStrmPT::N * sizeof(int64_t);
+  DefltStrmPT::HstA = reinterpret_cast<int64_t*> (malloc(DefltStrmPT::Sz));
+  HIP_CHECK(hipMalloc(&(DefltStrmPT::DevA), DefltStrmPT::Sz));
+  for (int64_t i = 0; i < DefltStrmPT::N; ++i) {
+    DefltStrmPT::HstA[i] = DefltStrmPT::CONST;
+  }
+  HIP_CHECK(hipMemcpy(DefltStrmPT::DevA, DefltStrmPT::HstA, DefltStrmPT::Sz,
+            hipMemcpyHostToDevice));
+  HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm)));
+  unsigned ThrdsPerBlk = 32;
+  unsigned Blocks = ((DefltStrmPT::N + ThrdsPerBlk - 1)/ThrdsPerBlk);
+  void *Args[] = {&(DefltStrmPT::DevA), &(DefltStrmPT::N)};
+  // launch Kernel
+  if (NullStrm) {
+    HIP_CHECK(hipLaunchKernel((const void*)DefltStrmPT_Square,
+             dim3(Blocks, 1, 1), dim3(ThrdsPerBlk, 1, 1), Args, 0,
+             0));
+    HIP_CHECK(hipStreamSynchronize(0));
+  } else {
+    HIP_CHECK(hipLaunchKernel((const void*)DefltStrmPT_Square,
+           dim3(Blocks, 1, 1), dim3(ThrdsPerBlk, 1, 1), Args, 0,
+           DefltStrmPT::Strm));
+    HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm));
+  }
+  HIP_CHECK(hipMemcpy(DefltStrmPT::HstA, DefltStrmPT::DevA, DefltStrmPT::Sz,
+            hipMemcpyDeviceToHost));
+  for (int64_t i = 0; i < DefltStrmPT::N; ++i) {
+    if (DefltStrmPT::HstA[i] != (DefltStrmPT::CONST * DefltStrmPT::CONST)) {
+      REQUIRE(false);
+    }
+  }
+  HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+  HIP_CHECK(hipFree(DefltStrmPT::DevA));
+  free(DefltStrmPT::HstA);
+}
+
+
+void DefaultPT2_LaunchCooperativeKernel(int NullStrm) {
+  bool IfTestPassed = true;
+  uint32_t *dA;
+  int64_t *dB, *dC;
+  uint32_t BufferSizeInDwords = 448 * 1024 * 1024;
+  uint32_t* init = new uint32_t[BufferSizeInDwords];
+  for (uint32_t i = 0; i < BufferSizeInDwords; ++i) {
+    init[i] = i;
+  }
+  size_t SIZE = BufferSizeInDwords * sizeof(uint);
+  HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm)));
+  hipDeviceProp_t deviceProp;
+  HIP_CHECK(hipGetDeviceProperties(&deviceProp, 0));
+  HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dA), SIZE));
+  HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(&dC), sizeof(int64_t)));
+  HIPCHECK(hipMemcpy(dA, init, SIZE, hipMemcpyHostToDevice));
+
+  dim3 dimBlock = dim3(1);
+  dim3 dimGrid  = dim3(1);
+  int numBlocks = 0;
+  dimBlock.x = 32;
+  // Calculate the device occupancy to know how many blocks can be run
+  //  concurrently
+  HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
+                  DefltStrmPT_Test_gws,
+                  dimBlock.x * dimBlock.y * dimBlock.z,
+                  dimBlock.x * sizeof(int64_t)));
+  dimGrid.x = deviceProp.multiProcessorCount * std::min(numBlocks, 32);
+  HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dB),
+                     dimGrid.x * sizeof(int64_t)));
+
+  void *params[4];
+  params[0] = reinterpret_cast<void*>(&dA);
+  params[1] = reinterpret_cast<void*>(&BufferSizeInDwords);
+  params[2] = reinterpret_cast<void*>(&dB);
+  params[3] = reinterpret_cast<void*>(&dC);
+  if (NullStrm) {
+    HIPCHECK(hipLaunchCooperativeKernel(
+             reinterpret_cast<void*>(DefltStrmPT_Test_gws),
+             dimGrid, dimBlock, params, dimBlock.x * sizeof(int64_t), 0));
+    HIP_CHECK(hipStreamSynchronize(0));
+  } else {
+    HIPCHECK(hipLaunchCooperativeKernel(
+             reinterpret_cast<void*>(DefltStrmPT_Test_gws),
+             dimGrid, dimBlock, params, dimBlock.x * sizeof(int64_t),
+             DefltStrmPT::Strm));
+    HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm));
+  }
+  HIPCHECK(hipMemcpy(init, dC, sizeof(int64_t), hipMemcpyDeviceToHost));
+
+  if (*dC != (((int64_t)(BufferSizeInDwords) * (BufferSizeInDwords - 1)) / 2)) {
+    std::cout << "Data validation failed for grid size = " << dimGrid.x <<
+    " and block size = " << dimBlock.x << "\n";
+        std::cout << "Test failed! \n";
+        IfTestPassed = false;
+  }
+  HIPCHECK(hipStreamDestroy(DefltStrmPT::Strm));
+  HIPCHECK(hipHostFree(dC));
+  HIPCHECK(hipFree(dB));
+  HIPCHECK(hipFree(dA));
+  delete [] init;
+  REQUIRE(IfTestPassed);
+}
+
+
+void DefaultPT2_StrmGetFlag() {
+  HIP_CHECK(hipStreamCreateWithFlags(&(DefltStrmPT::Strm), hipStreamDefault));
+  unsigned int flag = 9999;
+  HIP_CHECK(hipStreamGetFlags(DefltStrmPT::Strm, &flag));
+  if (flag != 0) {
+    INFO("Expected flag value: 0, Received flag value: %u\n" << flag);
+    REQUIRE(false);
+  }
+  HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+  HIP_CHECK(hipStreamCreateWithFlags(&(DefltStrmPT::Strm),
+                                     hipStreamNonBlocking));
+  flag = 9999;
+  HIP_CHECK(hipStreamGetFlags(DefltStrmPT::Strm, &flag));
+  if (flag != 1) {
+    INFO("Expected flag value: 1, Received flag value: %u\n" << flag);
+    REQUIRE(false);
+  }
+  HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+}
+
+void DefaultPT2_StrmGetPriority() {
+  int low, high, ObsrvdPriority;
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&low, &high));
+  INFO("Lowest possible priority: %d\n" << low);
+  INFO("Highest possible priority: %d\n" << high);
+  INFO("Creating streams with flag hipStreamNonBlocking\n");
+  // hipStrmFlg = 0 = hipStreamDefault
+  // hipStrmFlg = 1 = hipStreamNonBlocking
+  for (int hipStrmFlg = 0; hipStrmFlg < 2; ++hipStrmFlg) {
+    for (int Priority = low; Priority <= high; ++Priority) {
+      if (hipStrmFlg == 0) {
+        HIP_CHECK(hipStreamCreateWithPriority(&(DefltStrmPT::Strm),
+                  hipStreamDefault, Priority));
+      } else {
+        HIP_CHECK(hipStreamCreateWithPriority(&(DefltStrmPT::Strm),
+                  hipStreamNonBlocking, Priority));
+      }
+      HIP_CHECK(hipStreamGetPriority(DefltStrmPT::Strm, &ObsrvdPriority));
+      if (ObsrvdPriority != Priority) {
+        INFO("Expected priority: %d" << Priority << " Observed Priority: %d\n"
+             << ObsrvdPriority);
+        INFO("Test Failed!\n\n");
+        REQUIRE(false);
+      }
+      HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+    }
+  }
+  INFO("Checking priority on null stream!!\n");
+  HIP_CHECK(hipStreamGetPriority(0, &ObsrvdPriority));
+  if (ObsrvdPriority != 0) {
+        INFO("Expected priority: 0, Observed Priority: %d\n"
+             << ObsrvdPriority);
+        INFO("Test Failed!\n\n");
+        REQUIRE(false);
+  }
+}
+
+
+void DefaultPT2_hipMemcpyFromSymbol() {
+  int64_t *Hst = nullptr;
+  HIP_CHECK(hipHostMalloc(&(Hst), DefltStrmPT::OneMBSz));
+  HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm)));
+  for (int i = 0; i < DefltStrmPT::OneMB; ++i) {
+    Hst[i] = DefltStrmPT::CONST;
+  }
+  HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalInDStrmPT), Hst,
+            DefltStrmPT::OneMBSz, 0, hipMemcpyHostToDevice));
+  for (int i = 0; i < DefltStrmPT::OneMB; ++i) {
+    Hst[i] = 0;
+  }
+  HIP_CHECK(hipMemcpyFromSymbol(Hst, HIP_SYMBOL(globalInDStrmPT),
+            DefltStrmPT::OneMBSz, 0, hipMemcpyDeviceToHost));
+  for (int i = 0; i < DefltStrmPT::OneMB; ++i) {
+    if (Hst[i] != DefltStrmPT::CONST) {
+      REQUIRE(false);
+    }
+  }
+  HIP_CHECK(hipHostFree(Hst));
+  HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+}
+
+void DefaultPT2_hipMemcpy2D(int Async) {
+  DefltStrmPT::numH = 1024;
+  DefltStrmPT::numW = 1024;
+  DefltStrmPT::width = DefltStrmPT::numW * sizeof(int64_t);
+  HIP_CHECK(hipHostMalloc(&(DefltStrmPT::HstA),
+            (DefltStrmPT::numH * DefltStrmPT::numW * sizeof(int64_t))));
+  HIP_CHECK(hipHostMalloc(&(DefltStrmPT::HstRes),
+            (DefltStrmPT::numH * DefltStrmPT::numW * sizeof(int64_t))));
+  DefltStrmPT::width = DefltStrmPT::numW * sizeof(int64_t);
+  for (size_t row = 0; row < DefltStrmPT::numH; ++row) {
+    for (size_t column = 0; column < DefltStrmPT::numW; ++column) {
+      DefltStrmPT::HstA[(row * DefltStrmPT::numW) + column] =
+                   DefltStrmPT::CONST;
+      DefltStrmPT::HstRes[(row * DefltStrmPT::numW) + column] = 0;
+    }
+  }
+  HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&(DefltStrmPT::DevA)),
+            &(DefltStrmPT::pitch_A), DefltStrmPT::width, DefltStrmPT::numH));
+  if (Async) {
+    HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm)));
+    HIP_CHECK(hipMemcpy2DAsync(DefltStrmPT::DevA, DefltStrmPT::pitch_A,
+            DefltStrmPT::HstA, DefltStrmPT::numW*sizeof(int64_t),
+            DefltStrmPT::numW*sizeof(int64_t), DefltStrmPT::numH,
+            hipMemcpyHostToDevice, DefltStrmPT::Strm));
+    HIP_CHECK(hipMemcpy2DAsync(DefltStrmPT::HstRes,
+            DefltStrmPT::numW*sizeof(int64_t),
+            DefltStrmPT::DevA, DefltStrmPT::pitch_A,
+            DefltStrmPT::numW*sizeof(int64_t), DefltStrmPT::numH,
+            hipMemcpyDeviceToHost, DefltStrmPT::Strm));
+    HIP_CHECK(hipStreamSynchronize(DefltStrmPT::Strm));
+    HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+  } else {
+    HIP_CHECK(hipMemcpy2D(DefltStrmPT::DevA, DefltStrmPT::pitch_A,
+              DefltStrmPT::HstA, DefltStrmPT::numW*sizeof(int64_t),
+              DefltStrmPT::numW*sizeof(int64_t), DefltStrmPT::numH,
+              hipMemcpyHostToDevice));
+    HIP_CHECK(hipMemcpy2D(DefltStrmPT::HstRes,
+              DefltStrmPT::numW*sizeof(int64_t), DefltStrmPT::DevA,
+              DefltStrmPT::pitch_A, DefltStrmPT::numW*sizeof(int64_t),
+              DefltStrmPT::numH, hipMemcpyDeviceToHost));
+  }
+  for (size_t row = 0; row < DefltStrmPT::numH; ++row) {
+    for (size_t column = 0; column < DefltStrmPT::numW; ++column) {
+      if (DefltStrmPT::HstRes[(row * DefltStrmPT::numW) + column]
+          != DefltStrmPT::CONST) {
+        REQUIRE(false);
+      }
+    }
+  }
+  HIP_CHECK(hipFree(DefltStrmPT::DevA));
+  HIP_CHECK(hipHostFree(DefltStrmPT::HstA));
+}
+
+
+void DefaultPT2_hipMemcpy2DToArray() {
+  hipArray *Dptr = nullptr;
+  float *Hptr = nullptr, *HRes = nullptr;
+  DefltStrmPT::numH = 1024;
+  DefltStrmPT::numW = 1024;
+  DefltStrmPT::width = DefltStrmPT::numW * sizeof(float);
+  Hptr = new float[DefltStrmPT::width * DefltStrmPT::numH];
+  HRes = new float[DefltStrmPT::width * DefltStrmPT::numH];
+  for (size_t i = 0; i < DefltStrmPT::width * DefltStrmPT::numH; ++i) {
+    Hptr[i] = DefltStrmPT::CONST;
+  }
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
+  HIP_CHECK(hipMallocArray(&(Dptr), &desc, DefltStrmPT::numW,
+                           DefltStrmPT::numH, hipArrayDefault));
+  HIP_CHECK(hipMemcpy2DToArray(Dptr, 0, 0, Hptr, DefltStrmPT::width,
+                             DefltStrmPT::width,  DefltStrmPT::numH,
+                             hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy2DFromArray(HRes, DefltStrmPT::width, Dptr, 0, 0,
+            DefltStrmPT::width, DefltStrmPT::numH, hipMemcpyDeviceToHost));
+  // verifying the result
+  for (size_t i = 0; i < DefltStrmPT::numW * DefltStrmPT::numH; ++i) {
+    if (HRes[i] != DefltStrmPT::CONST) {
+      REQUIRE(false);
+    }
+  }
+  delete[] Hptr;
+  delete[] HRes;
+  HIP_CHECK(hipFreeArray(Dptr));
+}
+
+
+float DefaultPT2_hipMemcpy2DFromArray() {
+  hipArray *Dptr = nullptr;
+  float *Hptr_A = nullptr, *Hptr_B = nullptr;
+  DefltStrmPT::numH = 1024;
+  DefltStrmPT::numW = 1024;
+  DefltStrmPT::width = DefltStrmPT::numW * sizeof(float);
+  HIP_CHECK(hipDeviceGetAttribute(&(DefltStrmPT::clockrate),
+          hipDeviceAttributeMemoryClockRate, 0));
+  HIP_CHECK(hipHostMalloc(&(Hptr_A),
+            (DefltStrmPT::width * DefltStrmPT::numH * sizeof(float))));
+  HIP_CHECK(hipHostMalloc(&(Hptr_B),
+            (DefltStrmPT::width * DefltStrmPT::numH * sizeof(float))));
+  for (size_t i = 0; i < (DefltStrmPT::width * DefltStrmPT::numH); ++i) {
+    Hptr_A[i] = DefltStrmPT::CONST;
+  }
+  hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
+  HIP_CHECK(hipMallocArray(&(Dptr), &desc, DefltStrmPT::numW,
+                           DefltStrmPT::numH, hipArrayDefault));
+  HIP_CHECK(hipStreamCreate(&(DefltStrmPT::Strm)));
+  HIP_CHECK(hipMemcpy2DToArray(Dptr, 0, 0, Hptr_A, DefltStrmPT::width,
+                             DefltStrmPT::width,  DefltStrmPT::numH,
+                             hipMemcpyHostToDevice));
+    Wait_Kernel3 <<< 1, 1, 0, DefltStrmPT::Strm >>> (DefltStrmPT::clockrate,
+                                                      1);
+    HIP_CHECK(hipMemcpy2DFromArray(Hptr_B, DefltStrmPT::width, Dptr, 0, 0,
+              DefltStrmPT::width, DefltStrmPT::numH, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipStreamDestroy(DefltStrmPT::Strm));
+  HIP_CHECK(hipFreeArray(Dptr));
+  HIP_CHECK(hipHostFree(Hptr_A));
+  HIP_CHECK(hipHostFree(Hptr_B));
+  return true;
+}
+
+void DefaultPT2_hipMemcpy3D() {
+  int width = 8, height = 8, depth = 8;
+  unsigned int size = width * height * depth * sizeof(float);
+  float* Hptr = reinterpret_cast<float*>(malloc(size));
+  float* HRes = reinterpret_cast<float*>(malloc(size));
+  memset(Hptr, 0, size);
+  memset(HRes, 0, size);
+  hipExtent extent = make_hipExtent(width, height, depth);
+
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        Hptr[i*width*height + j*width +k] = i*width*height + j*width + k;
+      }
+    }
+  }
+  hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(float)*8, 0,
+                       0, 0, hipChannelFormatKindFloat);
+  hipArray *arr;
+
+  HIP_CHECK(hipMalloc3DArray(&arr, &channelDesc,
+            make_hipExtent(width, height, depth), hipArrayDefault));
+  hipMemcpy3DParms myparms{0, {0, 0, 0}, {0, 0, 0, 0}, 0, {0, 0, 0},
+                           {0, 0, 0, 0}, {0, 0, 0}, hipMemcpyDefault};
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.srcPtr = make_hipPitchedPtr(Hptr, width * sizeof(float), width,
+                                      height);
+  myparms.dstArray = arr;
+  myparms.extent = extent;
+
+#ifdef __HIP_PLATFORM_NVIDIA__
+  myparms.kind = cudaMemcpyHostToDevice;
+#else
+  myparms.kind = hipMemcpyHostToDevice;
+#endif
+  // Host to Device copy
+  HIP_CHECK(hipMemcpy3D(&myparms));
+
+  // Device to Host copy
+  memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
+  myparms.srcPos = make_hipPos(0, 0, 0);
+  myparms.dstPos = make_hipPos(0, 0, 0);
+  myparms.dstPtr = make_hipPitchedPtr(HRes, width * sizeof(float), width,
+                                      height);
+  myparms.srcArray = arr;
+  myparms.extent = extent;
+#ifdef __HIP_PLATFORM_NVIDIA__
+  myparms.kind = cudaMemcpyDeviceToHost;
+#else
+  myparms.kind = hipMemcpyDeviceToHost;
+#endif
+  HIP_CHECK(hipMemcpy3D(&myparms));
+
+  for (int i = 0; i < depth; i++) {
+    for (int j = 0; j < height; j++) {
+      for (int k = 0; k < width; k++) {
+        if (HRes[i*width*height + j*width +k] != i*width*height + j*width + k) {
+          REQUIRE(false);
+        }
+      }
+    }
+  }
+  HIP_CHECK(hipFreeArray(arr));
+  free(Hptr);
+  free(HRes);
+}
+
+
+TEST_CASE("Unit_hipStrmPerThrdDefault") {
+  SECTION("Testing hipMemset/Memcpy() and their async version") {
+    REQUIRE(DefaultPT2_Memcpy_MemSet(1, 0));
+    REQUIRE(DefaultPT2_Memcpy_MemSet(1, 1));
+    REQUIRE(DefaultPT2_Memcpy_MemSet(0, 1));
+    REQUIRE(DefaultPT2_Memcpy_MemSet(0, 0));
+  }
+  SECTION("Testing hipMemset2D() and its async version") {
+    DefaultPT2_Memset2D(0);
+    DefaultPT2_Memset2D(1);
+  }
+
+  SECTION("Testing_hipMemset3D() and its async version") {
+    PerThrdDefltStrm_Memset3D(0);
+    PerThrdDefltStrm_Memset3D(1);
+  }
+
+  SECTION("Testing_hipStreamQuery()") {
+    DefaultPT2_StrmQuery();
+  }
+
+  SECTION("Testing_hipStreamSynchronize()") {
+    DefaultPT2_StreamSync();
+  }
+
+  SECTION("Testing_hipLaunchKernel()") {
+    // launch with null stream
+    Default_LaunchKernel(1);
+    // launch with user created stream
+    Default_LaunchKernel(0);
+  }
+
+  hipDeviceProp_t deviceProp;
+  HIP_CHECK(hipGetDeviceProperties(&deviceProp, 0));
+  if (deviceProp.cooperativeLaunch) {
+    SECTION("Testing_hipLaunchCooperativeKernel()") {
+      // launching hipLaunchCooperativeKernel() with Null stream
+      DefaultPT2_LaunchCooperativeKernel(1);
+      // launching hipLaunchCooperativeKernel() with user created stream
+      DefaultPT2_LaunchCooperativeKernel(0);
+    }
+  } else {
+    INFO("Cooperative Launch feature is not supported, therefore skipping");
+    INFO(" the test Testing_hipLaunchCooperativeKernel()");
+  }
+
+  SECTION("Testing_StrmWaitEvent()") {
+    DefaultPT2_StrmWaitEvent();
+  }
+
+  SECTION("Testing_hipStreamGetFlag()") {
+    DefaultPT2_StrmGetFlag();
+  }
+
+  SECTION("Testing_hipStreamGetPriority()") {
+    DefaultPT2_StrmGetPriority();
+  }
+
+  SECTION("Testing_hipMemcpyFrom/To/Symbol()") {
+    DefaultPT2_hipMemcpyFromSymbol();
+  }
+
+  SECTION("Testing_hipMemcpy2D() & its Async version") {
+    DefaultPT2_hipMemcpy2D(0);
+    DefaultPT2_hipMemcpy2D(1);
+  }
+
+  SECTION("Testing_hipMemcpy2DToArray()") {
+    DefaultPT2_hipMemcpy2DToArray();
+  }
+
+  SECTION("Testing_hipMemcpy3D()") {
+    DefaultPT2_hipMemcpy3D();
+  }
+
+  SECTION("Testing_hipEventQuery()") {
+    DefaultPT2_EvtQuery();
+  }
+}
@@ -20,6 +20,9 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip_test_checkers.hh>

+
+#if CUDA_VERSION < CUDA_12000
+
 #define SIZE_H 8
 #define SIZE_W 12
 #define TYPE_t float
@@ -77,3 +80,6 @@ TEST_CASE("Unit_hipBindTexture2D_Pitch") {
  HIP_CHECK(hipFree(devPtrA));
  HIP_CHECK(hipFree(devPtrB));
 }
+
+#endif // CUDA_VERSION < CUDA_12000
+
@@ -19,6 +19,8 @@ THE SOFTWARE.

 #include <hip_test_common.hh>

+#if CUDA_VERSION < CUDA_12000
+
 #define N 512

 texture<float, 1, hipReadModeElementType> tex;
@@ -79,3 +81,7 @@ TEST_CASE("Unit_hipBindTexture_tex1DfetchVerification") {
  HIP_CHECK(hipFree(texBuf));
  HIP_CHECK(hipFree(devBuf));
 }
+
+
+#endif // CUDA_VERSION < CUDA_12000
+
@@ -19,6 +19,9 @@ THE SOFTWARE.

 #include <hip_test_common.hh>

+
+#if CUDA_VERSION < CUDA_12000
+
 #define SIZE          10
 #define EPSILON       0.00001
 #define THRESH_HOLD   0.01  // For filter mode
@@ -159,3 +162,5 @@ TEST_CASE("Unit_hipNormalizedFloatValueTex_CheckModes") {
    runTest_hipTextureFilterMode<hipFilterModeLinear>();
  }
 }
+
+#endif  // CUDA_VERSION < CUDA_12000
@@ -20,6 +20,8 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip_test_checkers.hh>

+#if CUDA_VERSION < CUDA_12000
+
 typedef float T;

 // Texture reference for 2D Layered texture
@@ -107,3 +109,6 @@ TEST_CASE("Unit_hipSimpleTexture2DLayered_Check") {
  free(hData);
  free(hOutputData);
 }
+
+#endif // CUDA_VERSION < CUDA_12000
+
@@ -20,6 +20,8 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip_test_checkers.hh>

+#if CUDA_VERSION < CUDA_12000
+
 // Texture reference for 3D texture
 texture<float, hipTextureType3D, hipReadModeElementType> texf;
 texture<int, hipTextureType3D, hipReadModeElementType>   texi;
@@ -119,3 +121,7 @@ TEST_CASE("Unit_hipSimpleTexture3D_Check_DataTypes") {
    runSimpleTexture3D_Check<char>(i, i+1, i, &texc);
  }
 }
+
+
+
+#endif // CUDA_VERSION < CUDA_12000
@@ -19,6 +19,9 @@ THE SOFTWARE.

 #include <hip_test_common.hh>

+
+#if CUDA_VERSION < CUDA_12000
+
 texture<float, 2, hipReadModeElementType> tex;

 __global__ void tex2DKernel(float* outputData, int width) {
@@ -90,3 +93,7 @@ TEST_CASE("Unit_hipTextureRef2D_Check") {
  HIP_CHECK(hipFree(dData));
  HIP_CHECK(hipFreeArray(hipArray));
 }
+
+
+#endif // CUDA_VERSION < CUDA_12000
+
@@ -18,7 +18,7 @@
 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 # THE SOFTWARE.

-cmake_minimum_required(VERSION 2.8.3)
+cmake_minimum_required(VERSION 3.10)
 if (NOT DEFINED ROCM_PATH )
     set ( ROCM_PATH "/opt/rocm"  CACHE STRING "Default ROCM installation directory." )
 endif ()
@@ -45,16 +45,15 @@ endif()
 set(MY_SOURCE_FILES MatrixTranspose.cpp)
 set(MY_TARGET_NAME MatrixTranspose)
 set(MY_HIPCC_OPTIONS)
-set(MY_HCC_OPTIONS)
 set(MY_CLANG_OPTIONS)
 set(MY_NVCC_OPTIONS)

 set_source_files_properties(${MY_SOURCE_FILES} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
-hip_add_executable(${MY_TARGET_NAME} ${MY_SOURCE_FILES} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} HCC_OPTIONS ${MY_HCC_OPTIONS} CLANG_OPTIONS ${MY_CLANG_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS})
+hip_add_executable(${MY_TARGET_NAME} ${MY_SOURCE_FILES} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} CLANG_OPTIONS ${MY_CLANG_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS})

 # Search for rocm in common locations
 list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
-find_package(hip QUIET)
+find_package(hip QUIET CONFIG)
 if(TARGET hip::host)
    message(STATUS "Found hip::host at ${hip_DIR}")
    target_link_libraries(${MY_TARGET_NAME} hip::host)